CUDA device pointer manipulation

Question

I've used:

float *devptr;
//...
cudaMalloc(&devptr, sizeofarray);
cudaMemcpy(devptr, hostptr, sizeofarray, cudaMemcpyHostToDevice);

in CUDA C to allocate and populate an array. Now I'm trying to run a cuda kernel, e.g.:

__global__ void kernelname(float *ptr)
{
   //...
}

in that array but with an offset value. In C/C++ it would be someting like this:

kernelname<<<dimGrid, dimBlock>>>(devptr+offset);

However, this doesn't seem to work.

Is there a way to do this without sending the offset value to the kernel in a separate argument and use that offset in the kernel code? Any ideas on how to do this?

talonmies talonmies · Accepted Answer · 2011-05-06T09:46:44

Pointer arithmetic does work just fine in CUDA. You can add an offset to a CUDA pointer in host code and it will work correctly (remembering the offset isn't a byte offset, it is a plain word or element offset).

EDIT: A simple working example:

#include <cstdio>
int main(void)
{

    const int na = 5, nb = 4;
    float a[na] = { 1.2, 3.4, 5.6, 7.8, 9.0 };
    float *_a, b[nb];

    size_t sza = size_t(na) * sizeof(float);
    size_t szb = size_t(nb) * sizeof(float);

    cudaFree(0);

    cudaMalloc((void **)&_a, sza );
    cudaMemcpy( _a, a, sza, cudaMemcpyHostToDevice);
    cudaMemcpy( b, _a+1, szb, cudaMemcpyDeviceToHost);

    for(int i=0; i<nb; i++)
        printf("%d %f\n", i, b[i]);

    cudaThreadExit();
}

Here, you can see a word/element offset has been applied to the device pointer in the second cudaMemcpy call to start the copy from the second word, not the first.

CUDA device pointer manipulation

2 Answers