From a672de89f564b323c653c2285e095bb0051e4a89 Mon Sep 17 00:00:00 2001
From: Abishek <52214183+r-abishek@users.noreply.github.com>
Date: Wed, 31 Jan 2024 13:40:50 -0800
Subject: [PATCH] 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                       |  6 +++---
 src/modules/hip/kernel/histogram.cpp | 24 ++++++++++++------------
 2 files changed, 15 insertions(+), 15 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 6f2b54622..f159179bd 100644
--- a/CMakeLists.txt
+++ b/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}")
 
diff --git a/src/modules/hip/kernel/histogram.cpp b/src/modules/hip/kernel/histogram.cpp
index aebcfe422..9b94bb1a7 100644
--- a/src/modules/hip/kernel/histogram.cpp
+++ b/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 ))