cuda nvcc 使 __device__ 有条件

Posted 2023-02-16

【中文标题】cuda nvcc 使 __device__ 有条件

【英文标题】：cuda nvcc make __device__ conditional

【发布时间】：2018-09-03 15:47:10

【问题描述】：

我正在尝试将 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;

在所有情况下，Wrapper 中的 call() 都是微不足道的，而 wrapper 本身是一个相当复杂的野兽（只有包含元信息的主机函数）。 条件编译不是一个选项，两个路径都打算并排使用。

我比“--disable-warnings”差了一步，因为老实说，复制和维护 10k loc 可怕的模板魔法的成本超过了警告的好处。

我会非常高兴有一种方法可以根据实现是针对 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

【问题讨论】：

This 可能感兴趣。

【参考方案1】：

这个问题其实是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();;
;

Similar question already asked

【参考方案2】：

很抱歉，您在滥用语言并误导读者。您的包装类具有__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。也许您需要在某处进行类型擦除。但这不行。

【参考方案3】：

我同时提出的解决方案是代码重复少得多，就是用函子级别替换调用：

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 Functorimpl.call_functor();;
    
;

//repeat for around 20 classes

Wrapper<ImplCPU> wrapCPU;
wrapCPU.call_functor()();