From 3fdb7f2c454b64aa0de9ded2e0d6a5288ee2271e Mon Sep 17 00:00:00 2001 From: archana-ramalingam Date: Sat, 10 Oct 2026 19:10:48 +0000 Subject: [PATCH 1/2] test(hipblaslt): skip gfx1250-strict where the tests can't run Four TensileLite configs skip gfx1250 but not gfx1250-strict, and a base skip-gfx1250 mark does not fire on gfx1250-strict by design, so they run there and fail before reaching the GPU: - custom_rocblas_gemv_bf16_tn_m{1,2,4}: gfx950 hand-written assembly (v_dot2c_f32_bf16, s_waitcnt) that doesn't assemble for gfx12. - sk_queue_selfclean: SGEMM with a 32x32x2 MFMA tile, rejected as an invalid WMMA configuration (0 solutions). Add skip-gfx1250-strict to each. The MX fast_check tests (#13031) run every fp8/bf8 pair on '(950|1250)', which also matches gfx1250-strict. gfx1250 and gfx1250-strict have MX kernels for f8 x f8 only, so the bf8 pairs fail with "NO solution found". Keep f8 x f8 on '(950|1250)' and run the bf8 pairs on '950', like the other MX tests. The entry names are unchanged, so gfx950 test IDs stay the same. Co-Authored-By: Claude --- .../clients/tests/data/matmul_gtest.yaml | 45 +++++++++++++++++-- .../common/streamk/sk_queue_selfclean.yaml | 2 +- .../custom_rocblas_gemv_bf16_tn_m1.yaml | 2 +- .../custom_rocblas_gemv_bf16_tn_m2.yaml | 2 +- .../custom_rocblas_gemv_bf16_tn_m4.yaml | 2 +- 5 files changed, 45 insertions(+), 8 deletions(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index 5ee9a17fecc0..d03d14c91621 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -4850,12 +4850,31 @@ Tests: # shape here. ############################################################################### -# Every fp8 format pair, each solution launched and checked twice. +# Every fp8 format pair, each solution launched and checked twice. gfx1250 has +# MX kernels for f8 x f8 only, so the pairs with bf8 run on gfx950 alone. - name: matmul_fast_check_mx_fp8 category: pre_checkin function: matmul: - { a_type: f8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} + matrix_size: + - { M: 256, N: 256, K: 512 } + - { M: 1040, N: 528, K: 1024 } + transA: T + transB: N + alpha: 2.0 + beta: -2.0 + requested_solution_num: 16 + fast_check_repeat: 2 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + gpu_arch: '(950|1250)' + +- name: matmul_fast_check_mx_fp8 + category: pre_checkin + function: + matmul: - { a_type: bf8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - { a_type: bf8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} @@ -4871,7 +4890,7 @@ Tests: initialization: integer_exact unit_check: 0 fast_check: 1 - gpu_arch: '(950|1250)' + gpu_arch: '950' # The A-scale and B-scale buffers placed across a 4 GiB boundary, with poison # 4 GiB away (AIHPBLAS-4994): a scale read whose carry was dropped returns the @@ -4894,13 +4913,12 @@ Tests: fast_check: 1 gpu_arch: '(950|1250)' -# Large K, where the scale tensors span many blocks. +# Large K, where the scale tensors span many blocks. f8 x bf8 on gfx950 alone, as above. - name: matmul_fast_check_mx_fp8_large_k category: nightly function: matmul: - { a_type: f8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} matrix_size: - { M: 3072, N: 3072, K: 16384 } - { M: 2560, N: 16384, K: 2560 } @@ -4915,6 +4933,25 @@ Tests: fast_check: 1 gpu_arch: '(950|1250)' +- name: matmul_fast_check_mx_fp8_large_k + category: nightly + function: + matmul: + - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} + matrix_size: + - { M: 3072, N: 3072, K: 16384 } + - { M: 2560, N: 16384, K: 2560 } + transA: T + transB: N + alpha: 2.0 + beta: -2.0 + requested_solution_num: 16 + fast_check_repeat: 2 + initialization: integer_exact + unit_check: 0 + fast_check: 1 + gpu_arch: '950' + - name: matmul_96_complex category: pre_checkin function: diff --git a/projects/hipblaslt/tensilelite/tensilelite/Tests/common/streamk/sk_queue_selfclean.yaml b/projects/hipblaslt/tensilelite/tensilelite/Tests/common/streamk/sk_queue_selfclean.yaml index 058f2fcad383..91c3f8f4cdad 100644 --- a/projects/hipblaslt/tensilelite/tensilelite/Tests/common/streamk/sk_queue_selfclean.yaml +++ b/projects/hipblaslt/tensilelite/tensilelite/Tests/common/streamk/sk_queue_selfclean.yaml @@ -13,7 +13,7 @@ # that do not block evenly over the queues. TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # NumXCD is 8 only on gfx942/gfx950; elsewhere the single queue is trivially balanced + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] # NumXCD is 8 only on gfx942/gfx950; elsewhere the single queue is trivially balanced GlobalParameters: # DynamicWorkQueue hybrid sub-path; the static sub-path has no work queue. diff --git a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m1.yaml b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m1.yaml index ad0eacad6363..b4b1bbdf3a73 100644 --- a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m1.yaml +++ b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m1.yaml @@ -2,7 +2,7 @@ # SPDX-License-Identifier: MIT TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) GlobalParameters: MinimumRequiredVersion: 5.0.0 diff --git a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m2.yaml b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m2.yaml index 58e312eef495..fcd698e5814e 100644 --- a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m2.yaml +++ b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m2.yaml @@ -2,7 +2,7 @@ # SPDX-License-Identifier: MIT TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) GlobalParameters: MinimumRequiredVersion: 5.0.0 diff --git a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m4.yaml b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m4.yaml index c6ed79c38656..71bce10328ce 100644 --- a/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m4.yaml +++ b/projects/hipblaslt/tensilelite/tensilelite/Tests/custom/custom_rocblas_gemv_bf16_tn_m4.yaml @@ -2,7 +2,7 @@ # SPDX-License-Identifier: MIT TestParameters: - marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] # wvSpltK BF16: gfx950 (v_dot2c_f32_bf16) GlobalParameters: MinimumRequiredVersion: 5.0.0 From d51015f597bca94043c5ce09155ad6b5442a4e65 Mon Sep 17 00:00:00 2001 From: archana-ramalingam Date: Mon, 12 Oct 2026 03:55:47 +0000 Subject: [PATCH 2/2] test(hipblaslt): limit the bf8 MX rows to gfx950 per row, not by split Instead of splitting matmul_fast_check_mx_fp8 and _large_k into separate gfx950 entries, tag the four rows with bf8 inputs with gpu_arch '950'. The generator merges each row over its entry, so the generated tests are the same, and the diff drops from 45 changed lines to 13. Co-Authored-By: Claude --- .../clients/tests/data/matmul_gtest.yaml | 48 +++---------------- 1 file changed, 6 insertions(+), 42 deletions(-) diff --git a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml index d03d14c91621..9aa07ef5d24c 100755 --- a/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml +++ b/projects/hipblaslt/clients/tests/data/matmul_gtest.yaml @@ -4850,13 +4850,16 @@ Tests: # shape here. ############################################################################### -# Every fp8 format pair, each solution launched and checked twice. gfx1250 has -# MX kernels for f8 x f8 only, so the pairs with bf8 run on gfx950 alone. +# Every fp8 format pair, each solution launched and checked twice. gfx1250 has MX +# kernels for f8 x f8 only, so the rows with bf8 run on gfx950 alone. - name: matmul_fast_check_mx_fp8 category: pre_checkin function: matmul: - { a_type: f8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} + - { a_type: bf8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r, gpu_arch: '950'} + - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r, gpu_arch: '950'} + - { a_type: bf8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r, gpu_arch: '950'} matrix_size: - { M: 256, N: 256, K: 512 } - { M: 1040, N: 528, K: 1024 } @@ -4871,27 +4874,6 @@ Tests: fast_check: 1 gpu_arch: '(950|1250)' -- name: matmul_fast_check_mx_fp8 - category: pre_checkin - function: - matmul: - - { a_type: bf8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - - { a_type: bf8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - matrix_size: - - { M: 256, N: 256, K: 512 } - - { M: 1040, N: 528, K: 1024 } - transA: T - transB: N - alpha: 2.0 - beta: -2.0 - requested_solution_num: 16 - fast_check_repeat: 2 - initialization: integer_exact - unit_check: 0 - fast_check: 1 - gpu_arch: '950' - # The A-scale and B-scale buffers placed across a 4 GiB boundary, with poison # 4 GiB away (AIHPBLAS-4994): a scale read whose carry was dropped returns the # poison, which E8M0 reads as 2^-87, and D comes out wrong instead of faulting. @@ -4919,6 +4901,7 @@ Tests: function: matmul: - { a_type: f8_r, b_type: f8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} + - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r, gpu_arch: '950'} matrix_size: - { M: 3072, N: 3072, K: 16384 } - { M: 2560, N: 16384, K: 2560 } @@ -4933,25 +4916,6 @@ Tests: fast_check: 1 gpu_arch: '(950|1250)' -- name: matmul_fast_check_mx_fp8_large_k - category: nightly - function: - matmul: - - { a_type: f8_r, b_type: bf8_r, c_type: f32_r, d_type: f32_r, compute_type: c_f32_r, scaleA: 3, scaleB: 3, scale_type: f32_r} - matrix_size: - - { M: 3072, N: 3072, K: 16384 } - - { M: 2560, N: 16384, K: 2560 } - transA: T - transB: N - alpha: 2.0 - beta: -2.0 - requested_solution_num: 16 - fast_check_repeat: 2 - initialization: integer_exact - unit_check: 0 - fast_check: 1 - gpu_arch: '950' - - name: matmul_96_complex category: pre_checkin function: