Pointer arithmetic in device memory

Accelerated Computing CUDA CUDA Programming and Performance
1

I was wondering if anyone had experience using pointer arithmetic with device memory.

What I’m trying to do, is something like this:

extern float* array1;

extern int array1Size;

extern float* array2;

extern int array2Size;

__host__ void someFn( void )

{

  float* d_array;

 cudaMalloc( (void**)&d_array, sizeof(float) * (array1Size + array2Size) );

  cudaMemcpy( d_array, array1, sizeof(float) * array1Size, cudaMemcpyHostToDevice );

  cudaMemcpy( d_array + (sizeof(float) * array1Size), array2, sizeof(float) * array2Size, cudaMemcpyHostToDevice );

 someKernel<<<gridSize, blockSize>>>( d_array );

}

I’ve had this code in my kernel for a while, and everything seemed to be working, but all of the sudden I get this error on the second cudaMemcpy call in debug emulation mode:

Cuda error in file: invalid device pointer.

Am I supposed to be able to do something like this? Compiling in release seems to be giving me good results, but it could be a silent failure.

edit: fixed wrong code

2

edit: this post is incorrect

Not to reply to my own message, but I wrote a minimal sample app that basically did the pseudo code above, and got the same error (using CUDA_SAFE_CALL() from the nvidia SDK)

So unless this is a mistaken error, I’m assuming that you can’t do it.

If someone from NVidia or with better hardware knowledge would like to comment it might be helpful to have clarified.

3

I’m sure device pointer arithmetic is valid in device functions. So i guess it is also valid within host code.

I think the mistake is, that your d_array is of type float*. If so, you should remove “sizeof(float)” from the product “(sizeof(float) * array1Size)” when adding to the device pointer.

4

Yeah, thank you for pointing this out! That was the problem a very simple mistake - too many years of working with void* and char* got me used to throwing the sizeof multiple in thered :)

To clarify my previous posts, you can in fact do this kind of pointer arithmetic. To avoid the mistake I made, just make sure you check the type of the pointer you’re manipulating and offset (or not) appropriately.