Skip to content

[MIOpen] Port LRN kernels from OpenCL to HIP #4887

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

Merged
BradPepersAMD merged 5 commits into from
Mar 12, 2026
Merged

[MIOpen] Port LRN kernels from OpenCL to HIP #4887

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
f492d08
[MIOpen] Port LRN kernels from OpenCL to HIP
EwanC Feb 25, 2026
279c5e6
Add in explicit loop unrolling for `MIOpenLRNWithinChannel_PS`
EwanC Mar 4, 2026
5854eb7
Add loop unrolling to Bwd Within LRN kernel
EwanC Mar 6, 2026
f39a0e0
Peel Bwd within loops
EwanC Mar 9, 2026
eb13453
Address PR comments
EwanC Mar 11, 2026
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
4 changes: 2 additions & 2 deletions projects/miopen/src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -574,8 +574,8 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
kernels/MIOpenGetitem.cpp
kernels/MIOpenKthvalue.cpp
kernels/MIOpenLayerNorm.cpp
kernels/MIOpenLRNBwd.cl
kernels/MIOpenLRNFwd.cl
kernels/MIOpenLRNBwd.cpp
kernels/MIOpenLRNFwd.cpp
kernels/MIOpenMultiMarginLoss.cpp
kernels/MIOpenNeuron.cpp
kernels/MIOpenPReLU.cpp
Expand Down
4 changes: 2 additions & 2 deletions projects/miopen/src/include/miopen/mlo_internal.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -557,8 +557,8 @@ struct mlo_construct_norm : mlo_construct_activ_lrn_pooling_common
void mloConstruct();

protected:
int mloConstructFwd();
int mloConstructBwd();
void mloConstructFwd();
void mloConstructBwd();
int _norm_region = 0;
int _norm_area = 0;
double _normAlpha = 0.0;
Expand Down
Loading
Loading