学步园

CUDA中三维数组的创建，读取，拷贝。以下创建float4类型的width*height*depth大小的数组，在线性存储器中。

Host Code:

...
cudaPitchedPtr data;
cudaExtent extent = make_cudaExtent(width* sizeof(float4), height, depth);
cudaMalloc3D(&(d_data), extent);
...

Kernel Code:

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

{...
float4 element = *((float4*) (data + (x*sizeof(float4) + y*pitch + z*pitch*height)));
...

}

unsigned int size = width*height*depth*sizeof(float4);

float4 * h_volume=(float4*)malloc(size);

cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = data;
copyParams.dstPtr = make_cudaPitchedPtr(h_volume,width*sizeof(float4),width,height);
copyParams.extent = extent;
copyParams.kind = cudaMemcpyDeviceToHost;
cutilSafeCall(cudaMemcpy3D(&copyParams));