일부 요소 쿠다

여기

에서 벡터 덧셈 계산되지 않은 코드이다 : 기본적일부 요소 쿠다

#include "common/book.h" 

#define N 36 

__global__ void add(int *a, int *b, int *c) { 
    int tid = blockIdx.x * gridDim.y * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x; 
    if(tid < N) { 
     c[tid] = a[tid] + b[tid]; 
    } 
} 

int main() { 
    int a[N], b[N], c[N]; 
    int *dev_a, *dev_b, *dev_c; 
    cudaMalloc((void**) &dev_a, N * sizeof(int)); 
    cudaMalloc((void**) &dev_b, N * sizeof(int)); 
    cudaMalloc((void**) &dev_c, N * sizeof(int)); 
    for (int i = 0; i < N; i++) { 
     a[i] = -1; 
     b[i] = i * i; 
    } 

    cudaMemcpy(
       dev_a, 
       a, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    cudaMemcpy(
       dev_b, 
       b, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    dim3 grid_dim(3, 2); 
    dim3 block_dim(3, 2); 
    add<<<grid_dim, block_dim>>>(dev_a, dev_b, dev_c); 
    cudaMemcpy(
       c, 
       dev_c, 
       N * sizeof(int), 
       cudaMemcpyDeviceToHost 
        ); 
    for (int i = 0; i < N; i++) { 
     printf("%d + %d = %d\n", a[i], b[i], c[i]); 
    } 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
}

, I는 3 × 2 레이아웃 그리드에 소자 현명한 개의 벡터를 추가하려고하고, 각 블록 그리드는 스레드의 3x2 레이아웃을 갖는다.

-1 + 0 = -1 
-1 + 1 = 0 
-1 + 4 = 3 
-1 + 9 = 8 
-1 + 16 = 15 
-1 + 25 = 24 
-1 + 36 = 0 
-1 + 49 = 0 
-1 + 64 = 0 
-1 + 81 = 0 
-1 + 100 = 0 
-1 + 121 = 0 
-1 + 144 = 143 
-1 + 169 = 168 
-1 + 196 = 195 
-1 + 225 = 224 
-1 + 256 = 255 
-1 + 289 = 288 
-1 + 324 = 0 
-1 + 361 = 0 
-1 + 400 = 0 
-1 + 441 = 0 
-1 + 484 = 0 
-1 + 529 = 0 
-1 + 576 = 575 
-1 + 625 = 624 
-1 + 676 = 675 
-1 + 729 = 728 
-1 + 784 = 783 
-1 + 841 = 840 
-1 + 900 = 0 
-1 + 961 = 0 
-1 + 1024 = 0 
-1 + 1089 = 0 
-1 + 1156 = 0 
-1 + 1225 = 0

은 분명히 일부 블록은 단지 무시된다 : 내가 컴파일 된 바이너리를 실행하면 다음

는 결과이다. 나는 또한 tid이 커널 함수 add에서 계산되는 방법을 가지고 놀려고 시도했지만 항상 누락 된 블록이 있습니다.

제안 사항?

출처

2017-12-20 qed

유일한 문제는 이미 추측 한대로 tid 계산에 있습니다.

매핑을 수행하고 산술을 작성하는 데는 여러 가지 방법이 있습니다. 범용 2D 그리드의 경우, x와 y에서 2D 인덱스를 생성하고 y 인덱스와 x 인덱스를 곱한 그리드 폭 (x 단위)을 사용하면 편리합니다 (즉, 기억하기 쉬운 방법론) 스레드 고유의 1-D 인덱스 만들 :

int idy = threadIdx.y+blockDim.y*blockIdx.y; // y-index 
int idx = threadIdx.x+blockDim.x*blockIdx.x; // x-index 
int tid = gridDim.x*blockDim.x*idy + idx;  // thread-unique 1D index

gridDim.x*blockDim.x는 X의 격자 폭을 스레드의 단위로 표시.

우리가 코드에서이 범용 2D 인덱싱 기법을 사용

, 나를 위해 제대로 작동하는 것 같다 :

$ cat t10.cu 
#include <stdio.h> 

#define N 36 

__global__ void add(int *a, int *b, int *c) { 
    int idy = threadIdx.y+blockDim.y*blockIdx.y; 
    int idx = threadIdx.x+blockDim.x*blockIdx.x; 
    int tid = gridDim.x*blockDim.x*idy + idx; 
    if(tid < N) { 
     c[tid] = a[tid] + b[tid]; 
    } 
} 

int main() { 
    int a[N], b[N], c[N]; 
    int *dev_a, *dev_b, *dev_c; 
    cudaMalloc((void**) &dev_a, N * sizeof(int)); 
    cudaMalloc((void**) &dev_b, N * sizeof(int)); 
    cudaMalloc((void**) &dev_c, N * sizeof(int)); 
    for (int i = 0; i < N; i++) { 
     a[i] = -1; 
     b[i] = i * i; 
    } 

    cudaMemcpy(
       dev_a, 
       a, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    cudaMemcpy(
       dev_b, 
       b, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    dim3 grid_dim(3, 2); 
    dim3 block_dim(3, 2); 
    add<<<grid_dim, block_dim>>>(dev_a, dev_b, dev_c); 
    cudaMemcpy(
       c, 
       dev_c, 
       N * sizeof(int), 
       cudaMemcpyDeviceToHost 
        ); 
    for (int i = 0; i < N; i++) { 
     printf("%d + %d = %d\n", a[i], b[i], c[i]); 
    } 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
} 
$ nvcc -arch=sm_35 -o t10 t10.cu 
$ cuda-memcheck ./t10 
========= CUDA-MEMCHECK 
-1 + 0 = -1 
-1 + 1 = 0 
-1 + 4 = 3 
-1 + 9 = 8 
-1 + 16 = 15 
-1 + 25 = 24 
-1 + 36 = 35 
-1 + 49 = 48 
-1 + 64 = 63 
-1 + 81 = 80 
-1 + 100 = 99 
-1 + 121 = 120 
-1 + 144 = 143 
-1 + 169 = 168 
-1 + 196 = 195 
-1 + 225 = 224 
-1 + 256 = 255 
-1 + 289 = 288 
-1 + 324 = 323 
-1 + 361 = 360 
-1 + 400 = 399 
-1 + 441 = 440 
-1 + 484 = 483 
-1 + 529 = 528 
-1 + 576 = 575 
-1 + 625 = 624 
-1 + 676 = 675 
-1 + 729 = 728 
-1 + 784 = 783 
-1 + 841 = 840 
-1 + 900 = 899 
-1 + 961 = 960 
-1 + 1024 = 1023 
-1 + 1089 = 1088 
-1 + 1156 = 1155 
-1 + 1225 = 1224 
========= ERROR SUMMARY: 0 errors 
$

가 위의 올바른 결과를 제공해야한다. 퍼포먼스와 관련해서는이 장난감 문제에 대해 가장 효율적인 매핑이 아닐 수도 있습니다. 이 문제는 일반적으로 효율적인 CUDA 프로그래밍에 권장되지 않는 32의 배수가 아닌 스레드 크기를가집니다. 이 경우에 대한 최적의 매핑 (성능/효율성면에서)을 시도하기보다는 스레드 블록을 재구성하여 적어도 블록 당 32 개의 스레드를 여러 개 제공하는 것이 좋습니다. 또한 적어도 고려할 것을 권장합니다. 16 또는 32 개의 스레드를 블록의 x- 차원에 배치하여 인덱싱을 쉽게 이해할 수있을뿐 아니라 대략적인 최적의 메모리 액세스 성능을 얻을 수 있습니다.

출처

2017-12-20 16:05:05

답변

관련 문제