header image
Finally, JET is back
August 18th, 2009 under Music. [ Comments: none ]

The new album “Shaka Rock” is supposed to be available on the 21st of August. 😀
CUDA 2D/2D Memcpy Issue and Pitched Pointers
August 18th, 2009 under OpenGL. [ Comments: none ]

In my application I wanted to copy 3D array data from linear memory (allocated with

Click to toggle codeblock
1

cudaMalloc
cudaMalloc
) to a
Click to toggle codeblock
1

cudaArray
cudaArray
with
Click to toggle codeblock
1

cudaMemcpy3D
cudaMemcpy3D
. Now, I had the problem that
Click to toggle codeblock
1

cudaMemcpy3D 
cudaMemcpy3D
reported
Click to toggle codeblock
1

cudaErrorInvalidValue
cudaErrorInvalidValue
. If you encounter a problem like that it might be because of the fact that
Click to toggle codeblock
1

cudaMalloc 
cudaMalloc
was used instead of
Click to toggle codeblock
1

cudaMalloc3D
cudaMalloc3D
for allocatimg the linear memory. If
Click to toggle codeblock
1

cudaMalloc 
cudaMalloc
was used then you are forced to create you own
Click to toggle codeblock
1

cudaPitchedPtr 
cudaPitchedPtr
with
Click to toggle codeblock
1

make_cudaPitchedPtr
make_cudaPitchedPtr
in order to call
Click to toggle codeblock
1

cudaMemcpy3D
cudaMemcpy3D
. According to the manual (this information is not extremely emphasized in it)
Click to toggle codeblock
1

make_cudaPitchedPtr
make_cudaPitchedPtr
is not guaranteed to construct a pitched pointer which is valid for memory copy procedures. One should always use
Click to toggle codeblock
1

cudaMalloc3D 
cudaMalloc3D
or
Click to toggle codeblock
1

cudaMallocPitched 
cudaMallocPitched
whenever we want to memcpy because these methods guarantee valid pitched pointers.

“For allocations of 2D and 3D objects, it is highly recommended that programmers perform allocations using cudaMalloc3D() or cudaMallocPitch(). Due to alignment restrictions in the hardware, this is especially true if the application will be performing memory copies involving 2D or 3D objects (whether linear memory or CUDA arrays).” [from the Reference Manual]

This, in fact, solved my problem of the

Click to toggle codeblock
1

cudaErrorInvalidValue
cudaErrorInvalidValue
when I tried to
Click to toggle codeblock
1

cudaMemcpy3D
cudaMemcpy3D
. However, the fact that we use pitched pointer complicates the array access in the kernels. First of all we have to keep track of the data type stored in a pitched pointer (ptr is a void pointer). Furthermore, we have to be careful when dealing with pointer arithmectics because the pitch is, obviously, a byte offset (width of the array).  Here is some sample code that shows how to access pitched pointer data. We assume we want to store
Click to toggle codeblock
1

float4 
float4
data in our 3D array which is of size (width, height, depth). In the kernel we want to access the element (x,y,z).

Host Code:

Click to toggle codeblock
1
2
3
4
5

...<br />
cudaPitchedPtr data;<br />
extent = make_cudaExtent(width* sizeof(float4), height, depth);<br />
cudaMalloc3D(&(d_data), extent);<br />
...
...<br />
cudaPitchedPtr data;<br />
extent = make_cudaExtent(width* sizeof(float4), height, depth);<br />
cudaMalloc3D(&(d_data), extent);<br />
...

Kernel Code:

Click to toggle codeblock
1

__global__ kernel(char* data, size_t pitch, ...) {
__global__ kernel(char* data, size_t pitch, ...) {

Click to toggle codeblock
1
2
3

...<br />
float4 element = *((float4*) (data + (x*sizeof(float4) + y*pitch + z*pitch*height)));<br />
...
...<br />
float4 element = *((float4*) (data + (x*sizeof(float4) + y*pitch + z*pitch*height)));<br />
...

Click to toggle codeblock
1

}
}

The one thing I’d like to point out is the fact that we have to use

Click to toggle codeblock
1

x*sizeof(float4)
x*sizeof(float4)
for accessing the data. This makes sense because
Click to toggle codeblock
1

data 
data
is a char/byte pointer at the time of the pointer arithmetic. Thus, we have to move the pointer by
Click to toggle codeblock
1

x*sizeof(float4)
x*sizeof(float4)
. This musn’t be done in the y and z direction because the
Click to toggle codeblock
1

sizeof(float4)
sizeof(float4)
is already considered in the pitch (the pitch is the width of the array in byte, probably padded for faster memory access).
CUDA 3D Array Pitch
August 17th, 2009 under OpenGL. [ Comments: none ]

The CUDA manual isn’t really precise in telling how to use the pitch with 3D Array (memory obtained by using

Click to toggle codeblock
1

cudaMallocPitch, cudaMalloc3D
cudaMallocPitch, cudaMalloc3D
, not 3D cudaArrays). So here is a brief summary of how (I think) everything works.

What nVidia means by a pitch is basically the width in byte of an array (2D or 3D). We can specify the width of the array in bytes like this

Click to toggle codeblock
1

volumeSize.width *sizeof(YOUR_TYPE)
volumeSize.width *sizeof(YOUR_TYPE)
. With this information we can access any element in a 3D array. We can compute the
Click to toggle codeblock
1

slicePitch
slicePitch
(for a 3D array) which is
Click to toggle codeblock
1

volumeSize.height*pitch
volumeSize.height*pitch
. The
Click to toggle codeblock
1

slicePitch
slicePitch
is the offset in bytes from one slice in a 3D array to the next slice.

Now let’s say we want to access the element (x,y,z) in a 3D array

Click to toggle codeblock
1

data
data
(which is the base pointer). This is how we can do it:

Click to toggle codeblock
1

(myType) (((char*)data)[x*sizeof(myType) + y*pitch + z*slicePitch])
(myType) (((char*)data)[x*sizeof(myType) + y*pitch + z*slicePitch])

So what we get from the methods

Click to toggle codeblock
1

cudaMallocPitch, cudaMalloc3D
cudaMallocPitch, cudaMalloc3D
is linear memory which need to be accessed in the way described above.
We need to know how pitches work when we want to copy arrays (when the copy method expects a pitched pointer which we can create with
Click to toggle codeblock
1

make_cudaPitchedPtr(baseAddress, pitch, width, height)
make_cudaPitchedPtr(baseAddress, pitch, width, height)
). This is also sufficient for the copy process. Let’s assume we use
Click to toggle codeblock
1

cudaMemcpy3D
cudaMemcpy3D
. In this case we must specify the extent, in other words the size of our array. All necessary information to execute the copy are now present (it can be ensured that no invalid memory will be accessed) . Each element can be accessed by
Click to toggle codeblock
1

cudaMemcpy3D
cudaMemcpy3D
as described above. So, this means that we do not need to use the depth (in case of a 3D array) when we
Click to toggle codeblock
1

make_cudaPitchedPtr
make_cudaPitchedPtr
.
Even more Classics :D
August 5th, 2009 under This And That. [ Comments: none ]

Wi buchet mir dr glutz :)
August 3rd, 2009 under This And That. [ Comments: none ]

Some classic swiss comedy.