Skip to content

Add initial support for gfx1151 #538

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
cgmillette merged 2 commits into from
May 9, 2025
Merged

Add initial support for gfx1151 #538

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ac73bb6
Add initial support for gfx1151
cgmillette May 7, 2025
6dafe44
Add more gfx1151 test cases
cgmillette May 9, 2025
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
2 changes: 1 addition & 1 deletion CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,7 +101,7 @@ if (ADDRESS_SANITIZER_ENABLED)
TARGETS "gfx90a:xnack+;gfx942:xnack+;gfx950:xnack+" )
else()
rocm_check_target_ids(DEFAULT_GPU_TARGETS
TARGETS "gfx908;gfx90a;gfx942;gfx950;gfx1100;gfx1101;gfx1102;gfx1200;gfx1201" )
TARGETS "gfx908;gfx90a;gfx942;gfx950;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201" )
endif()

# check if ROCm supports `__hip_fp8_e5m2` and `__hip_fp8_e4m3`
Expand Down
8 changes: 7 additions & 1 deletion library/include/rocwmma/internal/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,7 @@
/// ROCWMMA_ARCH_GFX1100
/// ROCWMMA_ARCH_GFX1101
/// ROCWMMA_ARCH_GFX1102
/// ROCWMMA_ARCH_GFX1151
/// ROCWMMA_ARCH_GFX1200
/// ROCWMMA_ARCH_GFX1201
///
Expand All @@ -63,6 +64,8 @@
#define ROCWMMA_ARCH_GFX1101 __gfx1101__
#elif defined(__gfx1102__) && ROCWMMA_DEVICE_COMPILE
#define ROCWMMA_ARCH_GFX1102 __gfx1102__
#elif defined(__gfx1151__) && ROCWMMA_DEVICE_COMPILE
#define ROCWMMA_ARCH_GFX1151 __gfx1151__
#elif defined(__gfx1200__) && ROCWMMA_DEVICE_COMPILE
#define ROCWMMA_ARCH_GFX1200 __gfx1200__
#elif defined(__gfx1201__) && ROCWMMA_DEVICE_COMPILE
Expand Down Expand Up @@ -94,6 +97,9 @@ static_assert(0, "Unsupported architecture");
#if !defined(ROCWMMA_ARCH_GFX1102)
#define ROCWMMA_ARCH_GFX1102 0
#endif
#if !defined(ROCWMMA_ARCH_GFX1151)
#define ROCWMMA_ARCH_GFX1151 0
#endif
#if !defined(ROCWMMA_ARCH_GFX1200)
#define ROCWMMA_ARCH_GFX1200 0
#endif
Expand Down Expand Up @@ -122,7 +128,7 @@ static_assert(0, "Unsupported architecture");
#define ROCWMMA_BLOCK_DIM_32_SUPPORTED 1
#endif

#if ROCWMMA_ARCH_GFX1100 || ROCWMMA_ARCH_GFX1101 || ROCWMMA_ARCH_GFX1102
#if ROCWMMA_ARCH_GFX1100 || ROCWMMA_ARCH_GFX1101 || ROCWMMA_ARCH_GFX1102 || ROCWMMA_ARCH_GFX1151
#define ROCWMMA_ARCH_GFX11 1
#define ROCWMMA_WAVE32_MODE 1
#define ROCWMMA_BLOCK_DIM_16_SUPPORTED 1
Expand Down
3 changes: 3 additions & 0 deletions library/include/rocwmma/internal/constants.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,7 @@ namespace rocwmma
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1100 = 0x1100;
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1101 = 0x1101;
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1102 = 0x1102;
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1151 = 0x1151;
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1200 = 0x1200;
static constexpr uint32_t AMDGCN_ARCH_ID_GFX1201 = 0x1201;
static constexpr uint32_t AMDGCN_ARCH_ID_NONE = 0x0000;
Expand Down Expand Up @@ -70,6 +71,8 @@ namespace rocwmma
static constexpr uint32_t AMDGCN_CURRENT_ARCH_ID = AMDGCN_ARCH_ID_GFX1101;
#elif ROCWMMA_ARCH_GFX1102
static constexpr uint32_t AMDGCN_CURRENT_ARCH_ID = AMDGCN_ARCH_ID_GFX1102;
#elif ROCWMMA_ARCH_GFX1151
static constexpr uint32_t AMDGCN_CURRENT_ARCH_ID = AMDGCN_ARCH_ID_GFX1151;
#elif ROCWMMA_ARCH_GFX1200
static constexpr uint32_t AMDGCN_CURRENT_ARCH_ID = AMDGCN_ARCH_ID_GFX1200;
#elif ROCWMMA_ARCH_GFX1201
Expand Down
4 changes: 3 additions & 1 deletion library/include/rocwmma/internal/wmma_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,8 @@ namespace rocwmma
TargetId,
Constants::AMDGCN_ARCH_ID_GFX1100,
Constants::AMDGCN_ARCH_ID_GFX1101,
Constants::AMDGCN_ARCH_ID_GFX1102> && Cond>;
Constants::AMDGCN_ARCH_ID_GFX1102,
Constants::AMDGCN_ARCH_ID_GFX1151> && Cond>;

// Enabler for all of gfx12
template <uint32_t TargetId, bool Cond = true>
Expand All @@ -66,6 +67,7 @@ namespace rocwmma
Constants::AMDGCN_ARCH_ID_GFX1100,
Constants::AMDGCN_ARCH_ID_GFX1101,
Constants::AMDGCN_ARCH_ID_GFX1102,
Constants::AMDGCN_ARCH_ID_GFX1151,
Constants::AMDGCN_ARCH_ID_GFX1200,
Constants::AMDGCN_ARCH_ID_GFX1201> && Cond>;

Expand Down
3 changes: 2 additions & 1 deletion samples/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,8 @@ bool isGfx11()

return ((deviceName.find("gfx1100") != std::string::npos)
|| (deviceName.find("gfx1101") != std::string::npos)
|| (deviceName.find("gfx1102") != std::string::npos));
|| (deviceName.find("gfx1102") != std::string::npos)
|| (deviceName.find("gfx1151") != std::string::npos));
}

bool isGfx12()
Expand Down
3 changes: 2 additions & 1 deletion test/dlrm/dlrm_kernel_base_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -108,7 +108,8 @@ namespace rocwmma
// Arch
auto isGfx908 = deviceArch == DeviceInfo::GFX908;
auto isGfx11 = (deviceArch == DeviceInfo::GFX1100) || (deviceArch == DeviceInfo::GFX1101)
|| (deviceArch == DeviceInfo::GFX1102);
|| (deviceArch == DeviceInfo::GFX1102)
|| (deviceArch == DeviceInfo::GFX1151);

auto isGfx12 = (deviceArch == DeviceInfo::GFX1200) || (deviceArch == DeviceInfo::GFX1201);

Expand Down
50 changes: 26 additions & 24 deletions test/gemm/gemm_kernel_base_dispatch_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -80,18 +80,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, SWITCH_BODY_TBLOCK_Y, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down Expand Up @@ -156,18 +157,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, SWITCH_BODY_TBLOCK_Y, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_KERNEL_FUNC_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_KERNEL_FUNC_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_KERNEL_FUNC_BODY

Expand Down
3 changes: 2 additions & 1 deletion test/gemm/gemm_test_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,12 +103,13 @@ namespace rocwmma
IsGfx1100 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1100),
IsGfx1101 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1101),
IsGfx1102 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1102),
IsGfx1151 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1151),

IsGfx1200 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1200),
IsGfx1201 = (ArchId == Constants::AMDGCN_ARCH_ID_GFX1201),

IsGfx9 = IsGfx908 || IsGfx90A || IsGfx942 || IsGfx950,
IsGfx11 = IsGfx1100 || IsGfx1101 || IsGfx1102,
IsGfx11 = IsGfx1100 || IsGfx1101 || IsGfx1102 || IsGfx1151,
IsGfx12 = IsGfx1200 || IsGfx1201,
};

Expand Down
4 changes: 4 additions & 0 deletions test/hip_device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,10 @@ namespace rocwmma
{
mGcnArch = hipGcnArch_t::GFX1102;
}
else if(deviceName.find("gfx1151") != std::string::npos)
{
mGcnArch = hipGcnArch_t::GFX1151;
}
else if(deviceName.find("gfx1200") != std::string::npos)
{
mGcnArch = hipGcnArch_t::GFX1200;
Expand Down
1 change: 1 addition & 0 deletions test/hip_device.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,7 @@ namespace rocwmma
GFX1100 = Constants::AMDGCN_ARCH_ID_GFX1100,
GFX1101 = Constants::AMDGCN_ARCH_ID_GFX1101,
GFX1102 = Constants::AMDGCN_ARCH_ID_GFX1102,
GFX1151 = Constants::AMDGCN_ARCH_ID_GFX1151,
GFX1200 = Constants::AMDGCN_ARCH_ID_GFX1200,
GFX1201 = Constants::AMDGCN_ARCH_ID_GFX1201,
UNSUPPORTED_ARCH = Constants::AMDGCN_ARCH_ID_NONE,
Expand Down
25 changes: 13 additions & 12 deletions test/unit/contamination_test/detail/load_contamination.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,18 +121,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, CASE_IMPL_ASSIGN2, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down
25 changes: 13 additions & 12 deletions test/unit/contamination_test/detail/store_contamination.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,18 +121,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, CASE_IMPL_ASSIGN2, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down
15 changes: 8 additions & 7 deletions test/unit/cross_lane_ops_test/detail/cross_lane_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,8 @@ namespace rocwmma
// gfx11 doesn't support dpp wave shift / rotate, bcast15x16 or bcast31x32
bool isGfx11 = (deviceArch == Base::DeviceInfo::GFX1100
|| deviceArch == Base::DeviceInfo::GFX1101
|| deviceArch == Base::DeviceInfo::GFX1102);
|| deviceArch == Base::DeviceInfo::GFX1102
|| deviceArch == Base::DeviceInfo::GFX1151);

bool isGfx12 = (deviceArch == Base::DeviceInfo::GFX1200)
|| (deviceArch == Base::DeviceInfo::GFX1201);
Expand Down Expand Up @@ -253,12 +254,12 @@ namespace rocwmma
// Map GTest params to Kernel params
using TestParamsT = std::tuple<Ts...>;
using KernelT = DppOpsKernel<
std::tuple_element_t<DataT, TestParamsT>, // DataT
std::tuple_element_t<CrossLaneOp, TestParamsT>, // CrossLaneOp
std::tuple_element_t<WriteRowMask, TestParamsT>::value, // WriteRowMask
std::tuple_element_t<WriteBankMask, TestParamsT>::value, // WriteBankMask
std::tuple_element_t<BoundCtrl, TestParamsT>::value // BoundCtrl
>;
std::tuple_element_t<DataT, TestParamsT>, // DataT
std::tuple_element_t<CrossLaneOp, TestParamsT>, // CrossLaneOp
std::tuple_element_t<WriteRowMask, TestParamsT>::value, // WriteRowMask
std::tuple_element_t<WriteBankMask, TestParamsT>::value, // WriteBankMask
std::tuple_element_t<BoundCtrl, TestParamsT>::value // BoundCtrl
>;

return std::make_shared<KernelT>();
}
Expand Down
25 changes: 13 additions & 12 deletions test/unit/fill_fragment_test/detail/fill_fragment.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,18 +103,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, CASE_IMPL_ASSIGN2, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down
25 changes: 13 additions & 12 deletions test/unit/layout_test/detail/col_layout.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -105,18 +105,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, CASE_IMPL_ASSIGN2, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down
25 changes: 13 additions & 12 deletions test/unit/layout_test/detail/colnt_layout.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -105,18 +105,19 @@ namespace rocwmma
ROCWMMA_SWITCH_BODY2_ARG2( \
waveSize, CASE_IMPL_ASSIGN2, HipDevice::Wave32, HipDevice::Wave64, ARCH_ID)

#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY9_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1200, \
HipDevice::GFX1201)
#define DISPATCH_GUARD_BODY \
ROCWMMA_SWITCH_BODY10_ARG1(deviceArch, \
SWITCH_BODY_WAVE_SIZE, \
HipDevice::GFX908, \
HipDevice::GFX90A, \
HipDevice::GFX942, \
HipDevice::GFX950, \
HipDevice::GFX1100, \
HipDevice::GFX1101, \
HipDevice::GFX1102, \
HipDevice::GFX1151, \
HipDevice::GFX1200, \
HipDevice::GFX1201)

DISPATCH_GUARD_BODY

Expand Down
Loading