Skip to content

GFX1250 changes with updated AITER #732

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

Open
ipanfilo wants to merge 10 commits into
base: dev
Choose a base branch
from
Open

GFX1250 changes with updated AITER #732

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
99d2207
Update to new QoLA/CK-JIT/AITER
ipanfilo Aug 19, 2026
5cced7a
Update CI prebuild CK blob list
ipanfilo Aug 23, 2026
d668d29
Merge remote-tracking branch 'origin/dev' into ipanfilo/update_aiter_…
ipanfilo Aug 24, 2026
88b9669
Merge remote-tracking branch 'origin/dev' into ipanfilo/update_aiter_…
ipanfilo Aug 28, 2026
4f651ad
Skip known failing test config
ipanfilo Aug 28, 2026
1e97485
Update QoLA with CK patches fix, restore skip test, add custom QoLA s…
ipanfilo Aug 31, 2026
12ad9e7
Update CK-JIT and QoLA to main branch-es commits
ipanfilo Sep 2, 2026
9048129
Gfx1250 changes on top of updated AITER
ipanfilo Sep 3, 2026
fe1cf36
Enable CK tests. Enable MXFP4, MXFP8 and NVFP4 in core, still disable…
ipanfilo Sep 3, 2026
5b657ff
Merge branch 'dev' into ipanfilo/gfx1250_aiter0813_a50a62a
ipanfilo Sep 3, 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
2 changes: 1 addition & 1 deletion 3rdparty/QoLA
Open in desktop
Submodule QoLA updated 9 files
+12 −12 patches/aiter/0002-hsa-only-aiter-hip-common.patch
+0 −114 patches/aiter/0003-ck-te-receipt-700.patch
+48 −14 patches/aiter/0005-bwd-ondevice-workspace-prepare.patch
+6 −6 patches/aiter/0007-mha-bwd-v3-api-check-dry-run.patch
+371 −0 patches/aiter/0008-mha-fwd-native-splitkv.patch
+254 −0 patches/aiter/0009-mha-fwd-with-sink.patch
+9 −1 qola/cpp_itfs/qola_mha_fwd.cu
+6 −1 qola/cpp_itfs/qola_mha_fwd.h
+0 −5 qola/cpp_itfs/registry.toml
2 changes: 1 addition & 1 deletion 3rdparty/ck_jit
Open in desktop
Submodule ck_jit updated 3 files
+7 −6 ck_build_interceptor.py
+98 −23 ck_jit_build.py
+100 −0 codegen_jit_hints_22ee91463.patch
78 changes: 67 additions & 11 deletions ci/ck_jit_prebuild.txt
View file
Open in desktop

Large diffs are not rendered by default.

2 changes: 1 addition & 1 deletion examples/jax/encoder/common.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,7 +56,7 @@ def is_nvfp4_supported():
gpu_arch = get_device_compute_capability(0)
if is_hip_extension():
# only GFX12.5 machines support nvfp4
return False #TODO add gfx1250 (gpu_arch == 125) when ready
return gpu_arch >= 120

@github-actions github-actions Bot Sep 3, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

The condition contradicts the comment directly above it. gpu_arch >= 120 also matches gfx1200/gfx1201 (RDNA4), which per the comment do not support NVFP4 — only gfx1250 does. On those parts the example would now claim NVFP4 support and fail at runtime.

Suggested change
return gpu_arch >= 120
return gpu_arch == 125

return gpu_arch >= 100


Expand Down
11 changes: 4 additions & 7 deletions tests/cpp/operator/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,21 +58,18 @@ if(USE_ROCM)
test_dequantize_nvfp4.cu
test_cublaslt_gemm.cu
test_ck_grouped_gemm.cu
test_ck_grouped_mxfp8.cu
test_cast_mxfp4_transpose.cu
test_multi_quantize_mxfp8.cu)

list(FIND CMAKE_HIP_ARCHITECTURES "gfx1250" _gfx1250_idx)
if(NOT _gfx1250_idx EQUAL -1)
list(APPEND test_cuda_sources test_ck_grouped_mxfp8.cu)
target_include_directories(test_operator BEFORE PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/../../../3rdparty/composable_kernel/include)
endif()

TE_GetHipifiedSources("${test_cuda_sources}" ${CMAKE_CURRENT_SOURCE_DIR} test_hip_sources)
TE_AddHipifyDeps("${test_cuda_sources}" ${CMAKE_CURRENT_SOURCE_DIR})
message("${message_line}")
message(STATUS "test_operator hipified sources: ${test_hip_sources}")
set_target_properties(test_operator PROPERTIES SOURCES "${test_hip_sources}")
target_include_directories(test_operator BEFORE PRIVATE
Comment on lines 68 to +70

@github-actions github-actions Bot Sep 3, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

This include directory was previously added only when gfx1250 was in CMAKE_HIP_ARCHITECTURES; it's now unconditional for every ROCm test build. Two consequences worth confirming:

  • 3rdparty/composable_kernel becomes a hard build requirement for all ROCm test builds, not just gfx1250 ones. Fine if git submodule update --init --recursive is always assumed, but it's a new failure mode for partial checkouts.
  • BEFORE puts CK's include/ ahead of every other include path for the whole test_operator target. CK ships generically-named headers, so this can shadow same-named headers from TE/HIP/gtest in translation units that never needed CK. Scoping it to the one consumer would be safer — test_ck_grouped_mxfp8.cu is the only file that includes ck_tile/*, and it already guards itself at runtime (GTEST_SKIP unless gfx1250, line 504-508). A per-source set_source_files_properties(... INCLUDE_DIRECTORIES ...) would avoid the global reordering.

${CMAKE_CURRENT_SOURCE_DIR}/../../../3rdparty/composable_kernel/include)

endif()

# Find required packages
Expand Down
2 changes: 2 additions & 0 deletions tests/cpp/operator/test_cast_mxfp4_transpose.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -463,11 +463,13 @@ TEST_P(FusedCastTransposeMXFP4TestSuite, TestFusedCastTransposeMXFP4) {
// Forward activations
auto OP = &identity;
switch (Act_type) {
case ActivationType::Identity: OP = &identity; break;
case ActivationType::GeLU: OP = &gelu; break;
case ActivationType::SiLU: OP = &silu; break;
case ActivationType::ReLU: OP = &relu; break;
case ActivationType::QGeLU: OP = &qgelu; break;
case ActivationType::SReLU: OP = &srelu; break;
default: GTEST_FAIL() << "Unsupported activation type"; break;
}

TRANSFORMER_ENGINE_TYPE_SWITCH_FP16_FP32_ONLY(input_type, InputType,
Expand Down
4 changes: 4 additions & 0 deletions tests/cpp/operator/test_cast_mxfp8.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -658,11 +658,13 @@ TEST_P(FusedCastMXFP8TestSuite, TestFusedCastMXFP8) {
// Forward activations
auto OP = &identity;
switch (Act_type) {
case ActivationType::Identity: OP = &identity; break;
case ActivationType::GeLU: OP = &gelu; break;
case ActivationType::SiLU: OP = &silu; break;
case ActivationType::ReLU: OP = &relu; break;
case ActivationType::QGeLU: OP = &qgelu; break;
case ActivationType::SReLU: OP = &srelu; break;
default: GTEST_FAIL() << "Unsupported activation type"; break;
}

TRANSFORMER_ENGINE_TYPE_SWITCH_FP16_FP32_ONLY(input_type, InputType,
Expand All @@ -681,11 +683,13 @@ TEST_P(FusedCastMXFP8TestSuite, TestFusedCastMXFP8) {
} else {
auto OP = &identity;
switch (Act_type) {
case ActivationType::Identity: OP = &identity; break;
case ActivationType::GeLU: OP = &dgelu; break;
case ActivationType::SiLU: OP = &dsilu; break;
case ActivationType::ReLU: OP = &drelu; break;
case ActivationType::QGeLU: OP = &dqgelu; break;
case ActivationType::SReLU: OP = &dsrelu; break;
default: GTEST_FAIL() << "Unsupported activation type"; break;
}
TRANSFORMER_ENGINE_TYPE_SWITCH_FP16_FP32_ONLY(input_type, InputType,
TRANSFORMER_ENGINE_TYPE_SWITCH_FP8_ONLY(output_type, OutputType,
Expand Down
2 changes: 2 additions & 0 deletions tests/cpp/operator/test_cast_nvfp4_transpose.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1453,11 +1453,13 @@ TEST_P(FusedCastTransposeNVFP4TestSuite, TestFusedCastTransposeNVFP4) {
// Forward activations
auto OP = &identity;
switch (Act_type) {
case ActivationType::Identity: OP = &identity; break;
case ActivationType::GeLU: OP = &gelu; break;
case ActivationType::SiLU: OP = &silu; break;
case ActivationType::ReLU: OP = &relu; break;
case ActivationType::QGeLU: OP = &qgelu; break;
case ActivationType::SReLU: OP = &srelu; break;
default: GTEST_FAIL() << "Unsupported activation type"; break;
}

TRANSFORMER_ENGINE_TYPE_SWITCH_FP16_FP32_ONLY(input_type, InputType,
Expand Down
89 changes: 49 additions & 40 deletions tests/cpp/operator/test_cublaslt_gemm.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -515,6 +515,47 @@ std::pair<double, double> getTestTolerances(const DType type, bool use_fp8, bool
}


void checkMxFP8Support(const TestParams& params, const cudaDeviceProp& prop, bool &use_mxfp8, bool &use_hipkittens_mxfp8) {

@github-actions github-actions Bot Sep 3, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

GTEST_SKIP() inside a helper does not abort the caller.

GTEST_SKIP() expands to return GTEST_MESSAGE_(msg, TestPartResult::kSkip) — it returns from the enclosing function only. Previously all of these skips lived directly in performTest/performDqTest, so hitting one aborted the test body. Now they return from checkMxFP8Support, and performTest continues executing the full GEMM afterwards.

Consequence: on a device where MXFP8 isn't supported (e.g. gfx942 running OperatorTestMXFP8), or on a shape failing the M/N/K alignment guards, the test is reported as skipped but the unsupported GEMM still runs — likely an NVTE_CHECK abort, HIP error, or bogus comparison rather than a clean skip. Same for performDqTest.

Suggested fix: have the helper report rather than skip, and skip at the call site, e.g. return a std::optional<std::string> (reason) or an out-param skip_reason, then in the caller:

if (auto reason = checkMxFP8Support(params, prop, use_mxfp8, use_hipkittens_mxfp8)) {
  GTEST_SKIP() << *reason;
}

use_mxfp8 = (params.scaling_mode == NVTEScalingMode::NVTE_MXFP8_1D_SCALING);
if (!use_mxfp8) {
use_hipkittens_mxfp8 = false;
return;
}
#ifdef __HIP_PLATFORM_AMD__
if (!(prop.major == 9 && prop.minor >= 5) && !(prop.major >= 12)) {
GTEST_SKIP() << "MXFP8 requires gfx950 or newer";
}
#endif
if (params.m % 16 || params.n % 16) {
GTEST_SKIP() << "MXFP8 requires M & N to be multiples of 16";
}

#ifdef __HIP_PLATFORM_AMD__
const size_t required_k_multiple = (prop.major == 12 && prop.minor == 5) ? 32 : 128;
#else
const size_t required_k_multiple = 128;
#endif
if (params.k % required_k_multiple) {
GTEST_SKIP() << "MXFP8 requires K to be a multiple of " << required_k_multiple;
}

use_hipkittens_mxfp8 = !params.force_hipblaslt;
if (!use_hipkittens_mxfp8) {
return;
}
#ifdef __HIP_PLATFORM_AMD__
if (!(prop.major == 9 && (prop.minor == 4 || prop.minor == 5))) {
GTEST_SKIP() << "HipKittens requires gfx942 or gfx950";
}
if (params.m % 256 || params.n % 256 || params.k < 256) {
GTEST_SKIP() << "HipKittens requires M and N 256-aligned, K >= 256";
}
#else
GTEST_SKIP() << "HipKittens requires ROCm";
#endif
}


template <typename A_Type, typename B_Type, typename Bias_Type, typename Gelu_Type, typename D_Type>
void performTest(const TestParams& params) {
DType atype = TypeInfo<A_Type>::dtype;
Expand All @@ -524,30 +565,19 @@ void performTest(const TestParams& params) {
DType dtype = TypeInfo<D_Type>::dtype;

const bool has_fp8 = isFp8Type(atype) || isFp8Type(btype);
const bool use_mxfp8 = params.scaling_mode == NVTEScalingMode::NVTE_MXFP8_1D_SCALING;
const bool use_hipkittens_mxfp8 = use_mxfp8 && !params.force_hipblaslt;

cudaDeviceProp prop;
(void)cudaGetDeviceProperties(&prop, 0);

bool use_mxfp8 = false;
bool use_hipkittens_mxfp8 = false;
checkMxFP8Support(params, prop, use_mxfp8, use_hipkittens_mxfp8);

if (use_mxfp8)
{
if (!has_fp8) {
GTEST_SKIP() << "MXFP8 scaling mode requires Float8 types";
}
if (params.m % 16 || params.n % 16) {
GTEST_SKIP() << "MXFP8 requires M & N to be multiples of 16";
}
size_t required_k_multiple = 128;
#ifdef __HIP_PLATFORM_AMD__
required_k_multiple = (prop.major == 12 && prop.minor == 5) ? 32 : 128;
#endif
if (params.k % required_k_multiple) {
GTEST_SKIP() << "MXFP8 requires K to be a multiple of " << required_k_multiple;
}
if (use_hipkittens_mxfp8 && (params.m % 256 || params.n % 256 || params.k < 256)) {
GTEST_SKIP() << "HipKittens requires M and N 256-aligned, K >= 256";
}
}

#ifdef __HIP_PLATFORM_AMD__
Expand Down Expand Up @@ -580,10 +610,6 @@ void performTest(const TestParams& params) {
if (!fp8_supported) {
GTEST_SKIP() << "FP8 is not supported in current config";
}
const bool mxfp8_supported = (prop.major == 9 && prop.minor >= 5) || prop.major >= 12;
if (use_mxfp8 && !mxfp8_supported) {
GTEST_SKIP() << "MXFP8 is not supported in current config";
}
if (!use_hipkittens_mxfp8 && params.use_bias) {
GTEST_SKIP() << "MXFP8 GEMM with bias is not supported by hipBLASLt";
}
Expand Down Expand Up @@ -759,28 +785,11 @@ void performDqTest(const TestParams &params) {
cudaDeviceProp prop;
(void)cudaGetDeviceProperties(&prop, 0);

if (params.m % 16 || params.n % 16) {
GTEST_SKIP() << "MXFP8 requires M & N to be multiples of 16";
}
size_t required_k_multiple = 128;
#ifdef __HIP_PLATFORM_AMD__
required_k_multiple = (prop.major == 12 && prop.minor == 5) ? 32 : 128;
#endif
if (params.k % required_k_multiple) {
GTEST_SKIP() << "MXFP8 requires K to be a multiple of " << required_k_multiple;
}

bool mxfp8_supported = (prop.major == 9 && prop.minor >= 5) || prop.major >= 12;
const bool use_hipkittens_mxfp8 = !params.force_hipblaslt;
if (!mxfp8_supported) {
GTEST_SKIP() << "MXFP8 is not supported in current config";
}
bool _unused = false;
checkMxFP8Support(params, prop, _unused, _unused);

@github-actions github-actions Bot Sep 3, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

Passing the same object for two independent non-const out-params aliases them. It happens to be harmless today only because DqGEMMTestSuite is instantiated with use mxfp8 = Values(true) (line 992), so the early if (!use_mxfp8) { use_hipkittens_mxfp8 = false; return; } branch is never taken. If a non-MXFP8 param is ever added to that suite, the aliasing makes the helper silently skip all of the alignment/arch guards that performDqTest used to do unconditionally.

Two separate locals would remove the coupling and read better:

Suggested change
checkMxFP8Support(params, prop, _unused, _unused);
bool dq_use_mxfp8 = false;
bool dq_use_hipkittens = false;
checkMxFP8Support(params, prop, dq_use_mxfp8, dq_use_hipkittens);

if (params.use_bias || params.use_gelu) {
GTEST_SKIP() << "DqGEMMTestSuite does not yet have reference for bias/gelu epilogues";
}
if (use_hipkittens_mxfp8 && (params.m % 256 || params.n % 256 || params.k % 128 || params.k < 256)) {
GTEST_SKIP() << "HipKittens requires M and N 256-aligned, K >= 256";
}

// hipBLASLt on gfx950 produces incorrect results for certain MXFP8
// GEMMs with non-TN layouts.
Expand Down Expand Up @@ -940,7 +949,7 @@ INSTANTIATE_TEST_SUITE_P(OperatorTest, GEMMTestSuite,
::testing::Values(false, true), //use_gelu
::testing::ValuesIn(kLayouts), //transa,transb
::testing::Values(false), //use mxfp8
::testing::Values(false)), //force hipblaslt
::testing::Values(true)), //force hipblaslt

@github-actions github-actions Bot Sep 3, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

Is this flip intentional (same on line 961)? Both suites pass use mxfp8 = false, and NVTE_ROCM_USE_HIPBLASLT_MXFP8 is only consulted under if (is_mxfp8 && kittens_mxfp8_supported()) (rocm_gemm.cu:2280), so the execution path is unchanged.

What does change is every test name in OperatorTest/OperatorTestFP8: GEMMTestName appends std::get<5>(...) ? "HB" : "HK", so all of these rename from …xHK to …xHB. Any saved --gtest_filter or downstream job pinned to the old names silently stops matching. If the goal is just to stop labelling non-MXFP8 tests as "HipKittens", it may be cleaner to make the suffix conditional on use_mxfp8 rather than flipping a parameter that has no effect.

GEMMTestName);

INSTANTIATE_TEST_SUITE_P(OperatorTestFP8, FP8GEMMTestSuite,
Expand All @@ -949,7 +958,7 @@ INSTANTIATE_TEST_SUITE_P(OperatorTestFP8, FP8GEMMTestSuite,
::testing::Values(false, true), //use_gelu
::testing::ValuesIn(kLayouts), //transa,transb
::testing::Values(false), //use mxfp8
::testing::Values(false)), //force hipblaslt
::testing::Values(true)), //force hipblaslt
GEMMTestName);

INSTANTIATE_TEST_SUITE_P(OperatorTestMXFP8, FP8GEMMTestSuite,
Expand Down
Loading
Loading