1

这是我的问题：

我的设备上有一个 3D 数组float3：

int size[3] = {416,464,512};
cudaExtent extent = make_cudaExtent(size[0]*sizeof(float3),size[1],size[2]);
cudaPitchedPtr renderedVolume;
int ret = cudaMalloc3D(&renderedVolume, extent);
size_t pitch = renderedVolume.pitch; //pitch = 5,120
size_t slicePitch = pitch * size[1]; //slicePitch = 2,375,680

然后我使用它，让它充满优秀的数据。

之后，我希望将其复制到主机上的一维线性内存中：

float *host_memory = (float*)malloc(size[0]*size[1]*size[2]*sizeof(float3));
cudaMemcpy3DParms p = {0};
p.srcPtr = renderedVolume;
p.dstPtr = make_cudaPitchedPtr(host_memory,size[0]*sizeof(float3),size[0],size[1]); 
p.extent = make_cudaExtent(size[0]*sizeof(float3),size[1],size[2]);
p.srcPos = make_cudaPos(0,0,0);
p.dstPos = make_cudaPos(0,0,0);
p.kind=cudaMemcpyDeviceToHost;
cudaMemcpy3D(&p);

我将 host_memory 中的结果与我最初写入 tu renderedVolume( my_data) 的数据以及我在 my 中读取的数据3Dmemory逐个切片进行比较：

float* test1 = (float*)malloc(size[0]*size[1]*sizeof(float3));
cudaMemcpy(test1, myData, size[0]*size[1]*sizeof(float3) , cudaMemcpyDeviceToHost);
float* test2 = (float*)malloc(size[0]*size[1]*sizeof(float3));
cudaMemcpy(test2,(char*)renderedVolume.ptr + slicePitch * i,size[0]*size[1]*sizeof(float3), cudaMemcpyDeviceToHost);

问题：

  • 第一个切片 ( i=0) 没问题，我在 host_memory 中有相同的数据，test1并且test2.
  • 在第二个切片中，我在test1和中有相同的数据test2。但是，我应该在host_memory+579072（=float每个切片的数量，也是heigth*pitch目标倾斜指针的数量）中找到这些数据，并且我在host_memory+577504. 它是按1568字节关闭的，这与我所知道的没有任何关系，这就是为什么如果你们中的任何人知道我的代码中可能存在什么问题，我将非常感激？
4

1 回答 1

5

这是一个较晚的答案，用于从未回答的列表中删除此问题。

下面，我提供了一个完整的代码，展示了如何分配3D内存cudaMalloc3D，将主机分配的1D内存移动到3D设备内存，通过函数对设备数据cudaMemcpy3D执行一些操作，并将结果数据移动回主机内存，再次通过。3Dtest_kernel_3D __global__3D1DcudaMemcpy3D

该__global__函数对设备内存test_kernel_3D的每个元素进行平方。3D特别是，网格的每个线程都负责沿着“深度”维度2D执行循环。for

#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<conio.h>

#define BLOCKSIZE_x 16
#define BLOCKSIZE_y 16

#define N 128
#define M 64
#define W 16

/*****************/
/* CUDA MEMCHECK */
/*****************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
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) { getch(); exit(code); }
    }
}

/*******************/
/* iDivUp FUNCTION */
/*******************/
int iDivUp(int a, int b){ return ((a % b) != 0) ? (a / b + 1) : (a / b); }

/******************/
/* TEST KERNEL 3D */
/******************/
__global__ void test_kernel_3D(cudaPitchedPtr devPitchedPtr)
{
    int tidx =  blockIdx.x*blockDim.x+threadIdx.x;
    int tidy =  blockIdx.y*blockDim.y+threadIdx.y;

    char* devPtr = (char*) devPitchedPtr.ptr;
    size_t pitch = devPitchedPtr.pitch;
    size_t slicePitch = pitch * N;

    for (int w = 0; w < W; w++) {
        char* slice = devPtr + w * slicePitch;
        float* row = (float*)(slice + tidy * pitch);
        row[tidx] = row[tidx] * row[tidx];
    }
}

/********/
/* MAIN */
/********/
int main()
{
    float a[N][M][W];

    for (int i=0; i<N; i++)
        for (int j=0; j<M; j++) 
            for (int w=0; w<W; w++) {
                a[i][j][w] = 3.f;
                //printf("row %i column %i depth %i value %f \n",i,j,w,a[i][j][w]);
            }

    // --- 3D pitched allocation and host->device memcopy
    cudaExtent extent = make_cudaExtent(M * sizeof(float), N, W);

    cudaPitchedPtr devPitchedPtr;
    gpuErrchk(cudaMalloc3D(&devPitchedPtr, extent));

    cudaMemcpy3DParms p = { 0 };
    p.srcPtr.ptr = a;
    p.srcPtr.pitch = M * sizeof(float);
    p.srcPtr.xsize = M;
    p.srcPtr.ysize = N;
    p.dstPtr.ptr = devPitchedPtr.ptr;
    p.dstPtr.pitch = devPitchedPtr.pitch;
    p.dstPtr.xsize = M;
    p.dstPtr.ysize = N;
    p.extent.width = M * sizeof(float);
    p.extent.height = N;
    p.extent.depth = W;
    p.kind = cudaMemcpyHostToDevice;
    gpuErrchk(cudaMemcpy3D(&p));

    dim3 GridSize(iDivUp(M,BLOCKSIZE_x),iDivUp(N,BLOCKSIZE_y));
    dim3 BlockSize(BLOCKSIZE_y,BLOCKSIZE_x);
    test_kernel_3D<<<GridSize,BlockSize>>>(devPitchedPtr);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    p.srcPtr.ptr = devPitchedPtr.ptr;
    p.srcPtr.pitch = devPitchedPtr.pitch;
    p.dstPtr.ptr = a;
    p.dstPtr.pitch = M * sizeof(float); 
    p.kind = cudaMemcpyDeviceToHost;
    gpuErrchk(cudaMemcpy3D(&p));

    for (int i=0; i<N; i++) 
        for (int j=0; j<M; j++) 
            for (int w=0; w<W; w++)
                printf("row %i column %i depth %i value %f\n",i,j,w,a[i][j][w]);

    getch();
    return 0;
}
于 2014-04-14T05:44:33.603 回答