[Issue]: cannot build NVSHMEM examples with clang

[JamesMBartlett]

How is this issue impacting you?

Application crash

Share Your Debug Logs

Clang seems to be stricter about device/host overloading than nvcc so some of the headers fail to compile:

Details

/usr/include/nvshmem/device/nvshmem_defines.h:42:58: error: __device__ function 'nvshmem_team_my_pe' cannot overload __host__ __device__ function 'nvshmem_team_my_pe'
   42 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_my_pe(nvshmem_team_t team) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:465:32: note: previous declaration is here
  465 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_my_pe(nvshmem_team_t team);
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:46:58: error: __device__ function 'nvshmem_team_n_pes' cannot overload __host__ __device__ function 'nvshmem_team_n_pes'
   46 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_n_pes(nvshmem_team_t team) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:466:32: note: previous declaration is here
  466 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_n_pes(nvshmem_team_t team);
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:50:58: error: __device__ function 'nvshmem_team_translate_pe' cannot overload __host__ __device__ function 'nvshmem_team_translate_pe'
   50 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_translate_pe(
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:469:32: note: previous declaration is here
  469 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_translate_pe(nvshmem_team_t src_team, int src_pe,
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:55:58: error: __device__ function 'nvshmem_my_pe' cannot overload __host__ __device__ function 'nvshmem_my_pe'
   55 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_my_pe(void) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:74:32: note: previous declaration is here
   74 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_my_pe();
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:59:58: error: __device__ function 'nvshmem_n_pes' cannot overload __host__ __device__ function 'nvshmem_n_pes'
   59 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_n_pes(void) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:75:32: note: previous declaration is here
   75 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_n_pes();
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:63:59: error: __device__ function 'nvshmem_info_get_name' cannot overload __host__ __device__ function 'nvshmem_info_get_name'
   63 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmem_info_get_name(char *name) {
      |                                                           ^
/usr/include/nvshmem/host/nvshmem_api.h:77:33: note: previous declaration is here
   77 | NVSHMEMI_HOSTDEVICE_PREFIX void nvshmem_info_get_name(char *name);
      |                                 ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:73:59: error: __device__ function 'nvshmem_info_get_version' cannot overload __host__ __device__ function 'nvshmem_info_get_version'
   73 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmem_info_get_version(int *major,
      |                                                           ^
/usr/include/nvshmem/host/nvshmem_api.h:76:33: note: previous declaration is here
   76 | NVSHMEMI_HOSTDEVICE_PREFIX void nvshmem_info_get_version(int *major, int *minor);
      |                                 ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_bfloat16_p' cannot overload __host__ __device__ function 'nvshmem_bfloat16_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_half_p' cannot overload __host__ __device__ function 'nvshmem_half_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_float_p' cannot overload __host__ __device__ function 'nvshmem_float_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]

In addition, clang doesn't make assert available by default, requiring a cassert.h include:

Details

In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
In file included from /usr/include/nvshmem/device/nvshmem_defines.h:29:
In file included from /usr/include/nvshmem/non_abi/device/pt-to-pt/nvshmemi_transfer_api.cuh:8:
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:56:13: error: use of undeclared identifier 'assert'
   56 |             assert(0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:79:13: error: use of undeclared identifier 'assert'
   79 |             assert(0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:101:13: error: use of undeclared identifier 'assert'
  101 |             assert((blockDim.x * blockDim.y * blockDim.z) % (4 * warpSize) == 0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:103:13: error: use of undeclared identifier 'assert'
  103 |             assert((barrierId < 16) && "Reduce blocksize");
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:111:13: error: use of undeclared identifier 'assert'
  111 |             assert(0);
      |             ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
In file included from /usr/include/nvshmem/device/nvshmem_defines.h:33:
In file included from /usr/include/nvshmem/non_abi/device/common/nvshmemi_common_device.cuh:28:
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:89:9: error: use of undeclared identifier 'assert'
   89 |         assert(rc == PROXY_GLOBAL_EXIT_INIT);
      |         ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:230:5: error: use of undeclared identifier 'assert'
  230 |     assert(0);
      |     ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:259:9: error: use of undeclared identifier 'assert'
  259 |         assert(pred_pe && pred_contigous);
      |         ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:530:17: error: use of undeclared identifier 'assert'
  530 |                 assert(false);
      |                 ^

Steps to Reproduce the Issue

I attempted to build one of the examples with clang:

clang++-20 -x cuda --cuda-gpu-arch=sm_100 \
  -I/usr/include/nvshmem/ \
  -I/usr/local/cuda/include/ \
  -L/usr/lib/x86_64-linux-gnu/nvshmem/12/ \
  -L/usr/local/cuda/lib64 \
  -lnvshmem_device -lnvshmem_host -lcudart -lcudadevrt \
  --cuda-path=/usr/local/cuda \
  examples/on-stream.cu -o a.out

NVSHMEM Version

3.4.5-1

Your platform details

Ubuntu linux.

This is a compilation issue so I don't think nvidia-smi/other GPU details are necessary.

Error Message & Behavior

Details

In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
In file included from /usr/include/nvshmem/device/nvshmem_defines.h:29:
In file included from /usr/include/nvshmem/non_abi/device/pt-to-pt/nvshmemi_transfer_api.cuh:8:
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:56:13: error: use of undeclared identifier 'assert'
   56 |             assert(0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:79:13: error: use of undeclared identifier 'assert'
   79 |             assert(0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:101:13: error: use of undeclared identifier 'assert'
  101 |             assert((blockDim.x * blockDim.y * blockDim.z) % (4 * warpSize) == 0);
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:103:13: error: use of undeclared identifier 'assert'
  103 |             assert((barrierId < 16) && "Reduce blocksize");
      |             ^
/usr/include/nvshmem/non_abi/device/threadgroup/nvshmemi_common_device_defines.cuh:111:13: error: use of undeclared identifier 'assert'
  111 |             assert(0);
      |             ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
In file included from /usr/include/nvshmem/device/nvshmem_defines.h:33:
In file included from /usr/include/nvshmem/non_abi/device/common/nvshmemi_common_device.cuh:28:
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:89:9: error: use of undeclared identifier 'assert'
   89 |         assert(rc == PROXY_GLOBAL_EXIT_INIT);
      |         ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:230:5: error: use of undeclared identifier 'assert'
  230 |     assert(0);
      |     ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:259:9: error: use of undeclared identifier 'assert'
  259 |         assert(pred_pe && pred_contigous);
      |         ^
/usr/include/nvshmem/non_abi/device/pt-to-pt/proxy_device.cuh:530:17: error: use of undeclared identifier 'assert'
  530 |                 assert(false);
      |                 ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:42:58: error: __device__ function 'nvshmem_team_my_pe' cannot overload __host__ __device__ function 'nvshmem_team_my_pe'
   42 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_my_pe(nvshmem_team_t team) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:465:32: note: previous declaration is here
  465 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_my_pe(nvshmem_team_t team);
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:46:58: error: __device__ function 'nvshmem_team_n_pes' cannot overload __host__ __device__ function 'nvshmem_team_n_pes'
   46 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_n_pes(nvshmem_team_t team) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:466:32: note: previous declaration is here
  466 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_n_pes(nvshmem_team_t team);
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:50:58: error: __device__ function 'nvshmem_team_translate_pe' cannot overload __host__ __device__ function 'nvshmem_team_translate_pe'
   50 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_team_translate_pe(
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:469:32: note: previous declaration is here
  469 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_team_translate_pe(nvshmem_team_t src_team, int src_pe,
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:55:58: error: __device__ function 'nvshmem_my_pe' cannot overload __host__ __device__ function 'nvshmem_my_pe'
   55 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_my_pe(void) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:74:32: note: previous declaration is here
   74 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_my_pe();
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:59:58: error: __device__ function 'nvshmem_n_pes' cannot overload __host__ __device__ function 'nvshmem_n_pes'
   59 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE int nvshmem_n_pes(void) {
      |                                                          ^
/usr/include/nvshmem/host/nvshmem_api.h:75:32: note: previous declaration is here
   75 | NVSHMEMI_HOSTDEVICE_PREFIX int nvshmem_n_pes();
      |                                ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:63:59: error: __device__ function 'nvshmem_info_get_name' cannot overload __host__ __device__ function 'nvshmem_info_get_name'
   63 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmem_info_get_name(char *name) {
      |                                                           ^
/usr/include/nvshmem/host/nvshmem_api.h:77:33: note: previous declaration is here
   77 | NVSHMEMI_HOSTDEVICE_PREFIX void nvshmem_info_get_name(char *name);
      |                                 ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:73:59: error: __device__ function 'nvshmem_info_get_version' cannot overload __host__ __device__ function 'nvshmem_info_get_version'
   73 | NVSHMEMI_DEVICE_PREFIX NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmem_info_get_version(int *major,
      |                                                           ^
/usr/include/nvshmem/host/nvshmem_api.h:76:33: note: previous declaration is here
   76 | NVSHMEMI_HOSTDEVICE_PREFIX void nvshmem_info_get_version(int *major, int *minor);
      |                                 ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_bfloat16_p' cannot overload __host__ __device__ function 'nvshmem_bfloat16_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_half_p' cannot overload __host__ __device__ function 'nvshmem_half_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
In file included from examples/on-stream.cu:15:
In file included from /usr/include/nvshmem/nvshmem.h:23:
/usr/include/nvshmem/device/nvshmem_defines.h:85:38: error: __device__ function 'nvshmem_float_p' cannot overload __host__ __device__ function 'nvshmem_float_p'
   85 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_TYPENAME_P_IMPL)
      |                                      ^
/usr/include/nvshmem/host/nvshmem_api.h:202:38: note: previous declaration is here
  202 | NVSHMEMI_REPT_FOR_STANDARD_RMA_TYPES(NVSHMEMI_DECL_TYPE_P)
      |                                      ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated when compiling for sm_100.

[JamesMBartlett]

If I attempt to build the host side example, hello.cpp. Then I get a different issue:

clang++-20 -v --cuda-path=/usr/local/cuda-12 -I /usr/include/nvshmem/ -I /usr/local/cuda/include/  -L/usr/lib/x86_64-linux-gnu/nvshmem/12/ -L /usr/local/cuda/lib64 -lnvshmem_device -lnvshmem_host -lcudart -fuse-ld=lld examples/hello.cpp

ld.lld: error: undefined symbol: __cudaRegisterLinkedBinary_8d6d1ab6_14_init_device_cu_4920c2c8
>>> referenced by tmpxft_00126787_00000000-6_init_device.compute_90.cudafe1.cpp
>>>               init_device.cu.o:(__sti____cudaRegisterAll()) in archive /usr/lib/x86_64-linux-gnu/nvshmem/12/libnvshmem_device.a
clang++-20: error: linker command failed with exit code 1 (use -v to see invocation)

which I think is related to clang not properly handling the relocatable device code that libnvshmem_device.a is built with.