Initializing cuda global variable

Accelerated Computing CUDA CUDA Programming and Performance
#1 
__constant__ const unsigned int *ff = (const unsigned int[]){90, 50, 100};


int main()
{
}

Compiling:

n

vcc ./test.cu
./test.cu(1): error: identifier "__T20" is undefined in device code

1 error detected in the compilation of "/tmp/tmpxft_0000785f_00000000-10_test.cpp2.i".

Verbose compiling:

nvcc --verbose ./test.cu
    #$ _SPACE_= 
    #$ _CUDART_=cudart
    #$ _HERE_=/usr/lib/nvidia-cuda-toolkit/bin
    #$ _THERE_=/usr/lib/nvidia-cuda-toolkit/bin
    #$ _TARGET_SIZE_=
    #$ _TARGET_DIR_=
    #$ _TARGET_SIZE_=64
    #$ NVVMIR_LIBRARY_DIR=/usr/lib/nvidia-cuda-toolkit/libdevice
    #$ PATH=/usr/lib/nvidia-cuda-toolkit/bin:/home/kasha/bin:/home/kasha/.local/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin
    #$ LIBRARIES=  -L/usr/lib/x86_64-linux-gnu/stubs
    #$ gcc -D__CUDA_ARCH__=200 -E -x c++        -DCUDA_DOUBLE_MATH_FUNCTIONS  -D__CUDACC__ -D__NVCC__  -D"__CUDACC_VER__=70517" -D"__CUDACC_VER_BUILD__=17" -D"__CUDACC_VER_MINOR__=5" -D"__CUDACC_VER_MAJOR__=7" -include "cuda_runtime.h" -m64 "./test.cu" > "/tmp/tmpxft_0000799b_00000000-9_test.cpp1.ii" 
    #$ cudafe --allow_managed --m64 --gnu_version=50400 -tused --no_remove_unneeded_entities --gen_c_file_name "/tmp/tmpxft_0000799b_00000000-4_test.cudafe1.c" --stub_file_name "/tmp/tmpxft_0000799b_00000000-4_test.cudafe1.stub.c" --gen_device_file_name "/tmp/tmpxft_0000799b_00000000-4_test.cudafe1.gpu" --nv_arch "compute_20" --gen_module_id_file --module_id_file_name "/tmp/tmpxft_0000799b_00000000-3_test.module_id" --include_file_name "tmpxft_0000799b_00000000-2_test.fatbin.c" "/tmp/tmpxft_0000799b_00000000-9_test.cpp1.ii" 
    #$ gcc -D__CUDA_ARCH__=200 -E -x c        -DCUDA_DOUBLE_MATH_FUNCTIONS  -D__CUDACC__ -D__NVCC__ -D__CUDANVVM__  -D__CUDA_PREC_DIV -D__CUDA_PREC_SQRT -m64 "/tmp/tmpxft_0000799b_00000000-4_test.cudafe1.gpu" > "/tmp/tmpxft_0000799b_00000000-10_test.cpp2.i" 
    #$ cudafe -w --allow_managed --m64 --gnu_version=50400 --c --gen_c_file_name "/tmp/tmpxft_0000799b_00000000-11_test.cudafe2.c" --stub_file_name "/tmp/tmpxft_0000799b_00000000-11_test.cudafe2.stub.c" --gen_device_file_name "/tmp/tmpxft_0000799b_00000000-11_test.cudafe2.gpu" --nv_arch "compute_20" --module_id_file_name "/tmp/tmpxft_0000799b_00000000-3_test.module_id" --include_file_name "tmpxft_0000799b_00000000-2_test.fatbin.c" "/tmp/tmpxft_0000799b_00000000-10_test.cpp2.i" 
    ./test.cu(1): error: identifier "__T20" is undefined in device code

    1 error detected in the compilation of "/tmp/tmpxft_0000799b_00000000-10_test.cpp2.i".
    # --error 0x2 --

During compilation cuda assign array (const unsigned int){90, 50, 100} to __T20 variable and declare it as static. Thus its unaccessible from the main file. How to initialize global pointer with array in cuda?

#2

How about like this:

__constant__ unsigned int ff[3];

int main(){

  unsigned int h_ff[] = {90, 50, 100};
  cudaMemcpyToSymbol(ff, h_ff, 3*sizeof(unsigned int));

}
#3

I note that it seems like you are satisfied with the answer given on your cross posting:

https://stackoverflow.com/questions/47945178/initializing-cuda-global-variable

I think it’s worth pointing out that putting a pointer in __constant__ memory as demonstrated there does not imply that items referenced using that pointer will obey uniform access rules for __constant__ memory or be cached in the associated constant cache. It strikes me as a questionable construct, even apart from the fix supplied in the answer.

#4

thank you, txbob