CUDA_EXCEPTION_14:Warp Illegal Address - when accessing device array

Hi all,

I have a situation where I have a kind of “hash table” with a simple hashing function that is guaranteed not to give collisions. The hash table is just an array that I have declared on the device. However, when trying to access it, I get this error:

transformEventPackage [0] [device 0 (GP104-A)]  (Signal) - CUDA_EXCEPTION_14:Warp Illegal Address ; ExceptionPC=0x16e4130 CUDA Thread (0,0,0) Block (0,48,0) All Kernel Threads (256 Blocks of 96 Threads)

I have set up a trivial example that does (roughly speaking) the same things, and in fact gives me the same error. Is there something wrong with declaring arrays on the device in this fashion? :

#include <cuda.h>
#include <cuda_runtime.h>

#define SIZE 32

__device__ int testarray[512000];

__global__
void testFillArray(int * values) {
	int idx = threadIdx.x+blockIdx.x*blockDim.x;
	int gridIdx = blockIdx.y;
	testarray[gridIdx][idx] = idx;
	if(testarray[gridIdx][idx]==100) {
		values[gridIdx] = testarray[gridIdx][idx];
	}
}

int main(void)
{
	printf("begin... \n");
	int * values;
	cudaMallocManaged(&values, sizeof(int)*SIZE);
	dim3 gridDim(SIZE,1000);
	testFillArray<<<gridDim,512>>>(values);
	cudaDeviceSynchronize();
	printf("end...\n");
	return 0;
}

since blockIdx.y ranges from 0…999:

dim3 gridDim(SIZE,1000);
                  ^^^^

and you are using it as the first index into testarray:

int gridIdx = blockIdx.y;
testarray[gridIdx][idx] = idx;
          ^^^^^^^

and you have declared the first index of testarray to range from 0…31:

#define SIZE 32

__device__ int testarray[512000];
                         ^^^^

I’m not surprised you are having trouble. An out-of-rage/illegal address is precisely what I would expect.

Oh wow. facepalm Thank you!