4

My goal is to take advantage of cache memory in my application and searching for online examples shows that using __ldg should be relatively straightforward.

NVIDIA has documentation for GPU optimization (found here: https://www.olcf.ornl.gov/wp-content/uploads/2013/02/GPU_Opt_Fund-CW1.pdf) which provides the straightforward example:

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

However when I try to compile this I get the following error message:

error: identifier "__ldg" is undefined.  

Searching Google for a solution to this error message has been unfortunately unhelpful. Any suggestions what may be wrong with this simple example?
Is there a compiler flag that I am missing?

For reference my device is compute capability 3.5 and I am working with CUDA 5.5.

Thank you.

3
  • 4
    What is your compile command line? To take advantage of a cc3.5 device, you need to compile for that architecture with -arch=sm_35 or similar. Jun 5, 2014 at 20:27
  • I'm fairly certain that this was the issue. I was using sample code from nVidia and their makefile calls multiple flags and I thought it would just use the highest capability available, but removing the lower compute capability flags resulted in correct compilation. Thanks. Jun 5, 2014 at 23:11
  • @user3280204 if compiling for the correct architecture solved the problem, I believe you should accept the answer given below.
    – BRabbit27
    Aug 7, 2014 at 8:12

2 Answers 2

10

The __ldg() intrinsic is only available on compute capability 3.5 (or newer) architecture.

That means:

  1. It must be run on a compute 3.5 (or newer) GPU
  2. It must be compiled for a compute 3.5 (or newer) GPU
  3. It cannot also be compiled for an older architecture.

That means:

  1. This won't work: nvcc -arch=sm_30 ...
  2. This will work: nvcc -arch=sm_35 ...
  3. This won't work: nvcc -gencode arch=compute30,code=sm_30 -gencode arch=compute_35,code=sm_35 ...
7

For an implementation of __ldg that generalizes to arbitrary types and correctly falls back on compute capability less than 3.5, see the BryanCatanzaro/generics Github project.

Here is a bare bones template:

template<typename T>
__device__ __forceinline__ T ldg(const T* ptr) {
#if __CUDA_ARCH__ >= 350
    return __ldg(ptr);
#else
    return *ptr;
#endif
}
1
  • But template<typename T> __device__ __forceinline__ T __ldg(const T* ptr) { Jan 2, 2022 at 12:35

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.