Multiple definitions of all __device__ __constant__ variables

[ewanbarr]

Throughout the astro-accelerate code, there are many uses of CUDA constant memory. For example:

/Users/ebarr/Soft/TDT/astro-accelerate/lib/headers/device_dedispersion_kernel.h:
   12  // Stores temporary shift values
   13  __device__ __constant__ float dm_shifts[15500];
   14: __device__ __constant__ int   i_nsamp, i_nchans, i_t_processed_s, i_t_processed_c;
   15  //__device__ __shared__ float fa_line[ARRAYSIZE+1];
   16  __device__ __shared__ float2 fa_line[ARRAYSIZE+1];

/Users/ebarr/Soft/TDT/astro-accelerate/lib/smem_works/device_dedispersion_kernel.cu:
    6  // Stores temporary shift values
    7  __device__ __constant__ float dm_shifts[15500];
    8: __device__ __constant__ int   i_nsamp, i_nchans, i_t_processed_s;
    9  __device__ __shared__ float2 fa_line[ARRAYSIZE+1];
   10  __device__ __shared__ float2 fb_line[ARRAYSIZE+1];

As can be seen above, there are global constant memory blocks that are defined with the same name in many different files. When compiling against the libastro-accelerate.a static library (such as with Cheetah) this results in many linking errors, like below:

../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x0): multiple definition of `c_sqrt_taps'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x20): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0xa0): multiple definition of `fd_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xc0): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0xfc0): multiple definition of `fc_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xfe0): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x1ee0): multiple definition of `fb_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x1f00): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x2e00): multiple definition of `fa_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x2e20): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x3d08): multiple definition of `i_t_processed_c'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d28): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x3d0c): multiple definition of `i_t_processed_s'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d2c): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x3d10): multiple definition of `i_nchans'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d30): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x3d14): multiple definition of `i_nsamp'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d34): first defined here
../../libcheetah.a(Dedispersion.cpp.o):(.bss+0x3d20): multiple definition of `dm_shifts'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d40): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x0): multiple definition of `c_sqrt_taps'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x20): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0xa0): multiple definition of `fd_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xc0): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0xfc0): multiple definition of `fc_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xfe0): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x1ee0): multiple definition of `fb_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x1f00): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x2e00): multiple definition of `fa_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x2e20): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x3d08): multiple definition of `i_t_processed_c'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d28): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x3d0c): multiple definition of `i_t_processed_s'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d2c): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x3d10): multiple definition of `i_nchans'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d30): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x3d14): multiple definition of `i_nsamp'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d34): first defined here
../../libcheetah.a(PipelineHandlerFactory.cpp.o):(.bss+0x3d20): multiple definition of `dm_shifts'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d40): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x0): multiple definition of `c_sqrt_taps'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x20): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0xa0): multiple definition of `fd_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xc0): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0xfc0): multiple definition of `fc_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0xfe0): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x1ee0): multiple definition of `fb_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x1f00): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x2e00): multiple definition of `fa_line'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x2e20): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x3d08): multiple definition of `i_t_processed_c'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d28): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x3d0c): multiple definition of `i_t_processed_s'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d2c): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x3d10): multiple definition of `i_nchans'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d30): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x3d14): multiple definition of `i_nsamp'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d34): first defined here
../../libcheetah.a(Sps.cpp.o):(.bss+0x3d20): multiple definition of `dm_shifts'
CMakeFiles/gtest_pipeline.dir/src/DedispersionTest.cpp.o:(.bss+0x3d40): first defined here
collect2: error: ld returned 1 exit status

These errors make it impossible to use libastro-accelerate.a To solve this problem these definitions need to either be moved into a common globals.h (for example) header with appropriate guards, or they need to be hidden in namespaces. Moving these common definitions to a global header would make most sense in this context.

[ewanbarr]

Looks like Wes's additions to the master branch have solved this. I have merged these changes into aa_interface and committed.