在编译时在cuda内核中使用__ldg时出错

Question

我的目标是利用我的应用程序中的缓存内存并搜索在线示例显示使用__ldg应该相对简单.

NVIDIA有GPU优化文档(在这里找到:https://www.olcf.ornl.gov/wp-content/uploads/2013/02/GPU_Opt_Fund-CW1.pdf),它提供了一个简单的例子:

__global__ void kernel ( int *output, int *input)
{
  ...
  output[idx] = __ldg( &input[idx] );
}

但是,当我尝试编译它时,我收到以下错误消息:

error: identifier "__ldg" is undefined.  

不幸的是,在Google上搜索此错误消息的解决方案无济于事.有什么建议这个简单的例子可能有什么问题吗？
是否有我遗漏的编译器标志？

作为参考,我的设备是计算能力3.5,我正在使用CUDA 5.5.

谢谢.

Answer 1

该__ldg() 固有只在计算能力3.5(或更新版本)架构可用.

这意味着:

1. 它必须在计算3.5(或更新)的GPU上运行
2. 必须针对计算3.5(或更新)GPU进行编译
3. 它也不能为旧架构编译.

这意味着:

1. 这不起作用: nvcc -arch=sm_30 ...
2. 这将有效: nvcc -arch=sm_35 ...
3. 这不起作用: nvcc -gencode arch=compute30,code=sm_30 -gencode arch=compute_35,code=sm_35 ...

Answer 2

对于__ldg一般实现任意类型并且正确地回退到小于3.5的计算能力的实现,请参阅BryanCatanzaro/generics Github项目.

这是一个裸骨模板:

template<typename T>
__device__ __forceinline__ T ldg(const T* ptr) {
#if __CUDA_ARCH__ >= 350
    return __ldg(ptr);
#else
    return *ptr;
#endif
}