cudaMalloc3D (mistake in reference manual?)

[Blake]

I’ve copied and pasted the following out of the CUDA reference manual v3.0:

The following code sample allocates a width×height×depth 3D array of floating-point values and shows how to loop over the array elements in device code:

// Host code 

cudaPitchedPtr devPitchedPtr; 

cudaExtent extent = make_cudaExtent(64, 64, 64); 

cudaMalloc3D(&devPitchedPtr, extent);

MyKernel<<<100, 512>>>(devPitchedPtr, extent); 

// Device code __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) 

{

  char* devPtr = devPitchedPtr.ptr;

  size_t pitch = devPitchedPtr.pitch;

  size_t slicePitch = pitch * extent.height;

  for (int z = 0; z < extent.depth; ++z)

  {

	char* slice = devPtr + z * slicePitch;

	for (int y = 0; y < extent.height; ++y)

	{

	  float* row = (float*)(slice + y * pitch);

	  for (int x = 0; x < extent.width; ++x)

	  {

		float element = row[x];

	  }

	}

  }

}

However, this is from the reference manual:

4.8.2.10 cudaError_t cudaMalloc3D (struct cudaPitchedPtr  pitchedDevPtr, struct cudaExtent extent)

Allocates at least width * height * depth bytes of linear memory on the device and returns a cudaPitchedPtr in

which ptr is a pointer to the allocated memory.

Thus either the reference manual or the programming guide is wrong. I would guess that the mistake is in the programming guide, (i.e. extent is the array size in bytes, not floats). If that’s the case then this would fix the mistake:

char * row = slice + y * pitch;

for (int x = 0; x < extent.width; x+=sizeof(float))

{

  float element = *(float *)(row + x);

}

Is this correct or am I missing something?

[YDD]

I’ve found that it’s the width field of the extent which needs to be in bytes:

cudaExtent extent = make_cudaExtent( 64*sizeof(T), 64, 64);

This has come up a few times on the forums (I’ve been one of them!), but I don’t think it’s been fixed yet in the manual. On the GPU, you then access elements via

__device__ T operator()( const unsigned int ix,

				 				 const unsigned int iy,

							 const unsigned int iz ) const {

  const char* data = reinterpret_cast<const char*>(this->pitchedPtr);

  // Rows are pitch apart

  const size_t pitch = this->dataPitch;

  // Slices are slicePitch apart

  const size_t slicePitch = pitch * this->dims.y;

	

  const char* slice = data + ( iz * slicePitch );

  const char* row = slice + ( iy * pitch );

	

  return( reinterpret_cast<const T*>(row)[ix] );

}

Where the [font=“Courier New”]dataPitch[/font] member is copied from the [font=“Courier New”]pitch[/font] field of the [font=“Courier New”]cudaPitchedPtr[/font] struct.