extern __device__ unsigned char deviceArray[] impossible?

Accelerated Computing CUDA CUDA Programming and Performance
#1

Hello Community,

currently I am working on a project, where I have to set bits on a device array to track the code coverage of the cude code.

I have an array of a given size, which I need in device and on host.

On host, it works fine.

I declare in the normal cu files:

extern unsigned char hostArray;

and in another cu file I define it:
unsiged char hostArray[7];

The same thing I want to do is for the device.
I tried
extern device unsigned char deviceArray;

but then my system tells me: size not known.

if I declare
extern device unsigned char deviceArray[7];

it seems to let the cudaMemCpyFromSymbol fail, if I want to retrieve the data.
If I let it copy from the defining .cu-file, The array is just empty, so it seems it is only touched by the defining cuda.

I suspect both .cu-Files get their own device unsigned charArray;

TL;DR;
How can I use an array allocated on device across multiple .cu-files (just like a global c-array)?

Best regards,
Sebastian.

#2

This sounds like a really sketchy way of implementing a testing framework.

#3

are you using separate compilation and linking?

#4

I did use directly a VS2013-Project and CUDA 7.5… if the seperate compile and link ain’t the default, then know. Is there a switch to change that?

BR,
Londran.

#5

Yes, there is a switch to change it. What you’re describing won’t work unless you use separate compilation and linking.

This describes something similar:

http://stackoverflow.com/questions/13597707/how-to-use-extern-cuda-device-variables

#6

Hi,

had time to try it and it works. Thank you very much. Also works now with Unified Memory (managed)

If someone runs into similiar Problems and uses VS2013 and CUDA 7.5:

Change Project-config:
CUDA C/C+±>Device->Code Generation: Switch of inheritance and crank it up to: compute_30,sm_30 (I think this is only needed because auf managed but who knows?)
CUDA C/C+±>Common->Generate Relocatable Device Code: Yes (-rdc=true)

For the ones using managed: only compile for x64.

For me, this is solved and can be closed.