constant memory variable with "__constant__" don't support "extern" keyword for CUDA?

Recently, I found that if I define a constant memory variable with “constant” in other *.cu file and then I use cudaMemcpyToSymbol to assign value for constant memory variable, the value of constant memeory varibale are always zero!!! I don’t know why. my os is ubuntu16.04. cuda is 9.0. Could you someone tell me the reason?
Best regards for you!
here is my detailed information:

foo.h is as follow:
struct align(16) Test {
int a ;
int b ;
float c;
}; is as follow:
#include <device_functions.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include “foo.h”
constant Test Test_3; is as follow:
#include <device_functions.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include “foo.h”
#define CHECK(res) { if(res != cudaSuccess){printf("Error :%s:%d , ", FILE,LINE);
printf(“code : %d , reason : %s \n”, res,cudaGetErrorString(res));exit(-1);}}
constant Test Test_1; // definition in
extern constant Test Test_3; //definition in
global void foo()


printf("a1 =%d\nb1 =%d\nc1 = %f\n",Test_1.a,Test_1.b,Test_1.c);
printf("a3 =%d\nb3 =%d\nc3 = %f\n",Test_3.a,Test_3.b,Test_3.c);


void useCUDA()
Test Test2;
Test2.a =70;
Test2.b =90;
Test2.c =900.9;

cudaMemcpyToSymbol(Test_1,&Test2, sizeof(Test2));
cudaMemcpyToSymbol(Test_3,&Test2, sizeof(Test2));


int main(){



CMakeLists.txt is as follows:
cmake_minimum_required(VERSION 2.8)
set(CUDA_NVCC_FLAGS -O3;-G;-g;-std=c++11)
file(GLOB_RECURSE CURRENT_HEADERS *.h *.hpp .cuh containers/.hpp)
file(GLOB CURRENT_SOURCES *.cpp .cu containers/.cpp )
source_group(“Include” FILES {CURRENT_HEADERS}) source_group("Source" FILES {CURRENT_SOURCES})

because this requires that you compile with relocatable device code and device linking

For CMake, you’ll need something like this:

set_target_properties( main

as discussed here:

Thank you txbob

I have read the materials, which are suggested by you. I use the command “nvcc -arch=sm_61 -rdc=true -o main” in the terminal in ubuntu16.04. I succeed in compiling my example above. I have read the material suggested by you in “” . But
cmake 3.8 is required. I want compile the example in cmake 3.5 for some reason.

the CmakeLists.txt about the above example is as follow:

cmake_minimum_required(VERSION 3.0)


find_package(CUDA REQUIRED)

set(CUDA_NVCC_FLAGS -arch=sm_61;-G;-g;-rdc=true)

file(GLOB_RECURSE source_file *.cu *.cuh)

cuda_add_executable(Test {source_file}) set_target_properties( Test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) message("{CUDA_LIBRARIES}")
target_link_libraries(Test cuda
I test cmake 3.5, I failed.

the error messages are as follows:
Scanning dependencies of target Test
[100%] Linking CXX executable Test
CMakeFiles/Test.dir/ In function __sti____cudaRegisterAll()': /tmp/tmpxft_00003ff7_00000000-5_main.cudafe1.stub.c:2: undefined reference to __cudaRegisterLinkedBinary_39_tmpxft_00003ff7_00000000_6_main_cpp1_ii_Test_1’
CMakeFiles/Test.dir/ In function __sti____cudaRegisterAll()': /tmp/tmpxft_00003fcd_00000000-5_foo.cudafe1.stub.c:11: undefined reference to __cudaRegisterLinkedBinary_38_tmpxft_00003fcd_00000000_6_foo_cpp1_ii_Test_3’
collect2: error: ld returned 1 exit status
CMakeFiles/Test.dir/build.make:90: recipe for target ‘Test’ failed
make[2]: *** [Test] Error 1
CMakeFiles/Makefile2:67: recipe for target ‘CMakeFiles/Test.dir/all’ failed
make[1]: *** [CMakeFiles/Test.dir/all] Error 2
Makefile:83: recipe for target ‘all’ failed
make: *** [all] Error 2

To my surprise, If i make the second time. It succeed and got right results!!! I don’t know Why? Could you give me some suggestions about my CMakeLists.txt. I think maybe CMakeLists is something wrong!!!

Best Regards for you!