Cuda Dynamic Parallelism trigger thrust error

I have a main function. it will call the two dynamic parallelism function. I use CMakeLists to compile the three functions separately. This is my project link to the github

when i run the command ( path main/build/ ) the executable ./test_rdc
this is my error(terminal):
hello1
----out—
terminate called after throwing an instance of ‘thrust::system::system_error’
what(): radix_sort: failed on 1st step: cudaErrorInvalidDeviceFunction: invalid device function
Aborted (core dumped)
this is my project tree
tree

CUDA version is 11.4 GPU:RTX3060 DRIVER: 470.141.03
Im not update thrust .The thrust is default in cuda-11.4/include

Invalid device function often means you are compiling for an incorrect architecture. I won’t be able to help with CMake.

If you want to create a reproducer that can fit in a forum post with ordinary compile commands, I’ll take a look as time permits.

In my CMakeLists.txt I use this command
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_86,code=sm_86;-std=c++14)
to set the arch and code sm_86. I think it is same like
nvcc x.cu --gpu-architecture=compute_86 --gpucode=sm_86.
thanks for your reply!

There isn’t much I can say based on the CMake input. I’m better at spotting trouble when I can see the actual compile commands.

FWIW I copied all your source files into a single directory, and built it as follows and had no runtime errors running on CUDA 11.4 on a V100:

$ ls
main      package_1      package_1.h          package_1_kernel.h  package_2.cpp  package_2_kernel.cu  README.md
main.cpp  package_1.cpp  package_1_kernel.cu  package_2           package_2.h    package_2_kernel.h
$ nvcc *.cu *.cpp -o test -arch=sm_70  -rdc=true -lcudadevrt
$ compute-sanitizer ./test
========= COMPUTE-SANITIZER
hello1
2parentkernel
hello2
----out---
========= ERROR SUMMARY: 0 errors
$

So my guess is there is a problem in your compilation sequence. Since that is generated by CMake in your case, I think this is basically a CMake question. I don’t use CMake. It might help if you compiled with CMake verbose output, and added that here. Someone may spot something.

Thanks for your reply!