LNK2005 Error when using Constant Memory in CUDA Header (.cuh) File

Question

I have a CUDA Header (.cuh) file that contains two constant float arrays. There are two files that include this header, one is a CPP file that tries to copy to this constant memory and one is a CUDA file that tries to use this constant memory. All three files are in a project that should compile to make a DLL.

I've tried to simplify things with the following code:

obj1.cuh

#pragma once

__constant__ float d_array1[5];
__constant__ float d_array2[5];

obj1.cu

#include "obj1.cuh"

//do random stuff
__global__ void kernel(float * d_array1, float * d_array2) {
  int id = threadIdx.x;
  float sum = d_array1[i] + d_array2[i];
}

ext.cpp

#include "obj1.cuh"

void function(float * array1, float * array2) {
  cudaMemcpyToSymbol(d_array1, array1, sizeof(float)*5);
  cudaMemcpyToSymbol(d_array2, array2, sizeof(float)*5);

  kernel<<<1,5>>>(d_array1,d_array2);
}

The build fails and outputs the following errors:

1>ext.obj : error LNK2005: "float * d_array1" (?d_array1@@3PAMA) already defined in obj1.cu.obj
1>ext.obj : error LNK2005: "float * d_array2" (?d_array2@@3PAMA) already defined in obj1.cu.obj

Before you ask, yes I have tried using include guards instead of the pragma once and it still outputs the same error.

Include guard version of obj1.cuh

#ifndef CONSTANTARRAYS
#define CONSTANTARRAYS
__constant__ float d_array1[5];
__constant__ float d_array2[5];
#endif

Some programmer dude Some programmer dude · Accepted Answer · 2014-09-08T09:42:04

The problem is exactly what the linker tells you: The variables are defined in multiple translation units.

When the preprocessor includes a header files, it quite literally pastes the text in the header into the place where the #include directive was, which means both generated objects files will have the definitions of the variables.

What you should do is to only declare the variables in the header file, then in a single source file you define them. To change the definition to a declaration, the simplest way is to prepend the extern keyword:

#pragma once

extern __constant__ float d_array1[5];
extern __constant__ float d_array2[5];

Then in a single source file you have the old definitions:

__constant__ float d_array1[5];
__constant__ float d_array2[5];

The #pragma once or header include guards prevents a header file to be included multiple times in the same translation unit.

LNK2005 Error when using Constant Memory in CUDA Header (.cuh) File

1 Answers