Build fix - ROCm 6.0.0 gfx942 (#295)

* modified histogram hip kernel to use atomicAdd instead of atomicInc

* Update CMakeLists.txt

Turn Advance build by default

* Update CMakeLists.txt

Note update for -D option

---------

Co-authored-by: sampath1117 <sampath.rachumallu@multicorewareinc.com>
Co-authored-by: Kiriti Gowda <kiriti.nageshgowda@amd.com>

CMakeLists.txt

@@ -54,7 +54,7 @@ endif(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
 
 # RPP Default Options
 set(DEFAULT_BUILD_TYPE "Release")
-option(BUILD_WITH_AMD_ADVANCE "Build RPP for advanced AMD GPU Architecture"    OFF)
+option(BUILD_WITH_AMD_ADVANCE "Build RPP for advanced AMD GPU Architecture"    ON)
 
 # Set message options
 if(NOT WIN32)
@@ -122,8 +122,8 @@ else()
 endif()
 
 message("-- ${Cyan}RPP Developer Options${ColourReset}")
-message("-- ${Cyan}     -D BUILD_WITH_AMD_ADVANCE=${BUILD_WITH_AMD_ADVANCE} [Turn ON/OFF Build for AMD advanced GPUs(default:OFF)]${ColourReset}")
-message("-- ${Yellow}   NOTE: For ROCm Version 6.0.0+ Use -D BUILD_WITH_AMD_ADVANCE=ON to support MI300+${ColourReset}")
+message("-- ${Cyan}     -D BUILD_WITH_AMD_ADVANCE=${BUILD_WITH_AMD_ADVANCE} [Turn ON/OFF Build for AMD advanced GPUs(default:ON)]${ColourReset}")
+message("-- ${Yellow}     NOTE: For ROCm Version less than 6.0.0+ Use -D BUILD_WITH_AMD_ADVANCE=OFF${ColourReset}")
 message("-- ${Cyan}     -D BACKEND=${BACKEND} [Select RPP Backend [options:CPU/OPENCL/HIP](default:HIP)]${ColourReset}")
 message("-- ${Cyan}     -D CMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} [Select RPP build type [options:Debug/Release](default:Release)]${ColourReset}")
 

src/modules/hip/kernel/histogram.cpp

@@ -39,9 +39,9 @@ void partial_histogram_pln( unsigned char *input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + width * height];
         unsigned char pixelB = input[pixId + 2 * width * height];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -98,9 +98,9 @@ void partial_histogram_pkd( unsigned char *input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + 1];
         unsigned char pixelB = input[pixId + 2];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -166,9 +166,9 @@ void partial_histogram_batch( unsigned char* input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + inc[id_z]];
         unsigned char pixelB = input[pixId + 2 * inc[id_z]];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -234,9 +234,9 @@ void partial_histogram_semibatch( unsigned char* input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + inc];
         unsigned char pixelB = input[pixId + 2 * inc];
-       atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+       atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))