Take the 2-minute tour ×

I have a small piece of code which runs perfectly on Nvidia old architecture (Tesla T10 processor) but not on Fermi (Tesla M2090)

I learned that Fermi behaves slightly differently. Due to which unsafe code might work correctly on old architectures, while on Fermi it catches the bug.

But I don't know how to resolve it.

Here is my code:

__global__ void exec (int *arr_ptr, int size, int *result) {

    int tx = threadIdx.x;
    int ty = threadIdx.y;

    *result = arr_ptr[-2];


void run(int *arr_dev, int size, int *result) {

    cudaStream_t stream = 0;
    int *arr_ptr = arr_dev + 5;

    dim3 threads(1,1,1);
    dim3 grid (1,1);

    exec<<<grid, threads, 0, stream>>>(arr_ptr, size, result);


since I am accessing arr_ptr[-2], the fermi throws CUDA_EXCEPTION_10, Device Illegal Address. But it is not. The address is legal.

Can anyone help me on this.

My driver code is

int main(){
    int *arr;
    int *arr_dev = NULL;
    int result = 1;

    arr = (int*)malloc(10*sizeof(int));

            arr[i] = i;

    if(arr_dev == NULL)
            cudaMalloc((void**)&arr_dev, 10);
            cudaMemcpy(arr_dev, arr, 10*sizeof(int), cudaMemcpyHostToDevice);

    run(arr_dev, 10, &result);
    printf("%d \n", result);
    return 0;


share|improve this question
I don't see why it shouldn't work. Could you verify / show us the code, how arr_dev and result pointers are generated? –  CygnusX1 Jun 28 '11 at 20:20
The cudaMalloc call is using an invalid size. It should be 10*sizeof(int). –  talonmies Jun 28 '11 at 20:44
add comment

1 Answer

up vote 5 down vote accepted

Fermi cards have much better memory protection on the device and will detect out of bounds conditions which will appear to "work" on older cards. Use cuda-memchk (or the cuda-memchk mode in cuda-gdb) to get a better handle on what is going wrong.


This is the culprit:


which should be

cudaMalloc((void**)&arr_dev, 10*sizeof(int));

This will result in this code

int *arr_ptr = arr_dev + 5;

passing a pointer to the device which is out of bounds.

share|improve this answer
Can you explain on "Fermi cards have much better memory protection on the device" –  veda Jun 28 '11 at 20:56
The Fermi memory management unit is a more sophisticated design which can detect a lot of protection fault conditions which older hardware cannot. –  talonmies Jun 28 '11 at 20:59
Thanks, The above problem is solved. But now, I tried the mem-check and found Invalid global read of size 4 == at 0x00000bf8 in == by thread (6,3,0) in block (0,0) == Address 0x6004b9ef8 is out of bounds ----Does this mean that I caused an error. Because the same code works on older hardware. So now I am not sure whether the error is due to me or Fermi –  veda Jun 28 '11 at 21:04
You have some out of bounds memory global memory access in your code. Just because the code appears to "work" on older hardware doesn't mean the problem isn't occurring on both, only that the older hardware isn't detecting the problem. –  talonmies Jun 28 '11 at 21:18
oh ok... So I am making some invalid global memory access. The mem-check is providing me the line in which the error happens and when I tried to calculate the address manually, I am not able to figure it out. Do you know some tools other than cuda-gdb to find the error. Do you suggest any ideas?? –  veda Jun 28 '11 at 21:29
show 1 more comment

Your Answer


