cuda nvcc使__device__有条件

内容来源于 Stack Overflow,并遵循CC BY-SA 3.0许可协议进行翻译与使用

  • 回答 (1)
  • 关注 (0)
  • 查看 (252)

我正在尝试将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__方法。您应该将警告视为更多错误。

因此,您不能只使用示例模板实例化ImplCPUImplGPU; 但是 - 你可以这样做吗?

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类,在这里我告诉您,您不能这样做。我怀疑你的问题是一个XY问题,你应该考虑使用该包装器的整个方法。也许你需要让使用它的代码模板化不同的CPU或GPU。也许你需要在某处进行类型擦除。但这不行。

所属标签

可能回答问题的人

  • 应用案例分享

    1 粉丝490 提问5 回答
  • o o

    4 粉丝495 提问5 回答
  • uncle_light

    5 粉丝518 提问4 回答
  • 找虫虫

    5 粉丝0 提问3 回答

扫码关注云+社区

领取腾讯云代金券