global variables not working?

Accelerated Computing CUDA CUDA Programming and Performance
1

I am experiencing a strange problems which seems really stupid as I am sure it has worked before. CUDA 2.3. See code below:

#include <cuda_runtime.h>

#pragma comment(lib, "cudart.lib")

__constant__ int* globalVar;

__global__ void runme(int *a) {

	a[0] = globalVar[0];

	a[1] = globalVar[1];

}

int main() {

	cudaError err;

	int a[2];

	int *test;

	err = cudaMalloc((void**)&test, sizeof(int) * 2);

	err = cudaMalloc((void**)&globalVar, sizeof(int) * 2);

	runme<<<1, 1>>>(test);

	err = cudaMemcpy(&a, globalVar, sizeof(int) * 2, cudaMemcpyDeviceToHost);

	return 0;

}

This code compiles and links fine, but fails to run, exiting with a kernel error. It seems I cannot access the ‘globalVar’ from within the kernel??

Changing the code to the following runs just fine:

__global__ void runme(int *a, const int *b) {

	a[0] = b[0];

	a[1] = b[1];

}

...

runme<<<1, 1>>>(test, globalVar);

So how does this work? Has CUDA 2.3 changed something?

2

You can’t dynamically allocate global variables in CUDA. I also very much doubt you can passed device pointers at runtime to global memory via the cudaMemcpyToSymbol() mechanism either, although I haven’t tried it.