블록 단위로 평균을 취함 Cuda

Question

Cuda는 공유 메모리가 동일한 블록의 데이터로만 공유 될 수 있다고 말했다. 그러나 블록에는 최대 1024 개의 스레드 만있을 수 있습니다. 거대한 행렬을 가지고 있고 최대화 된 스레드로 평균을 취하고 싶습니다.

예를 들어 설명합니다. 나는 4 개 개의 블록을 만들

#include <iostream> 
#include <stdio.h> 

__global__ void 
kernel(int *a, int dimx, int dimy) 
{ 
int ix = blockDim.x * blockIdx.x + threadIdx.x; 
int iy = blockDim.y * blockIdx.y + threadIdx.y; 

int idx = iy * dimx + ix; 

__shared__ int array[64]; 

a[idx] = a[idx] + 1; 

array[idx] = a[idx]; 

__syncthreads(); 

int sum=0; 
for(int i=0; i<dimx*dimy; i++) 
{ 
    sum += array[i]; 
} 

int average = sum/(dimx*dimy+1.0f); 

a[idx] = average; 

} 

int 
main() 
{ 
int dimx = 8; 
int dimy = 8; 
int num_bytes = dimx*dimy*sizeof(int); 

int *d_a=0, *h_a=0; // device and host pointers 
h_a = (int*)malloc(num_bytes); 

for (int i=0; i < dimx*dimy; i++){ 
    *(h_a+i) = i; 
} 
cudaMalloc((void**)&d_a, num_bytes); 

//cudaMemset(d_a, 0, num_bytes); 

cudaMemcpy(d_a, h_a, num_bytes, cudaMemcpyHostToDevice); 

dim3 grid, block; 
block.x = 4; 
block.y = 4; 
grid.x = dimx/block.x; 
grid.y = dimy/block.y; 

kernel<<<grid, block>>>(d_a, dimx, dimy); 

cudaMemcpy(h_a, d_a, num_bytes, cudaMemcpyDeviceToHost); 


std::cout << "the array a is:" << std::endl; 
for (int row = 0; row < dimy; row++) 
{ 
    for (int col =0; col < dimx; col++) 
    { 
     std::cout << h_a[row * dimx + col] << " "; 
    } 
    std::cout << std::endl; 
} 

free(h_a); 
cudaFree(d_a); 
} 

(난 그냥 데모로, 한 블록의 최대 스레드를 사용하지 않았다), 그리고 그들 모두의 평균으로 결과를 원한다. 이제 결과는 다음과 같습니다.

the array a is: 
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12 

각 블록의 평균은 오히려 전체 평균입니다. 모든 블록에 대해 평균을 취하는 방법은 무엇입니까?

저는 Cuda를 처음 사용합니다. 모든 관련 답변을 환영합니다.

Answer 1

가장 쉬운 방법은 블록 단위 평균을 수행하고 전역 메모리에 쓰고 다른 커널을 시작하여 이전 커널의 블록 단위 결과를 처리하는 등 여러 커널을 시작하는 것입니다. 데이터 차원에 따라이 과정을 여러 번 반복해야 할 수도 있습니다.

(의사 코드에서)

template <typename T> 
__global__ reduce(T* data, T* block_avgs) 
{ 
    //find the per-block average, write it out to block_avgs 
    //... 
} 

//in your caller: 
loop while you have more than 1 block: 
    call kernel using result from prev. iteration 
    update grid_dim and block_dim 

이것은 CUDA에서 블록 간 동기화가 필요하지 않기 때문에 필요합니다. 귀하의 문제는 감소의 아주 간단한 응용 프로그램입니다. 감축에 대한 더 나은 느낌을 얻으려면 평행 감소 샘플 at the nvidia samples page을보십시오.