From 80ea7794495974b137fa927e51f0ac7e41c259a5 Mon Sep 17 00:00:00 2001 From: sampath1117 Date: Mon, 29 Jan 2024 06:11:44 +0000 Subject: [PATCH 1/3] modified histogram hip kernel to use atomicAdd instead of atomicInc --- src/modules/hip/kernel/histogram.cpp | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) 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 )) From 7f6c08d495df371b5413244bdd3780f5b1d73533 Mon Sep 17 00:00:00 2001 From: Kiriti Gowda Date: Wed, 31 Jan 2024 10:42:44 -0800 Subject: [PATCH 2/3] Update CMakeLists.txt Turn Advance build by default --- CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6f2b54622..0a1b88610 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,7 +122,7 @@ 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("-- ${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 6.0.0+ Use -D BUILD_WITH_AMD_ADVANCE=ON to support MI300+${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}") From b5d66f0061ed5de1e95db886035e7394c2601735 Mon Sep 17 00:00:00 2001 From: Kiriti Gowda Date: Wed, 31 Jan 2024 12:26:17 -0800 Subject: [PATCH 3/3] Update CMakeLists.txt Note update for -D option --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 0a1b88610..f159179bd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -123,7 +123,7 @@ 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:ON)]${ColourReset}") -message("-- ${Yellow} NOTE: For ROCm Version 6.0.0+ Use -D BUILD_WITH_AMD_ADVANCE=ON to support MI300+${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}")