cudaMemcpy3D를 사용하여 *** 포인터를 전송하십시오.

Question

cudaMemcpy3D을 사용하여 동적으로 할당 된 3 차원 행렬 (텐서)을 전송하려고합니다. Tensor는 인접한 메모리 블록으로 할당됩니다 (아래 코드 참조). 내가 cudaExtent과 cudaMemcpy3DParms의 다양한 조합을 시도했지만, 요소의 순서가 뒤섞여 있습니다. 내가 잘못 뭐하는 거지

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 
... 

처럼 보이는, 호스트 변수 (h_tensor) 및 장치 (d_tensor) 차이에 대한

#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 
} 

출력 : 나는이 문제를 설명하기 위해 다음 예제를 만들어? cudaMemcpy3D을 사용하는 올바른 방법은 무엇입니까?

Answer 1

1. 언제든지 cuda 코드에 문제가있는 경우 proper cuda error checking을 수행하는 것이 좋습니다. 여기에 게시 한 코드는 적어도 나를 위해 올바르게 실행되지 않습니다. cudaMemcpy3D 줄에서 오류가 발생합니다. 이것은 아래의 2 번 항목 때문입니다. (난 당신이 출력이 여기에 표시 한 코드와 동일하지 않았다 생성하는 데 사용되는 코드를 의심하지만 그건 그냥 추측입니다.)
는 make_cudaPitchedPtr 귀하의 사용은 올바른

2. 하지 :

   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 차원의 잘못된 이해에 속성 발견 메모리의 값은 인접한 인덱스 값을 차지합니다. 이 때문에 제 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 
}