Repository navigation
[rocprofiler-compute] Add gfx arch. based pre-processor guards and runtime checks in rocflop.cpp - #2487
Conversation
There was a problem hiding this comment.
Pull request overview
This pull request removes MFMA (Matrix Fused Multiply-Add) functionality from the rocflop sample to fix build issues on MI50 GPUs, which do not support the __builtin_amdgcn_mfma_f32* intrinsics. The removal is clean and comprehensive, eliminating all MFMA-related code paths.
Key Changes:
- Removed MFMA kernel functions (
matmul_fp16_throughputandmatmul_fp32_throughput) that used unsupported intrinsics - Removed MFMA-related test infrastructure and result tracking
- Removed command-line options for MFMA tests (
--matfp16and--matfp32)
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
feizheng10
left a comment
There was a problem hiding this comment.
Technically, the code should be designed as detecting arch properly, right?
|
@benrichard-amd , this might be useful for you to expand support for MI50 and MI100 |
|
@feizheng10 can you have a look again |
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 1 out of 1 changed files in this pull request and generated 4 comments.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Hi. I updated https://github.com/benrichard-amd/rocflop with improved compile time and runtime detection. |
Motivation
Add gfx arch. based pre-processor guards and runtime checks in rocflop.cpp for following support matrix:
MFMA (Matrix operation) : Only supported > gfx906 (greater than MI 50)
SMFMAC (sparse MFMA): Only supported > gfx908 (greater than MI 100)
This solves the build issue on MI 50 and MI 100 where the following intrinsics dont work respectively
__builtin_amdgcn_mfma_f32*and__builtin_amdgcn_smfmac_f32*This issue was introduced in #2418
This PR is related to #2456
Technical Details
JIRA ID
Test Plan
In
tests/CMakeLists.txtsetset(CMAKE_HIP_ARCHITECTURES "gfx906" CACHE STRING "GPU architectures to compile for" FORCE)to build rocprofiler-compute for MI 50 even though the GPU on the system might be different.In
tests/CMakeLists.txtsetset(CMAKE_HIP_ARCHITECTURES "gfx908" CACHE STRING "GPU architectures to compile for" FORCE)to build rocprofiler-compute for MI 50 even though the GPU on the system might be different.Test Result
Ensure build works fine in both the cases
Submission Checklist