2

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!

thequanticlad
  • 81
  • 2
  • 9

1 Answers1

9

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.

Community
  • 1
  • 1
Clifford
  • 88,407
  • 13
  • 85
  • 165
  • 2
    Correct, `__device__` (and `__constant__`) variables [have file/translation unit scope](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-declarations). – Robert Crovella Dec 02 '14 at 20:36
  • @RobertCrovella : Thanks; so they have the similar semantics to `static` and `const`? – Clifford Dec 02 '14 at 20:43
  • It seems to be a [complicated topic](http://stackoverflow.com/questions/3709207/c-semantics-of-static-const-vs-const). But I would say yes, similar. In the absence of other specifics, CUDA should behave similarly to a C++ compiler. `__device__` and `__constant__` variables are not referencable in ordinary host C++ code, but that is a separate topic. – Robert Crovella Dec 02 '14 at 20:58
  • Thanks to you too, guys! As I commented above, most of my doubts clearly stemmed from my misunderstanding of the concepts of declaration and definition of a variable. For the sake of completeness, [a little navigation](http://bit.ly/1rSDpH5) here on Stack Overflow also answered to question 2 in my original post: `const`ants have default internal linkage in C++ (whereas they default to external linkage in plain C), so my solution (i.e. defining them in the header) works fine just because CUDA C is actually CUDA C/C++; a pure C compiler would have rejected the constants, too. – thequanticlad Dec 02 '14 at 22:17
  • @thequanticlad : If you declare const's `static const` the semantics are then the same in C and C++. However if your code takes the address of the const, they will be different in each module, however in C++ at least if you don't take the address of a const, all references are replaced with literals. – Clifford Dec 02 '14 at 22:55