Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Build fix for ROCm 6.0.0 gfx942 #295

Merged
merged 4 commits into from
Jan 31, 2024
Merged

Build fix for ROCm 6.0.0 gfx942 #295

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
80ea779
modified histogram hip kernel to use atomicAdd instead of atomicInc
sampath1117 Jan 29, 2024
363e305
Merge pull request #227 from sampath1117/sr/histogram_build_fix
r-abishek Jan 29, 2024
7f6c08d
Update CMakeLists.txt
kiritigowda Jan 31, 2024
b5d66f0
Update CMakeLists.txt
kiritigowda Jan 31, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -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}")

Expand Down
24 changes: 12 additions & 12 deletions src/modules/hip/kernel/histogram.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is atomicAdd good option. I thought we can't use any atomics

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@rrawther Ubuntu build seems to pass with atomicAdd. Ideally, both should work as they did on openCL/cuda/hip, but for some reason only atomicInc() fails on Rocm6.0.0 HIP

atomicAdd(&tmp_histogram[pixelG], 1);
atomicAdd(&tmp_histogram[pixelB], 1);
}
__syncthreads();
if (local_size >= (256 ))
Expand Down Expand Up @@ -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 ))
Expand Down Expand Up @@ -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 ))
Expand Down Expand Up @@ -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 ))
Expand Down