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

Question

我的目标是利用应用程序中的缓存内存，搜索在线示例表明使用__ldg应该相对简单。

NVIDIA提供了用于GPU优化的文档(见此处：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

有关泛化为任意类型并正确返回到小于3.5的计算能力的__ldg的实现，请参见BryanCatanzaro/泛型 Github项目。

下面是一个裸露的骨头模板：

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