CUDA finding the max value in given array

Question

I tried to develop a small CUDA program for find the max value in the given array,

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

max_value initialized by the first value of the input_data[0], The final answer is stored in result[0]. The kernel is giving 0 as the max value. I don't know what the problem is. I executed by 1 block 50 threads.

__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 
}

Even though there are in-built functions, just I am practicing small problems.

Answer 1

You are trying to set up a "critical section", but this approach on CUDA can lead to hang of your whole program - try to avoid it whenever possible.

Why your code hangs?

Your kernel (__global__ function) is executed by groups of 32 threads, called warps. All threads inside a single warp execute synchronously. So, the warp will stop in your do{} while(atomicCAS(&lock,0,1)) until all threads from your warp succeed with obtaining the lock. But obviously, you want to prevent several threads from executing the critical section at the same time. This leads to a hang.

Alternative solution

What you need is a "parallel reduction algorithm". You can start reading here:

• Parallel prefix sum @ wikipedia
• Parallel Reduction @ CUDA website
• NVIDIA's Guide to Reduction

Answer 2

Your code has potential race. I'm not sure if you defined the 'max_value' variable in shared memory or not, but both are wrong.

1) If 'max_value' is just a local variable, then each thread holds the local copy of it, which are not the actual maximum value (they are just the maximum value between input_data[0] and input_data[tid]). In the last line of code, all threads write to result[0] their own max_value, which will result in undefined behavior.

2) If 'max_value' is a shared variable, 49 threads will enter the if-statements block, and they will try to update the 'max_value' one at a time using locks. But the order of executions among 49 threads is not defined, and therefore some threads may overwrite the actual maximum value to smaller values. You would need to compare the maximum value again within the critical section.