cuda nvcc使__device__有条件

问题描述 投票:1回答:3

我正在尝试将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
cuda nvcc
3个回答
1
投票

对不起,你是在滥用语言和误导读者。你的包装类有一个__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类,在这里我告诉您,您不能这样做。我怀疑你的问题提出了X-Y problem,你应该考虑使用该包装器的整个方法。也许你需要让使用它的代码模板化不同的CPU或GPU。也许你需要在某处进行类型擦除。但这不行。


0
投票

我提出的解决方案平均时间远远少于代码重复,是用仿函数级别替换调用:

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()();

0
投票

这个问题实际上是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

© www.soinside.com 2019 - 2024. All rights reserved.