extern global device variable previously defined in cuda c++

Question

I have three files: a.cu, b.cu, and c.h. I want to have debug variables accessible to all compilation units, so I declared them as extern in c.h, and defined only in a.cu, per the answer here Global Variable within Multiple Files:

a.cu

#include "c.h"
#include <stdio.h>
#include <iostream>

bool h_debug;                                                                       //works fine
//__device__ bool d_debug;                     //POINT A

int main(int argc, char* argv[])
{
    h_debug = (argc > 1 ? true : false);
    std::cout<<"Host: "<<(h_debug ? "true" : "false")<<std::endl;
    cudaMemcpyToSymbol(d_debug, &h_debug, sizeof(bool));  
    std::cout<<cudaGetErrorString(cudaGetLastError())<<std::endl; 
    cudaDeviceSynchronize();
}

b.cu

#include "c.h"
#include "stdio.h"

__global__ 
void myKernel(){
    if(d_debug){
        printf("device debug on\n");
    }
    else{
        printf("device debug off\n");
    }
}

c.h

#ifndef MAIN_H
#define MAIN_H

extern __device__ bool d_debug;
extern bool h_debug;

#endif  /* MAIN_H */

At the line marked POINT A, if I have it commented, the code compiles but I get a cuda error when running, as one might expect:

$ nvcc a.cu b.cu -o globalTest
$ globalTest
Host: true
invalid device symbol

If I un-comment the line to define the d_debug, I get a compiler error that doesn't make sense to me ...

$ nvcc a.cu b.cu -o globalTest
a.cu:1:32: warning: unknown option after ‘#pragma GCC diagnostic’ kind [-Wpragmas]
a.cu:6:13: error: redefinition of ‘bool d_debug’
c.h:4:13: error: ‘bool d_debug’ previously declared here

Why is it not working like the global host variable? How should I create a global device variable which is accessible by all compilation units?

you'll need to incorporate device linking into your compilation sequence. Otherwise a.cu and b.cu have *separate* copies of the `d_debug` device variable. There are many questions here on the `cuda` tag that discuss this. — Robert Crovella, Jul 07 '16 at 16:30
This [question](http://stackoverflow.com/questions/24617318/cuda-constant-memory-value-not-correct) and linked duplicates generally covers the same ground. `__device__` variables have the same kind of default *compilation-unit scope* as `__constant__` variables, and the solution approach is the same. — Robert Crovella, Jul 07 '16 at 16:32
That answers it perfectly. It's not really a duplicate question but it is a duplicate answer ... should I close the question? — Daniel B., Jul 07 '16 at 16:36
Your answer is fine IMO. If someone wants to mark this a duplicate they can. In the meantime, you can come back at some point and accept your answer. I have upvoted it anyway. — Robert Crovella, Jul 07 '16 at 16:51

score 1 · Answer 1 · edited May 23 '17 at 10:28

From comments by Robert Crovella and the answer here: CUDA constant memory value not correct, The problem is that a.cu and b.cu have separate copies of d_debug. To prevent this, separate compilation and linking must be used; to that end, create relocatable device code (rather than executable device code).

--relocatable-device-code {true|false}
-rdc Enable (disable) the generation of relocatable device code. If disabled, executable device code is generated. Relocatable device code must be linked before it can be executed.

Allowed values for this option: true, false.

Default value: false

Changing nvcc arguments to:

nvcc -rdc=true a.cu b.cu -o globalTest

Resolves the issue.

extern global device variable previously defined in cuda c++

1 Answers1