CUDA에서 최대 행렬 찾기

Question

방금 ​​CUDA로 시작했습니다. 이제 나는 질문이있다. N * N 행렬이 있고 창 크기는 8x8입니다. 이 행렬을 여러 개의 부분 행렬로 세분화하고이 최대 값을 찾으십시오. 예를 들어 64 * 64 행렬을 가지고 있으므로 8 * 8 크기의 8 개의 작은 행렬을 가지며 8 개의 최대 값을 찾습니다. 마지막으로 모든 최대 값을 새 배열에 저장하지만 순서는 항상 변경됩니다. 내가 바로 순서대로 유지하는 해결책을 찾기 원하는 것은

__global__ void calculate_emax_kernel(float emap[],float emax[], int img_height, int img_width,int windows_size) 
{ 
    int x_index = blockIdx.x*blockDim.x+threadIdx.x; 
    int y_index = blockIdx.y*blockDim.y+threadIdx.y; 

    int num_row_block = img_height/windows_size; 
    int num_col_block = img_width/windows_size; 
    __shared__ float window_elements[256]; 
    __shared__ int counter; 
    __shared__ int emax_count; 

    if (threadIdx.x == 0) emax_count = 0; 
    __syncthreads(); 
    int index; 
    int emax_idx = 0; 


    if(y_index >= img_height|| x_index >= img_width) return; 
    for(int i = 0; i < num_row_block; i++) 
    { 
     for(int j = 0; j < num_col_block; j++) 
     { 
      counter = 0; 
      if(y_index >= i*windows_size && y_index < (i+1)*windows_size 
        && x_index >= j*windows_size && x_index < (j+1)*windows_size) 
      { 
       int idx = y_index*img_height + x_index; 
       index = atomicAdd(&counter, 1); 

       window_elements[index] = emap[idx]; 
       __syncthreads(); 


       // reduction 
       unsigned int k = (windows_size*windows_size)/2; 
       while(k != 0) 
       { 
        if(index < k) 
        { 
         window_elements[index] = fmaxf(window_elements[index], window_elements[index+k]); 

        } 
        k /= 2; 
       } 
       if(index == 0) 
       { 
        emax[i*num_row_block+j] = window_elements[index]; 
       } 
      } 
      __syncthreads(); 
     } 
     __syncthreads(); 
    } 
    __syncthreads(); 
} 

이 내 구성 CUDA와

void construct_emax(float *input,float *output, int img_height, int img_width) 
{ 
    int windows_size = 4; 
    float * d_input, * d_output; 
    cudaMalloc(&d_input, img_width*img_height*sizeof(float)); 
    cudaMalloc(&d_output, img_width*img_height*sizeof(float)); 

    cudaMemcpy(d_input, input, img_width*img_height*sizeof(float), cudaMemcpyHostToDevice); 
    dim3 blocksize(16,16); 
    dim3 gridsize; 

    gridsize.x=(img_width+blocksize.x-1)/blocksize.x; 
    gridsize.y=(img_height+blocksize.y-1)/blocksize.y; 

    calculate_emax_kernel<<<gridsize,blocksize>>>(d_input,d_output,img_height,img_width,windows_size); 

} 

Answer 1

, parallel reduction이 까다 롭습니다이다; segmented parallel reduction은 더 까다 롭습니다. 이제 2D로 작업하고 있으며 세그먼트/창은 스레드 블록보다 작습니다.

큰 창 크기의 경우 문제가 아닌 것 같습니다. 하나의 스레드 블록을 사용하여 하나의 창을 줄일 수 있습니다. 예를 들어 16x16 창을 사용하는 경우 16x16 스레드 블록을 사용할 수 있습니다. 예를 들어 64x64와 같이 더 큰 창 크기가있는 경우에도 여전히 16x16 스레드 블록을 사용할 수 있습니다. 먼저 데이터로드 중에 64x64 창을 16x16 요소로 줄인 다음 스레드 블록 내에서 1 스칼라로 줄입니다.

블록 크기보다 작은 창에 대해서는 성능 향상을 위해 스레드 블록 당 여러 개의 창을 줄여야합니다. 현재 256 개의 스레드 블록 (16x16)이 16 개의 4x4 창을 담당하는 현재 블록/그리드 구성을 사용할 수 있습니다. 그러나 각 32 스레드 랩이 2 부분 (2x16)으로 구성되어 있기 때문에 이는 최적이 아닙니다. 이것은 coalesced global memory access에 좋지 않으며 효율적인 병렬 축소를 위해 2x16 워프를 하나 이상의 4x4 창에 매핑하기가 어렵습니다.

또는 256 스레드의 1-D 스레드 블록을 사용하는 것이 좋습니다. m 개의 스레드는 모두 m x m 창을 줄입니다. 그런 다음 2 차원 그리드를 사용하여 전체 이미지를 커버 할 수 있습니다. 커널 함수에서

const int m = window_size; 
dim3 blocksize(256); 
dim3 gridsize((img_width+255)/256, (img_height+m-1)/m); 

, 당신은

1. 글로벌 데이터 로딩 중 1 배 m 벡터 각 m X m 창을 감소시킬 수있다;
2. 1x m 벡터를 스칼라로 줄이려면 트리 감소 방법을 사용하십시오.

다음 코드는 m이 2의 제곱수이고 m <= 32 일 때 작동하는 개념적 데모입니다. 임의로 m 및 더 나은 경계 확인을 위해이를 수정할 수 있습니다.

#include <assert.h> 
#include <cuda.h> 
#include <thrust/device_vector.h> 

__global__ void calculate_emax_kernel(const float* input, float* output, 
             int height, int width, int win_size, 
             int out_width) { 
    const int tid = threadIdx.x; 
    const int i = blockIdx.y * win_size; 
    const int j = blockIdx.x * 256 + tid; 
    const int win_id = j % win_size; 

    __shared__ float smax[256]; 

    float tmax = -1e20; 
    if (j < width) { 
    for (int tile = 0; tile < win_size; tile++) { 
     if (i + tile < height) { 
     tmax = max(tmax, input[(i + tile) * width + j]); 
     } 
    } 
    } 
    smax[tid] = tmax; 
    for (int shift = win_size/2; shift > 0; shift /= 2) { 
    if (win_id < shift) { 
     smax[tid] = max(smax[tid], smax[tid + shift]); 
    } 
    } 
    if (win_id == 0 && j < width) { 
    output[blockIdx.y * out_width + (j/win_size)] = smax[tid]; 
    } 
} 

int main() { 
    const int height = 1024; 
    const int width = 1024; 
    const int m = 4; 
    thrust::device_vector<float> in(height * width); 
    thrust::device_vector<float> out(
     ((height + m - 1)/m) * ((width + m - 1)/m)); 

    dim3 blocksize(256); 
    dim3 gridsize((width + 255)/256, (height + m - 1)/m); 

    assert(m == 2 || m == 4 || m == 8 || m == 16 || m == 32); 
    calculate_emax_kernel<<<gridsize, blocksize>>>(
     thrust::raw_pointer_cast(in.data()), 
     thrust::raw_pointer_cast(out.data()), 
     height, width, m, (width + m - 1)/m); 

    return 0; 
}