공유 메모리를 사용하는 3 차원 배열의 1 차원 이상의 합계

Question

다음과 같은 계산을 수행해야합니다. A [x] [y] = 합계 (z = 0에서 z = n까지) {B [x] [y] [ 행렬 A는 차원 [높이] [너비]와 행렬 B를 가지며, C는 차원 [높이] [너비] [n]을 갖는다] [z] + C [x] [y] [z] 나는 각 블록은 공유 메모리를 자신이 있기 때문에 각 블록은 하나 개의 합계를 계산하고 싶은

index = 0; 
for (z = 0; z<n; ++z) 
    for(y = 0; y<width; ++y) 
     for(x = 0; x<height; ++x) { 
      matrix[index] = value; 
      index++; 
     } 

:

값은 같은과 메모리에 매핑됩니다.

dim3 block (n, 1, 1); 
dim grid (height, width, 1); 

커널 : 글로벌 메모리의 코드

부 : 나는 atomicAdd,이 같은 것을 사용 데이터 경주을 방지하기 위해

atomicAdd(&(A[blockIdx.x + blockIdx.y*gridDim.y]), 
      B[blockIdx.x + blockIdx.y*gridDim.y+threadIdx.x*blockDim.x*blockDim.y] 
      + C[blockIdx.x + blockIdx.y*gridDim.y+threadIdx.x*blockDim.x*blockDim.y]); 

나는 합계를 계산하는 공유 메모리를 사용하고 싶습니다 이 결과를 전역 메모리로 복사하십시오.

공유 메모리로 작업하는 방법을 잘 모르겠습니다. 각 블록의 공유 메모리에는 하나의 숫자 (합계 결과) 만 저장됩니다. 이 숫자를 전역 메모리에있는 A 행렬의 올바른 위치에 어떻게 복사해야합니까?

Answer 1

사용자가 요구하는 합계를 수행하기 위해 공유 메모리 또는 원자 적 메모리 액세스가 필요하지 않을 수도 있습니다. 이것이 올바르게 이해 되었다면, 데이터가 열의 주요 순서에 있으므로 논리적 연산은 출력 행렬에 행렬 항목당 하나의 스레드를 가지며 각 스레드가 입력 행렬의 z 축을 가로 질러 가면서 합산하는 것입니다. 이 방법 width * height >= n와 열 주요 데이터에 대한 최적해야

__global__ void kernel(float *A, const float *B, const float *C, 
     const int width, const int height, const int n) 
{ 
    int tidx = threadIdx.x + blockDim.x * blockIdx.x; 
    int tidy = threadIdx.y + blockDim.y * blockIdx.y; 

    if ((tidx < height) && (tidy < width)) { 
     int stride = width * height; 
     int ipos = tidx + tidy * height; 

     float * oval = A + ipos; 
     float sum = 0.f; 
     for(int z=0; z<n; z++, ipos+=stride) { 
      sum += B[ipos] + C[ipos]; 
     } 
     *oval = sum; 
    } 
} 

이에 대한 커널은 같은 것을 볼 수 있었다. 이를 위해 공유 메모리를 사용하면 성능상의 이점이 없으며 아토믹 메모리 작업을 사용할 필요가 없습니다. 문제가 있다면 width * height << n 합계 당 블록 현명한 병렬 감소를 시도하는 것이 좋습니다. 그러나 당신은 문제의 전형적인 차원이 무엇인지 밝히지 않았다. 문제가 후자와 더 비슷하다면 의견을 남기고 답장에 기반한 샘플 커널을 추가 할 수 있습니다.