Проблемы с уменьшением в 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.