CUDA 在给定数组中查找最大值

Question

我试图开发一个小的 CUDA 程序来查找给定数组中的最大值，

int input_data[0...50] = 1,2,3,4,5....,50

max_value由 input_data[0] 的第一个值初始化,
最终答案存储在 result[0] .
内核给出 0 作为最大值。我不知道是什么问题。
我由 1 个块 50 个线程执行。

__device__ int lock=0;

__global__ void max(float *input_data,float *result)
{
     float max_value = input_data[0];
     int  tid = threadIdx.x;

     if( input_data[tid] > max_value)
     {
         do{} while(atomicCAS(&lock,0,1));
         max_value=input_data[tid];
         __threadfence();
         lock=0;
      }

    __syncthreads();
    result[0]=max_value;  //Final result of max value 
}

尽管有内置函数，但我只是在练习小问题。

Answer 1

您正在尝试设置一个“关键部分”，但是 CUDA 上的这种方法可能会导致整个程序挂起 - 尽可能避免它。

为什么你的代码挂起？

您的内核( __global__ 函数)由 32 个线程组执行，称为扭曲。单个 warp 中的所有线程同步执行。因此，经纱将停止在您的 do{} while(atomicCAS(&lock,0,1)) 中。直到您的经纱中的所有线程都成功获得锁。但很明显，您希望防止多个线程同时执行临界区。这会导致挂起。

替代解决方案

您需要的是“并行缩减算法”。你可以从这里开始阅读:

Parallel prefix sum @ wikipedia

Parallel Reduction @ CUDA website

NVIDIA's Guide to Reduction