Compatibility with Thrust?

jiapei-nexera commented 2 years ago

Describe the bug Current stdgpu seems to be NOT compatible with NVidia Thurst?

Steps to reproduce

git clone https://github.com/stotko/stdgpu.git
mkdir build
cd build
cmake ../

Expected behavior Successfully built and run.

Actual behavior No matter GCC or Clang, both failed.

GCC 11.2

[  1%] Building CXX object src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o
cd ....../stdgpu/build/src/stdgpu && /usr/bin/c++ -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -Dstdgpu_EXPORTS -I....../stdgpu/src/stdgpu/.. -I....../stdgpu/build/src/stdgpu/include -isystem /usr/local/cuda/include -fPIC -Wall -pedantic -Wextra -Wshadow -Wsign-compare -Wconversion -Wfloat-equal -Wundef -Wdouble-promotion -MD -MT src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o -MF CMakeFiles/stdgpu.dir/impl/iterator.cpp.o.d -o CMakeFiles/stdgpu.dir/impl/iterator.cpp.o -c ....../stdgpu/src/stdgpu/impl/iterator.cpp
In file included from /usr/local/cuda/include/nv/detail/__target_macros:13,
             from /usr/local/cuda/include/nv/target:195,
             from /usr/local/cuda/include/cub/detail/device_synchronize.cuh:23,
             from /usr/local/cuda/include/thrust/system/cuda/detail/util.h:36,
             from /usr/local/cuda/include/thrust/system/cuda/detail/malloc_and_free.h:26,
             from /usr/local/cuda/include/thrust/system/detail/adl/malloc_and_free.h:42,
             from /usr/local/cuda/include/thrust/system/detail/generic/memory.inl:22,
             from /usr/local/cuda/include/thrust/system/detail/generic/memory.h:69,
             from /usr/local/cuda/include/thrust/detail/reference.h:28,
             from ....../stdgpu/src/stdgpu/../stdgpu/iterator.h:30,
             from ....../stdgpu/src/stdgpu/impl/iterator.cpp:16:
/usr/local/cuda/include/cub/util_device.cuh: In function ‘cudaError_t cub::PtxVersionUncached(int&)’:
/usr/local/cuda/include/cub/util_device.cuh:368:15: error: invalid conversion from ‘EmptyKernelPtr’ {aka ‘void (*)()’} to ‘const void*’ [-fpermissive]
368 |           if (CubDebug(result = cudaFuncGetAttributes(&empty_kernel_attrs,
  |               ^~~~~~~~
  |               |
  |               EmptyKernelPtr {aka void (*)()}
In file included from /usr/local/cuda/include/thrust/system/cuda/detail/guarded_cuda_runtime_api.h:38,
             from /usr/local/cuda/include/thrust/system/cuda/detail/malloc_and_free.h:19,
             from /usr/local/cuda/include/thrust/system/detail/adl/malloc_and_free.h:42,
             from /usr/local/cuda/include/thrust/system/detail/generic/memory.inl:22,
             from /usr/local/cuda/include/thrust/system/detail/generic/memory.h:69,
             from /usr/local/cuda/include/thrust/detail/reference.h:28,
             from ....../stdgpu/src/stdgpu/../stdgpu/iterator.h:30,
             from ....../stdgpu/src/stdgpu/impl/iterator.cpp:16:
/usr/local/cuda/include/cuda_runtime_api.h:4337:125: note:   initializing argument 2 of ‘cudaError_t cudaFuncGetAttributes(cudaFuncAttributes*, const void*)’
4337 | extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const void *func);
  |                                                                                                                 ~~~~~~~~~~~~^~~~
make[2]: *** [src/stdgpu/CMakeFiles/stdgpu.dir/build.make:93: src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o] Error 1
make[2]: Leaving directory '....../stdgpu/build'
make[1]: *** [CMakeFiles/Makefile2:318: src/stdgpu/CMakeFiles/stdgpu.dir/all] Error 2
make[1]: Leaving directory '....../stdgpu/build'
make: *** [Makefile:149: all] Error 2

clang 14.0

[  1%] Building CXX object src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o
cd ....../build/src/stdgpu && /usr/bin/clang++ -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -I....../src/stdgpu/.. -I....../build/src/stdgpu/include -isystem /usr/local/cuda/include -Wall -pedantic -Wextra -Wshadow -Wsign-compare -Wconversion -Wfloat-equal -Wundef -Wdouble-promotion -MD -MT src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o -MF CMakeFiles/stdgpu.dir/impl/iterator.cpp.o.d -o CMakeFiles/stdgpu.dir/impl/iterator.cpp.o -c ....../src/stdgpu/impl/iterator.cpp
In file included from ....../src/stdgpu/impl/iterator.cpp:16:
In file included from ....../src/stdgpu/../stdgpu/iterator.h:30:
In file included from /usr/local/cuda/include/thrust/detail/reference.h:28:
In file included from /usr/local/cuda/include/thrust/system/detail/generic/memory.h:69:
In file included from /usr/local/cuda/include/thrust/system/detail/generic/memory.inl:22:
In file included from /usr/local/cuda/include/thrust/system/detail/adl/malloc_and_free.h:42:
In file included from /usr/local/cuda/include/thrust/system/cuda/detail/malloc_and_free.h:26:
In file included from /usr/local/cuda/include/thrust/system/cuda/detail/util.h:38:
/usr/local/cuda/include/cub/util_device.cuh:368:33: error: no matching function for call to 'cudaFuncGetAttributes'
      if (CubDebug(result = cudaFuncGetAttributes(&empty_kernel_attrs,
                            ^~~~~~~~~~~~~~~~~~~~~
/usr/local/cuda/include/cub/util_debug.cuh:115:64: note: expanded from macro 'CubDebug'
#define CubDebug(e) CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)
                                                           ^
/usr/local/cuda/include/nv/detail/__target_macros:455:78: note: expanded from macro 'NV_IF_TARGET'
#  define NV_IF_TARGET(cond, t, ...)    _NV_BLOCK_EXPAND(_NV_TARGET_IF(cond, t, __VA_ARGS__))
                                                                         ^
/usr/local/cuda/include/nv/detail/__target_macros:419:74: note: expanded from macro '_NV_TARGET_IF'
#    define _NV_TARGET_IF(cond, t, ...) _NV_IF( _NV_ARCH_COND_CAT(cond), t, __VA_ARGS__)
                                                                     ^
note: (skipping 24 expansions in backtrace; use -fmacro-backtrace-limit=0 to see all)
/usr/local/cuda/include/nv/detail/__preprocessor:83:47: note: expanded from macro '_NV_STRIP_PAREN'
#define _NV_STRIP_PAREN(...) _NV_STRIP_PAREN1(__VA_ARGS__)
                                          ^~~~~~~~~~~
/usr/local/cuda/include/nv/detail/__preprocessor:82:48: note: expanded from macro '_NV_STRIP_PAREN1'
#define _NV_STRIP_PAREN1(...) _NV_STRIP_PAREN2 __VA_ARGS__
                                           ^~~~~~~~~~~
/usr/local/cuda/include/nv/detail/__preprocessor:81:31: note: expanded from macro '_NV_STRIP_PAREN2'
#define _NV_STRIP_PAREN2(...) __VA_ARGS__
                          ^~~~~~~~~~~
/usr/local/cuda/include/cuda_runtime_api.h:4337:58: note: candidate function not viable: no known conversion from 'EmptyKernelPtr' (aka 'void (*)()') to 'const void *' for 2nd argument; take the address of the argument with &
extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const void *func);
                                                     ^
1 error generated.
make[2]: *** [src/stdgpu/CMakeFiles/stdgpu.dir/build.make:93: src/stdgpu/CMakeFiles/stdgpu.dir/impl/iterator.cpp.o] Error 1
make[2]: Leaving directory '....../build'
make[1]: *** [CMakeFiles/Makefile2:403: src/stdgpu/CMakeFiles/stdgpu.dir/all] Error 2
make[1]: Leaving directory '....../build'
make: *** [Makefile:149: all] Error 2

System (please complete the following information):

OS: Ubuntu 22.04
Compiler: gcc 11.2 and clang 14.0, both tried
Backend: CUDA 11.6
Library version: stdgpu 1.3.0 ??

stotko commented 2 years ago

I locally tested building with GCC 11.2 + CUDA 11.7 on Ubuntu 22.04 which worked without any issues. According to the CUDA download website (network install method), CUDA 11.6 is only available for Ubuntu 18.04 and 20.04 which means that the most recent Ubuntu 22.04 is not supported by this version.

Could you upgrade to CUDA 11.7 and try building the libary again?

jiapei-nexera commented 2 years ago

@stotko

Looks like somebody else met the same issue as mine ?

https://github.com/allisonvacanti/cub/commit/800c92a219121804009ddfc388ce7151ac5251d9

And, tested... By this single modification, stdgpu built successfully...

stotko commented 2 years ago

Thanks for testing and the pointer to cub. Looking at the list of cub releases, it seems like version 1.15 is affected by this bug and shipped with CUDA 11.6, whereas version 1.16 contains the mentioned fix and seems to be part of CUDA 11.7. So this supports the observed behavior.

Closing this issue as this bug is not directly caused by stdgpu.

stotko / stdgpu

Compatibility with Thrust? #310