CUDA의 축소 구현 관련 문제

Question

배열의 최대 요소를 찾을 수있는 cuda를 구현하려고합니다. 이 커널을 사용하여 최소값을 찾았지만 작동하지만 최대 값을 찾으려고하면 작동하지 않습니다. 반복적으로 알고리즘을 밟았으므로 버그를 찾을 수 없습니다. 어떤 도움이라도 정말로 감사 할 것입니다. 코드 몇 가지 문제가 있습니다

__global__ 
void findMaxAndMin(const float* const d_logLuminance, float* reduceCopy, int length, float* min_logLum, float* max_logLum){ 
    int idx = threadIdx.x + blockDim.x*blockIdx.x; 
    if(idx >= length){ 
     return; 
    } 
    reduceCopy[idx] = d_logLuminance[idx]; 
    __syncthreads(); 

    //do a reduction with max 

    for(int offset = 1;offset < length;offset = offset*2){ 
     if(idx % (offset*2) == 0){ 
      int compIdx = idx + offset; 
      if(compIdx < length){ 
       float newVal = a_max(reduceCopy[idx], reduceCopy[compIdx]); 
       if(idx == 0){ 
        //printf("val %f \n", newVal); 
       } 
       __syncthreads(); 
       reduceCopy[idx] = newVal; 
       __syncthreads(); 
      } 
     } 
     __syncthreads(); 
    } 
    __syncthreads(); 
    if(idx == 0){ 
     *max_logLum = reduceCopy[0]; 
    } 

} 

Answer 1

(내가 거기에 인쇄 문 주석을 제거 할 때 추가로, 나는 ....도 꽤 두통이다, 다른 출력을 얻을). 최소한 효과가 있었다면 운이 좋았을 것입니다.

• 두 블록 이상을 실행한다고 가정합니다 (blockIdx.x). 한 블록의 결과를 다른 블록에서 반복적으로 사용하면 reduceCopy[compIdx]이 다른 블록에 의해 설정 될 수 있습니다. 블록의 실행 순서를 예측하거나 동기화 할 수는 없습니다. __syncthreads()은 단일 블록 내에서만 작동하는 장벽입니다!

• if(idx >= length) return은 모든 스레드가 다음에 도달하지 않을 수도 있기 때문에 위험합니다. __syncthreads.

• __syncthreads()은 분기 조건 if(compIdx < length) 내에 있습니다.

• a_max은 정의되지 않습니다. 항상 최소 작업 예제를 포함해야합니다. 나는 그 기능이 무엇을 해야하는지 짐작할 수있다. 그러나 그 안에 또 다른 버그가 있을지도 모른다.

이론상 병렬 감소에 대해 잘 알고 있지만 구현이 CUDA 특정 동작으로 인해 실패한 것 같습니다.

CUDA를 사용하여 병렬 감소를 수행하는 방법에 대한 몇 가지 예를 읽어 보시기 바랍니다.