Link error: multiple definition error of associate_descriptor when including <cuda/annotated_ptr>

[xuantengh]

When I try to include the header <cuda/annotated_ptr> in multiple compile units, the linker returns multiple definition errors indicating something gets repeatedly defined.

• source files

// a.cu
#include <cuda/annotated_ptr>

__global__ void hello();

int main() {
  hello<<<1, 1>>>();
  return 0;
}

// b.cu
#include <cstdio>
#include <cuda/annotated_ptr>

__global__ void hello() {
  printf("hello from (%d, %d)\n", blockIdx.x, threadIdx.x); 
}

NVIDIA GPU driver version: 470.82.01, nvcc version: 11.6.124 .
compilation command & errors:

nvcc a.cu b.cu -arch=sm_80 -o a.out

/tmp/tmpxft_00035cb1_00000000-18_b.o: In function `void* cuda::__4::__detail_ap::__associate_descriptor<unsigned long>(void*, unsigned long)':
tmpxft_00035cb1_00000000-10_b.cudafe1.cpp:(.text+0x16): multiple definition of `void* cuda::__4::__detail_ap::__associate_descriptor<unsigned long>(void*, unsigned long)'
/tmp/tmpxft_00035cb1_00000000-15_a.o:tmpxft_00035cb1_00000000-6_a.cudafe1.cpp:(.text+0x16): first defined here
/tmp/tmpxft_00035cb1_00000000-18_b.o: In function `void* cuda::__4::__detail_ap::__associate_descriptor<cuda::__4::access_property::shared>(void*, cuda::__4::access_property::shared)':
tmpxft_00035cb1_00000000-10_b.cudafe1.cpp:(.text+0x37): multiple definition of `void* cuda::__4::__detail_ap::__associate_descriptor<cuda::__4::access_property::shared>(void*, cuda::__4::access_property::shared)'
/tmp/tmpxft_00035cb1_00000000-15_a.o:tmpxft_00035cb1_00000000-6_a.cudafe1.cpp:(.text+0x37): first defined here
/tmp/tmpxft_00035cb1_00000000-18_b.o: In function `cuda::__4::discard_memory(void volatile*, unsigned long)':
tmpxft_00035cb1_00000000-10_b.cudafe1.cpp:(.text+0x54): multiple definition of `cuda::__4::discard_memory(void volatile*, unsigned long)'
/tmp/tmpxft_00035cb1_00000000-15_a.o:tmpxft_00035cb1_00000000-6_a.cudafe1.cpp:(.text+0x54): first defined here
collect2: error: ld returned 1 exit status

[wmaxey]

It looks like this is from a toolkit. Thanks for the report!

[gonzalobg]

#274 fixes these and some other ODR issues in the annotated_ptr and access_property implementation.