我正在尝试将cuda后端添加到20k loc c ++表达式模板库中。到目前为止它工作得很好,但我完全被淹没了“警告:不允许从__host__
函数调用__host__ __device__
函数”警告。
大多数代码可以这样总结:
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;
在所有情况下,Wrapper中的call()都是微不足道的,而包装器本身是一个相当复杂的野兽(只有包含元信息的主机函数)。条件编译不是一种选择,两种路径都是并排使用的。
我只是“ - 禁用 - 警告”的一步,因为老实说,复制和维护10k loc的可怕模板魔法的成本超过了警告的好处。
根据实现是针对gpu还是cpu(因为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
对不起,你是在滥用语言和误导读者。你的包装类有一个__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();
}
话虽如此 - 您期望使用单个Wrapper类,在这里我告诉您,您不能这样做。我怀疑你的问题提出了X-Y problem,你应该考虑使用该包装器的整个方法。也许你需要让使用它的代码模板化不同的CPU或GPU。也许你需要在某处进行类型擦除。但这不行。
我提出的解决方案平均时间远远少于代码重复,是用仿函数级别替换调用:
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()();
这个问题实际上是CUDA语言扩展中非常不幸的缺陷。
处理这些警告的标准方法(在Thrust和类似的模板化CUDA库中)是通过使用#pragma hd_warning_disable
或更新的CUDA(9.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();};
};