cuda nvcc使__device__具有条件
我正在尝试添加一个cuda后端到一个20k的loc c++表达式模板库。到目前为止，它工作得很好，但我完全被“警告:从__host__ __device__函数调用__host__函数是不允许的”警告淹没了。
大部分代码可以总结如下：
template<class Impl>
struct Wrapper{
Impl impl;
// lots and lots of decorator code
__host__ __device__ void call(){ impl.call();};
};
//Guaranteed to never ever be used on gpu.
struct ImplCPU{
void call();
};
//Guaranteed to never ever be used on cpu.
struct ImplGPU{
__host__ __device__ void call();//Actually only __device__, but needed to shut up the compiler as well
};
Wrapper<ImplCPU> wrapCPU;
Wrapper<ImplGPU> wrapGPU;在所有情况下，包装器中的call()都是微不足道的，而包装器本身是一个相当复杂的野兽(只有包含元信息的主机函数)。条件编译不是一个选项，两个路径都应该同时使用。
我离“--禁用-警告”还差一步，因为老实说，复制和维护10k个可怕的模板魔法的成本超过了警告的好处。
我非常高兴有一种方法可以根据实现是针对gpu还是针对cpu来有条件地调用device或host (因为Impl知道它的用途)。
只是为了显示它是坏的。一条警告：
/home/user/Remora/include/remora/detail/matrix_expression_classes.hpp(859): warning: calling a __host__ function from a __host__ __device__ function is not allowed
detected during:
instantiation of "remora::matrix_matrix_prod<MatA, MatB>::size_type remora::matrix_matrix_prod<MatA, MatB>::size1() const [with MatA=remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, MatB=remora::matrix<float, remora::column_major, remora::hip_tag>]"
/home/user/Remora/include/remora/cpu/../assignment.hpp(258): here
instantiation of "MatA &remora::assign(remora::matrix_expression<MatA, Device> &, const remora::matrix_expression<MatB, Device> &) [with MatA=remora::dense_matrix_adaptor<float, remora::row_major, remora::continuous_dense_tag, remora::hip_tag>, MatB=remora::matrix_matrix_prod<remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, remora::matrix<float, remora::column_major, remora::hip_tag>>, Device=remora::hip_tag]"
/home/user/Remora/include/remora/cpu/../assignment.hpp(646): here
instantiation of "remora::noalias_proxy<C>::closure_type &remora::noalias_proxy<C>::operator=(const E &) [with C=remora::matrix<float, remora::row_major, remora::hip_tag>, E=remora::matrix_matrix_prod<remora::dense_triangular_proxy<const float, remora::row_major, remora::lower, remora::hip_tag>, remora::matrix<float, remora::column_major, remora::hip_tag>>]"
/home/user/Remora/Test/hip_triangular_prod.cpp(325): here
instantiation of "void Remora_hip_triangular_prod::triangular_prod_matrix_matrix_test(Orientation) [with Orientation=remora::row_major]"
/home/user/Remora/Test/hip_triangular_prod.cpp(527): here回答 3
Stack Overflow用户
发布于 2018-09-06 21:42:40
这个问题实际上是CUDA语言扩展中非常不幸的缺陷。
处理这些警告的标准方法(在推力和类似的模板化CUDA库中)是通过使用#pragma hd_warning_disable或在较新的CUDA9.0或更高版本的#pragma nv_exec_check_disable中禁用导致警告的函数/方法。
因此在您的情况下，它将是：
template<class Impl>
struct Wrapper{
Impl impl;
// lots and lots of decorator code
#pragma nv_exec_check_disable
__host__ __device__ void call(){ impl.call();};
};Stack Overflow用户
发布于 2018-09-04 07:30:17
很抱歉，你在滥用语言，误导读者。包装类有一个__host__ __device__方法并不是真的；您的意思是说它有一个__host__方法或一个__device__方法。您应该更多地将警告视为错误。
因此，您不能只使用ImplCPU和ImplGPU的示例模板实例化；但是-您可以这样做吗？
template<typename Impl> struct Wrapper;
template<> struct Wrapper<ImplGPU> {
ImplGPU impl;
__device__ void call(){ impl.call();};
}
template<> struct Wrapper<ImplCPU> {
ImplGPU impl;
__host__ void call(){ impl.call();};
}或者，如果你想变得更加书生气，它可以是：
enum implementation_device { CPU, GPU };
template<implementation_device ImplementationDevice> Wrapper;
template<> Wrapper<CPU> {
__host__ void call();
}
template<> Wrapper<GPU> {
__device__ void call();
}话虽如此，您原本希望使用单个包装器类，但在这里我要告诉您，您不能这样做。我怀疑您的问题是一个X-Y problem，您应该真正考虑使用该包装器的整个方法。也许你需要将使用它的代码模板化，不同于CPU或GPU。也许你需要在某个地方擦除类型。但这并不管用。
Stack Overflow用户
发布于 2018-09-04 16:25:51
我在此期间提出的解决方案是用函数器级别替换call，从而减少代码重复：
template<class Impl, class Device>
struct WrapperImpl;
template<class Impl>
struct WrapperImpl<Impl, CPU>{
typename Impl::Functor f;
__host__ operator()(){ f();}
};
//identical to CPU up to __device__
template<class Impl>
struct WrapperImpl<Impl, GPU>{
typename Impl::Functor f;
__device__ operator()(){ f();}
};
template<class Impl>
struct Wrapper{
typedef WrapperImpl<Impl, typename Impl::Device> Functor;
Impl impl;
// lots and lots of decorator code that i now do not need to duplicate
Functor call_functor()const{
return Functor{impl.call_functor();};
}
};
//repeat for around 20 classes
Wrapper<ImplCPU> wrapCPU;
wrapCPU.call_functor()();https://stackoverflow.com/questions/52152898
复制相似问题