我們來回顧一下爲cudaMemcpy3D
的文件說:
程度字段定義 元素的轉移區域的尺寸。如果CUDA陣列參與複製的程度是在該數組的元素來定義 。如果沒有CUDA數組是 參與複製然後的區段在 無符號字符的元素來定義。
同樣地,對於cudaMalloc3DArray
註釋的文檔:
所有值均以元素指定
所以,你需要形成要呼叫的程度需要有在元素的第一個維度(因爲cudaMemcpy3D
中的一個分配是一個數組)。
但是,您的代碼中可能有其他問題,因爲您正在使用cudaMalloc
分配線性內存來源d_volumeMem
。 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;
}
您可以自己確認的紋理輸出的主機上讀取原始的源內存匹配。
我很好奇你怎麼寫「實際上,我的程序運行良好,但我不確定結果是否正確」。這不是矛盾嗎?你怎麼能不確定結果是「正確的」?你當然知道正確的結果應該是什麼? – talonmies
我的意思是程序可以運行並輸出結果,但我不確定結果是否正確。事實上,我需要一個3D紋理來繪製圖片,但我不知道圖片應該是什麼。關鍵是我將數據複製到我上面描述的3D數組中。如果不對,我認爲結果可能會有一些錯誤。 – TonyLic