Dynamic initialization is not supported for a __device__ variable when using libcuda

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello everyone! I want to run the code snippet which using cuda::binary_semaphore, but I got the error: dynamic initialization is not supported for a device variable.
nvcc: 11.1
device: rtx3090
Can anyone help? Thanks in advance!

2

compile with cuda 12.2. Make sure you specify -arch=sm_86 when compiling for your cc8.6 GPU.

3

Thanks for your reply!

4

In fact a little fiddling with compiler explorer suggest to me that CUDA 11.3 or newer will work.

5

I used to believe 11.1 is enough because of the libcu++ documentation.
By the way, I want to know how to create a device dynamic array of binary_semaphore? And static array of binary_semaphore like:

__device__ cuda::binary_semaphore<cuda::thread_scope_device> s[2]{1, 1};

seems to work well.
Many thanks!

6

I think that is probably correct for libcu. This particular issue is associated with dynamic initialization of a __device__ variable, which isn’t specific to libcu.

with CUDA 11.3, you could not initialize the variable:

__device__ cuda::binary_semaphore<cuda::thread_scope_device> s[2];

and instead initialize it (from host code) with cudaMemcpyToSymbol (before usage):

cuda::binary_semaphore<cuda::thread_scope_device> hs[2]{1,1};
cudaMemcpyToSymbol(s, hs, 2*sizeof(hs[0]));
7

Thanks for your timely reply! In fact, even statement like:

__device__ cuda::binary_semaphore<cuda::thread_scope_device> s

can cause error on CUDA11.1.
And now I am using CUDA11.7, I think my question about dynamic array can be solved by:

__device__ cuda::binary_semaphore<cuda::thread_scope_device>* s;
auto s_arr_h = (cuda::binary_semaphore<cuda::thread_scope_device>*)malloc(num * sizeof(cuda::binary_semaphore<cuda::thread_scope_device>));
cudaMemcpyToSymbol(s, s_arr_h, num*sizeof(s_arr_h[0]));

Am I right?

8

Correct, my suggestion was wrong. There is apparently a constructor that is initializing the object.

no, that won’t work. The only thing you have created a static allocation for via the __device__ variable is a pointer. You cannot copy an array of semaphores into that space allocated for a single pointer. You’ll need to revisit how pointers work. It’s a bit involved to have a bare pointer as a __device__ variable. To make any sensible use of it from the host requires a multi-stage copy.

9

Thanks for your suggestion.

10

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.