使用cudaMemcpy3D傳輸***指針

Question

我正在嘗試使用cudaMemcpy3D傳輸動態分配的3D矩陣（張量）。張量被分配爲連續的內存塊（見下面的代碼）。我嘗試了cudaExtent和cudaMemcpy3DParms的各種組合，但元素的順序變得混亂起來。我創建了下面的例子來說明這個問題：

#include <stdio.h> 

int ***alloc_tensor(int Nx, int Ny, int Nz) { 
    int i, j; 
    int ***tensor; 

    tensor = (int ***) malloc((size_t) (Nx * sizeof(int **))); 
    tensor[0] = (int **) malloc((size_t) (Nx * Ny * sizeof(int *))); 
    tensor[0][0] = (int *) malloc((size_t) (Nx * Ny * Nz * sizeof(int))); 

    for(j = 1; j < Ny; j++) 
     tensor[0][j] = tensor[0][j-1] + Nz; 
    for(i = 1; i < Nx; i++) { 
     tensor[i] = tensor[i - 1] + Ny; 
     tensor[i][0] = tensor[i - 1][0] + Ny * Nz; 
     for(j = 1; j < Ny; j++) 
     tensor[i][j] = tensor[i][j - 1] + Nz; 
    } 

    return tensor; 
} 

__global__ void kernel(cudaPitchedPtr tensor, int Nx, int Ny, int Nz) { 
    int i, j, k; 
    char *tensorslice; 
    int *tensorrow; 

    for (i = 0; i < Nx; i++) { 
     for (j = 0; j < Ny; j++) { 
     for (k = 0; k < Nz; k++) { 
      tensorslice = ((char *)tensor.ptr) + k * tensor.pitch * Nx; 
      tensorrow = (int *)(tensorslice + i * tensor.pitch); 
      printf("d_tensor[%d][%d][%d] = %d\n", i, j, k, tensorrow[j]); 
     } 
     } 
    } 
} 

int main() { 
    int i, j, k, value = 0; 
    int Nx = 2, Ny = 6, Nz = 4; 

    int ***h_tensor; 
    struct cudaPitchedPtr d_tensor; 

    h_tensor = alloc_tensor(Nx, Ny, Nz); 
    cudaMalloc3D(&d_tensor, make_cudaExtent(Nx * sizeof(int), Ny, Nz)); 

    for(i = 0; i < Nx; i++) { 
     for(j = 0; j < Ny; j++) { 
     for(k = 0; k < Nz; k++) { 
      h_tensor[i][j][k] = value++; 
      printf("h_tensor[%d][%d][%d] = %d\n", i, j, k, h_tensor[i][j][k]); 
     } 
     } 
    } 

    cudaMemcpy3DParms cpy = { 0 }; 
    cpy.srcPtr = make_cudaPitchedPtr(h_tensor[0][0], Nx * sizeof(int), Ny, Nz); 
    cpy.dstPtr = d_tensor; 
    cpy.extent = make_cudaExtent(Nx * sizeof(int), Ny, Nz); 
    cpy.kind = cudaMemcpyHostToDevice; 

    cudaMemcpy3D(&cpy); 

    kernel<<<1, 1>>>(d_tensor, Nx, Ny, Nz); 

    // ... clean-up 
} 

輸出的主變量（h_tensor）和設備（d_tensor）不同，看上去就像

h_tensor[0][0][0] = 0 
h_tensor[0][0][1] = 1 
h_tensor[0][0][2] = 2 
h_tensor[0][0][3] = 3 
h_tensor[0][1][0] = 4 
h_tensor[0][1][1] = 5 
h_tensor[0][1][2] = 6 
... 

d_tensor[0][0][0] = 0 
d_tensor[0][0][1] = 12 
d_tensor[0][0][2] = 24 
d_tensor[0][0][3] = 36 
d_tensor[0][1][0] = 1 
d_tensor[0][1][1] = 13 
d_tensor[0][1][2] = 25 
... 

我在做什麼錯？什麼是使用cudaMemcpy3D的正確方法？

Answer 1

1. 任何時候如果遇到cuda代碼有問題，最好做proper cuda error checking。您在此發佈的代碼至少不會正確運行 - cudaMemcpy3D行會引發錯誤。這是由於下面的項目2。 （我懷疑你用來生成輸出是不相同的，你在這裏顯示的代碼的代碼，但是這只是一個猜測。）

2. 你的make_cudaPitchedPtr用法是不正確的：

   cpy.srcPtr = make_cudaPitchedPtr(h_tensor[0][0], Nx * sizeof(int), Ny, Nz); 
   

   審查API文檔。以這種方式製作CUDA投球指針在2D和3D之間沒有區別。所以，如你所做的那樣傳遞3個不同的維度是沒有意義的。代替這樣做：

   cpy.srcPtr = make_cudaPitchedPtr(h_tensor[0][0], Nx * sizeof(int), Nx, Ny); 
   

3. 遺留問題，我發現我的屬性，以C. 3個維的不正確的認識乘法標數組的最後一個標爲快速變化的尺寸，即它是一個在相鄰內存中的值佔用相鄰的索引值。由於這個原因，您在第三維中使用Z會讓我感到困惑。您的主機分配在第一個下標位置使用Nx，但您的設備索引不匹配。顯然有多種方式來處理這個問題。如果你不喜歡我的安排，你可以改變它，但主機和設備索引必須匹配。

不管怎樣，下面的代碼修改爲我工作：

#include <stdio.h> 

int ***alloc_tensor(int Nx, int Ny, int Nz) { 
    int i, j; 
    int ***tensor; 

    tensor = (int ***) malloc((size_t) (Nx * sizeof(int **))); 
    tensor[0] = (int **) malloc((size_t) (Nx * Ny * sizeof(int *))); 
    tensor[0][0] = (int *) malloc((size_t) (Nx * Ny * Nz * sizeof(int))); 

    for(j = 1; j < Ny; j++) 
     tensor[0][j] = tensor[0][j-1] + Nz; 
    for(i = 1; i < Nx; i++) { 
     tensor[i] = tensor[i - 1] + Ny; 
     tensor[i][0] = tensor[i - 1][0] + Ny * Nz; 
     for(j = 1; j < Ny; j++) 
     tensor[i][j] = tensor[i][j - 1] + Nz; 
    } 

    return tensor; 
} 

__global__ void kernel(cudaPitchedPtr tensor, int Nx, int Ny, int Nz) { 
    int i, j, k; 
    char *tensorslice; 
    int *tensorrow; 

    for (i = 0; i < Nx; i++) { 
     for (j = 0; j < Ny; j++) { 
     for (k = 0; k < Nz; k++) { 
      tensorslice = ((char *)tensor.ptr) + k * tensor.pitch * Ny; 
      tensorrow = (int *)(tensorslice + j * tensor.pitch); 
      printf("d_tensor[%d][%d][%d] = %d\n", i, j, k, tensorrow[i]); 
     } 
     } 
    } 
} 

int main() { 
    int i, j, k, value = 0; 
    int Nx = 2, Ny = 6, Nz = 4; 

    int ***h_tensor; 
    struct cudaPitchedPtr d_tensor; 

    h_tensor = alloc_tensor(Nz, Ny, Nx); 
    cudaMalloc3D(&d_tensor, make_cudaExtent(Nx * sizeof(int), Ny, Nz)); 

    for(i = 0; i < Nx; i++) { 
     for(j = 0; j < Ny; j++) { 
     for(k = 0; k < Nz; k++) { 
      h_tensor[k][j][i] = value++; 
      //printf("h_tensor[%d][%d][%d] = %d\n", i, j, k, h_tensor[i][j][k]); 
     } 
     } 
    } 
    for(i = 0; i < Nx; i++) { 
     for(j = 0; j < Ny; j++) { 
     for(k = 0; k < Nz; k++) { 
      //h_tensor[i][j][k] = value++; 
      printf("h_tensor[%d][%d][%d] = %d\n", i, j, k, h_tensor[k][j][i]); 
     } 
     } 
    } 

    cudaMemcpy3DParms cpy = { 0 }; 
    cpy.srcPtr = make_cudaPitchedPtr(h_tensor[0][0], Nx * sizeof(int), Nx, Ny); 
    cpy.dstPtr = d_tensor; 
    cpy.extent = make_cudaExtent(Nx * sizeof(int), Ny, Nz); 
    cpy.kind = cudaMemcpyHostToDevice; 

    cudaMemcpy3D(&cpy); 

    kernel<<<1, 1>>>(d_tensor, Nx, Ny, Nz); 
    cudaDeviceSynchronize(); 
    // ... clean-up 
}