我想有一个CUDA 3D float数组,这里是我的代码:
#define SIZE_X 128 //numbers in elements
#define SIZE_Y 128
#define SIZE_Z 128
typedef float VolumeType;
cudaExtent volumeSize = make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z); //The first argument should be SIZE_X*sizeof(VolumeType)??
float *d_volumeMem;
cutilSafeCall(cudaMalloc((void**)&d_volumeMem, SIZE_X*SIZE_Y*SIZE_Z*sizeof(float)));
.....//assign value to d_volumeMem in GPU
cudaArray *d_volumeArray = 0;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
cutilSafeCall( cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize) );
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = make_cudaPitchedPtr((void*)d_volumeMem, SIZE_X*sizeof(VolumeType), SIZE_X, SIZE_Y); //
copyParams.dstArray = d_volumeArray;
copyParams.extent = volumeSize;
copyParams.kin = cudaMemcpyDeviceToDevice;
cutilSafeCall( cudaMemcpy3D(©Params) );
其实,我的程序运行良好。 但我不知道结果是正确的。 这里是我的问题,在CUDA liberay,它表示make_cudaExtent的第一个参数是“以字节为单位宽度”和另外两个是在元素的高度和深度。 所以我觉得在我上面的代码,第五个行应该是
cudaExtent volumeSize = make_cudaExtent(SIZE_X*sizeof(VolumeType), SIZE_Y, SIZE_Z);
但这种方式,就在cutilSafeCall(cudaMemcpy3D(copyParams))错误“无效的参数”; 为什么?
和另一个难题是strcut cudaExtent,作为CUDA库所述,其组分宽度“指的是阵列中的存储器的情况下,以字节为单位指的线性存储器时宽度中的元素”代表。 所以我觉得在我的代码时,我指的元素volumeSize.width它应该是数量。 但是,如果我用
cudaExtent volumeSize = make_cudaExtent(SIZE_X*sizeof(VolumeType), SIZE_Y, SIZE_Z);
所述volumeSize.width将SIZE_X *的sizeof(VolumeType)(128 * 4),即在元件中的字节数数代替。
在许多CUDA SDK,他们使用碳作为VolumeType,所以他们只是用SIZE_X作为make_cudaExtent第一个参数。 但是,我的是浮动,所以,任何人都可以告诉我这是创建一个cudaExtent如果我需要用它来创建一个三维阵列以正确的方式? 非常感谢!
让我们来回顾一下在文档cudaMemcpy3D
说:
程度字段定义在元件转移区的尺寸。 如果CUDA阵列参与复制的程度是在该数组的元素来定义的。 如果没有CUDA阵列参与复制,则盘区在无符号字符的元素来定义。
同样地,对于文档cudaMalloc3DArray
指出:
所有值都在指定的元素
所以,你需要形成用于两个呼叫的程度需要具有在元件的第一维度(因为在该分配的一个cudaMemcpy3D
是阵列)。
但你可能在你的代码不同的问题,因为你是分配的线性存储器源d_volumeMem
使用cudaMalloc
。 cudaMemcpy3D
期望线性源存储器已被分配与兼容节距。 您的代码只是使用尺寸的线性分配
SIZE_X*SIZE_Y*SIZE_Z*sizeof(float)
现在,它可能是你所选择的尺寸产生了您正在使用的硬件兼容的间距,但不能保证它会这么做。 我建议使用cudaMalloc3D
来分配所述线性源存储器。 这在你小的代码段建立了一个扩大示范可能是这样的:
#include <cstdio>
typedef float VolumeType;
const size_t SIZE_X = 8;
const size_t SIZE_Y = 8;
const size_t SIZE_Z = 8;
const size_t width = sizeof(VolumeType) * SIZE_X;
texture<VolumeType, cudaTextureType3D, cudaReadModeElementType> tex;
__global__ void testKernel(VolumeType * output, int dimx, int dimy, int dimz)
{
int tidx = threadIdx.x + blockIdx.x * blockDim.x;
int tidy = threadIdx.y + blockIdx.y * blockDim.y;
int tidz = threadIdx.z + blockIdx.z * blockDim.z;
float x = float(tidx)+0.5f;
float y = float(tidy)+0.5f;
float z = float(tidz)+0.5f;
size_t oidx = tidx + tidy*dimx + tidz*dimx*dimy;
output[oidx] = tex3D(tex, x, y, z);
}
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
template<typename T>
void init(char * devPtr, size_t pitch, int width, int height, int depth)
{
size_t slicePitch = pitch * height;
int v = 0;
for (int z = 0; z < depth; ++z) {
char * slice = devPtr + z * slicePitch;
for (int y = 0; y < height; ++y) {
T * row = (T *)(slice + y * pitch);
for (int x = 0; x < width; ++x) {
row[x] = T(v++);
}
}
}
}
int main(void)
{
VolumeType *h_volumeMem, *d_output, *h_output;
cudaExtent volumeSizeBytes = make_cudaExtent(width, SIZE_Y, SIZE_Z);
cudaPitchedPtr d_volumeMem;
gpuErrchk(cudaMalloc3D(&d_volumeMem, volumeSizeBytes));
size_t size = d_volumeMem.pitch * SIZE_Y * SIZE_Z;
h_volumeMem = (VolumeType *)malloc(size);
init<VolumeType>((char *)h_volumeMem, d_volumeMem.pitch, SIZE_X, SIZE_Y, SIZE_Z);
gpuErrchk(cudaMemcpy(d_volumeMem.ptr, h_volumeMem, size, cudaMemcpyHostToDevice));
cudaArray * d_volumeArray;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<VolumeType>();
cudaExtent volumeSize = make_cudaExtent(SIZE_X, SIZE_Y, SIZE_Z);
gpuErrchk( cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize) );
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = d_volumeMem;
copyParams.dstArray = d_volumeArray;
copyParams.extent = volumeSize;
copyParams.kind = cudaMemcpyDeviceToDevice;
gpuErrchk( cudaMemcpy3D(©Params) );
tex.normalized = false;
tex.filterMode = cudaFilterModeLinear;
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.addressMode[2] = cudaAddressModeWrap;
gpuErrchk(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));
size_t osize = 64 * sizeof(VolumeType);
gpuErrchk(cudaMalloc((void**)&d_output, osize));
testKernel<<<1,dim3(4,4,4)>>>(d_output,4,4,4);
gpuErrchk(cudaPeekAtLastError());
h_output = (VolumeType *)malloc(osize);
gpuErrchk(cudaMemcpy(h_output, d_output, osize, cudaMemcpyDeviceToHost));
for(int i=0; i<64; i++)
fprintf(stdout, "%d %f\n", i, h_output[i]);
return 0;
}
您可以为自己确认纹理的输出读取匹配的主机上的原始来源记忆。
作为cudaArray涉及您的代码是正确的。 定到阵列的channelDesc保存有关浮体大小(4个字节)的信息。 您程度的规范。 有“*的sizeof(VolumeType)”将是两个存储器指针之间右复制(与srcPtr,dstPtr使用)。 此外srcPos和dstPos然后将不得不以字节为单位给出,即第一个参数“*的sizeof(VolumeType)”。
间距问题还可能与视GPU /驱动3D OPS出现。 我已经看到了这一点,但罕(2 ^ n维应该是罚款)。 你也可以把它分解使用cudaMemCpy2DToArray在一个for循环,因为它应该更宽容的间距。 没有cudaMalloc2D,所以对于2D操作的任何总是正确的球场是由SDK发出。