Duplicate symbol issue with C headers

Question

This is my first time tackling a CUDA project that's slightly more complex than the simple write-single-source-file-and-compile routine. As expected, I'm facing some issues with C headers, namely duplicated symbols.

According to the linker, conflicts arise over the inclusion of the following header file in multiple .cu files:

env_vars.h

#ifndef ENV_VARS_H_
#define ENV_VARS_H_

/*** GLOBAL VARIABLES ***/
unsigned int h_n_osc;
__device__ unsigned int d_n_osc;

/*** CONSTANTS ***/
const double OMEGA_0 = 6.447421494058077e+09;

/* other constants defined in the middle */

#endif

multigpu.cu

#include "env_vars.h"
/* assigns h_n_osc */

adm_matrix.cu

#include "env_vars.h"
/* uses h_n_osc */

Building the project in Nsight Eclipse Edition results in the linker complaining about the h_n_osc variable being defined twice:

duplicate symbol _h_n_osc in:
    ./adm_matrix.o
    ./multigpu.o
ld: 1 duplicate symbol for architecture x86_64

Searching through the Internet, I've realized that moving the declaration of the h_n_osc variable to multigpu.cu and re-declaring it as an extern variable in adm_matrix.cu (and wherever I might need it later) solves the problem, which in fact it does.

Problem solved, but I'd like to take a deeper look into this:

1. Why doesn't the linker complain about the d_n_osc variable as well? And why are the constants (such as OMEGA_0) equally not a problem?
2. Does this mean that it is not possible to place global variables in header files?
3. What puzzles me most is that a number of sources over the Internet state that duplicated symbol errors should happen only when the header file contains a definition of the variable, while its simple declaration shouldn't constitute a problem. The reason I have a hard time believing this is that I'm facing the issue even though my header only contains a declaration! Am I missing something?

Thanks in advance for your patience, folks!

Answer 1

Header files should normally contain only declarative code. h_n_osc should be declared here, not defined.

extern unsigned int h_n_osc;

In at least one of your modules, or a new one of its own you will need a definition; for example:

env_vars.cu

#include "env_vars.h"
unsigned int h_n_osc;

Then link that. Alternatively you could of course place the definition in one of the existing modules multigpu.cu or adm_matrix.cu.

I am not sure of the semantics of the The CUDA __device__ extension, while it may link, it is not necessarily correct; you may end up with each module referencing a separate copy of the device variable; it may be necessary to qualify that with extern as well. This question appears to deal with that issue.