错误：在 __device__/__global__ 函数中不允许引用纹理/表面变量

Question

我在使用 CUDA 的 .cu 文件中有以下代码：

#include "gpu_stgauss2.h"
#include "gpu_st.h"
#include "gpu_sampler.h"

static texture<float, 2, cudaReadModeElementType> s_texSRC1;
static texture<float4, 2, cudaReadModeElementType> s_texSRC4;

inline __host__ __device__ texture<float,2>& texSRC1() { return s_texSRC1; }
inline __host__ __device__ texture<float4,2>& texSRC4() { return s_texSRC4; }

static texture<float4, 2, cudaReadModeElementType> s_texST;
inline __host__ __device__ texture<float4,2>& texST() { return s_texST; }

它们稍后在同一个文件中使用如下：

gpu_image<float> gpu_stgauss2_filter( const gpu_image<float>& src, const gpu_image<float4>& st, 
                                      float sigma, float max_angle, bool adaptive,
                                      bool src_linear, bool st_linear, int order, float step_size,
                                      float precision )
{     
    if (sigma <= 0) return src;
    gpu_image<float> dst(src.size());

    gpu_sampler<float, texSRC1> src_sampler(src, src_linear? cudaFilterModeLinear : cudaFilterModePoint);
    float cos_max = cosf(radians(max_angle));

    if (src.size() == st.size()) {
        gpu_sampler<float4, texST> st_sampler(st, st_linear? cudaFilterModeLinear : cudaFilterModePoint);
        if (order == 1) imp_stgauss2_filter<1,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 2) imp_stgauss2_filter<2,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 4) imp_stgauss2_filter<4,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
    } else {
        float2 s = make_float2((float)st.w() / src.w(), (float)st.h() / src.h());
        gpu_resampler<float4, texST> st_sampler(st, s, st_linear? cudaFilterModeLinear : cudaFilterModePoint);
        if (order == 1) imp_stgauss2_filter<1,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 2) imp_stgauss2_filter<2,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 4) imp_stgauss2_filter<4,float><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
    }
    GPU_CHECK_ERROR();
    return dst;
}


gpu_image<float4> gpu_stgauss2_filter( const gpu_image<float4>& src, const gpu_image<float4>& st, 
                                       float sigma, float max_angle, bool adaptive,
                                       bool src_linear, bool st_linear, int order, float step_size,
                                       float precision )
{     
    if (sigma <= 0) return src;
    gpu_image<float4> dst(src.size());

    gpu_sampler<float4, texSRC4> src_sampler(src, src_linear? cudaFilterModeLinear : cudaFilterModePoint);
    float cos_max = cosf(radians(max_angle));

    if (src.size() == st.size()) {
        gpu_sampler<float4, texST> st_sampler(st, st_linear? cudaFilterModeLinear : cudaFilterModePoint);
        if (order == 1) imp_stgauss2_filter<1,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 2) imp_stgauss2_filter<2,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 4) imp_stgauss2_filter<4,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
    } else {
        float2 s = make_float2((float)st.w() / src.w(), (float)st.h() / src.h());
        gpu_resampler<float4, texST> st_sampler(st, s, st_linear? cudaFilterModeLinear : cudaFilterModePoint);
        if (order == 1) imp_stgauss2_filter<1,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 2) imp_stgauss2_filter<2,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
        else if (order == 4) imp_stgauss2_filter<4,float4><<<dst.blocks(), dst.threads()>>>(dst, src_sampler, st_sampler, sigma, cos_max, adaptive, step_size, precision);
    }
    GPU_CHECK_ERROR();
    return dst;
}

但是会导致如下错误：

error : taking reference of texture/surface variable not allowed in __device__/__global__ functions

我在 CUDA 方面的经验很少。任何人都可以帮助解决它吗？谢谢。

Answer 1

我在尝试编译完全相同的代码时遇到了同样的问题。事实证明，使用 ennetws 建议的“gpu_stbf2.cu”中的技巧毕竟没有必要在这里返回参考。

这3个函数实际上只是在这个文件中调用，所以将gpu_sampler.h中的struct definiton移回这里，不用调用这些函数获取纹理，直接使用即可。我已经把代码放在github上here。

Answer 2

对于任何有同样问题的人，在这种情况下来自here 的 GPU 库，我设法通过采用其他地方使用的相同策略来解决它，例如“gpu_stbf2.cu”。我设法使用 Cuda 6.0 和 Visual Studio 2012 x64 成功编译。

Answer 3

尝试将您的 CUDA 降级到 4.0。对于这样的代码语法，在 CUDA 4.0 中是可以的。我曾经遇到过类似的问题，CUDA 4.0 适合我。

Answer 4

我强烈建议在 cc 3.0 ++ 中使用无绑定纹理，因为取消绑定纹理命令不必同步主机线程

第二，你应该考虑使用CC 3.0+中提出的新的cash memory，为此请简单地将内存指定为

 const float* pArray; 

3，如果你坚持使用老式纹理，它对插值操作很强大。 在全局范围内：

   texture     <float, cudaTextureType1D> textureFloat32_1D; 

在你的代码中绑定纹理

cudaBindTexture (NULL, textureFloat32_1D, ...);

在内核内部随意使用纹理...

float a = tex1Dfatch(textureFloat32_1D , location) ;

内核之外

cudaUnbindTexture(textureFloat32_1D);

请注意，使用 CUDA 代码的多线程应用程序在使用与案例三中提到的相同纹理变量时会出现问题（它不受保护！）

Answer 5

编译器错误说明了一切：不允许你做你尝试过的事情。我建议直接使用变量（而不是通过texSRC1() 等访问它们，或者返回指针而不是引用。