Skip to content

Revert "fix(hipblaslt): correct MI300A kernel selection in the WGMXCC predicate and KD-tree search" #13479

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
arjun-raj-kuppala merged 1 commit into from
Oct 12, 2026
Merged

Revert "fix(hipblaslt): correct MI300A kernel selection in the WGMXCC predicate and KD-tree search" #13479

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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/hipblaslt/clients/tests/data/matmul_common.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ Definitions:
- { M: 24000, N: 256, K: 256, lda: 24000, ldb: 256, ldc: 24020, ldd: 24020 }

- &grid_limit_matrix_size_real
- { M: 2, N: 2097152, K: 1 }
- { M: 2, N: 2097152, K:1 }
- { M: 64, N: 1048576, K: 1 }
- { M: 256, N: 1048576, K: 1 }
- { M: 1024, N: 524288, K: 1 }
Expand All @@ -68,7 +68,7 @@ Definitions:
- { M: 4096, N: 131072, K: 1 }

- &grid_limit_matrix_size_double
- { M: 2, N: 1048576, K: 1 }
- { M: 2, N: 1048576, K:1 }
- { M: 3, N: 2, K: 1, lda: 4, ldb: 5, ldc: 6, ldd: 6}
- { M: 64, N: 1048576, K: 1 }
- { M: 256, N: 1048576, K: 1 }
Expand Down
29 changes: 0 additions & 29 deletions projects/hipblaslt/clients/tests/data/matmul_gtest.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,35 +127,6 @@ Tests:
- {name: alpha_beta_zero_NaN, category: pre_checkin, precision: *real_precisions,
function: matmul, transA: N, transB: N, M: 256, N: 128, K: 64, alpha: [ .NaN, 2 ], beta: [ .NaN, 2 ] }

# Port of the rocBLAS gemm_ex case hipblaslt_f16 (gemm_gtest.yaml, use_hipblaslt: 1),
# instance NN_2_2_2_5_2_2_0_2_2: a 2x2x2 NN GEMM with alpha != 1 and beta == 0, run
# directly through hipBLASLt instead of via the rocBLAS backend.
- name: matmul_hipblaslt_f16_2x2x2_NN
category: pre_checkin
function:
matmul: *hpa_half_precision
matrix_size:
- { M: 2, N: 2, K: 2, lda: 2, ldb: 2, ldc: 2, ldd: 2 }
transA: N
transB: N
alpha: 5
beta: 0

# Port of the rocBLAS gemm_ex case gemm_ex_h2f (f16 in, f32 out), instance
# TT_1_1_64_1_64_1_1_1_1. A 1x1x64 TT GEMM: a single dot product of length 64, which
# the GridBased heuristic answers from the (M=1,N=1) cell whose only tabulated K is 1.
- name: matmul_hipblaslt_h2f_1x1x64_TT
category: pre_checkin
function:
matmul: { a_type: f16_r, b_type: f16_r, c_type: f32_r, d_type: f32_r,
compute_type: c_f32_r, scale_type: f32_r }
matrix_size:
- { M: 1, N: 1, K: 64, lda: 64, ldb: 1, ldc: 1, ldd: 1 }
transA: T
transB: T
alpha: 1
beta: 1

# Special floating-point inputs on 64x64, FP types only. The inf, neg_inf and nan cases exercise
# check_special_value_consistency(): CPU and GPU must agree on the IEEE class (no Inf->NaN).
# The zero and neg_zero cases are ordinary initialization coverage that flows through the normal
Expand Down
27 changes: 2 additions & 25 deletions projects/hipblaslt/tensilelite/include/Tensile/ContractionProblemPredicates.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3048,25 +3048,6 @@ namespace TensileLite
// value = [XCC, XCCG]
std::array<int, 2> value;
size_t cuCount;
AMDGPU::Processor processor;

/// Generic gfx942 solutions are CU-fallbacks on MI300A (228 CUs, vs 304 for
/// MI300X) and always launch at WGMXCC=1, but this runs before
/// setFallbackStatus() marks them, so fallbackStatus() alone misses them.
bool isMI300A() const
{
static constexpr size_t MI300ACuCount = 228;
return processor == AMDGPU::Processor::gfx942 && cuCount == MI300ACuCount;
}

/// True when this solution will be launched at WGMXCC=1 regardless of the
/// XCC baked into it, so the predicate must be evaluated with XCC=1.
/// On MI300A this is unconditionally true: every solution there takes the
/// WGMXCC=1 launch path, so isMI300A() short-circuits the whole check.
bool fallbackXCC(ContractionProblemGemm const& problem) const
{
return problem.getParams().fallbackStatus() || isMI300A();
}

WorkgroupMappingXCCCheck()
{
Expand All @@ -3075,7 +3056,6 @@ namespace TensileLite
Hardware const& hardware = *pHardware;
AMDGPU const* pAMDGPU = dynamic_cast<AMDGPU const*>(&hardware);
cuCount = pAMDGPU->computeUnitCount;
processor = pAMDGPU->processor;
}
WorkgroupMappingXCCCheck(std::array<int, 2> value)
: value(value)
Expand All @@ -3085,16 +3065,13 @@ namespace TensileLite
Hardware const& hardware = *pHardware;
AMDGPU const* pAMDGPU = dynamic_cast<AMDGPU const*>(&hardware);
cuCount = pAMDGPU->computeUnitCount;
processor = pAMDGPU->processor;
}

/// Constructor for testing: inject cuCount so selection logic can be
/// unit-tested without a GPU (e.g. ROCM-2963: 38-CU partition alignment).
/// Assumes gfx942, so injecting MI300ACuCount exercises isMI300A().
WorkgroupMappingXCCCheck(std::array<int, 2> value, size_t cuCountForTest)
: value(value)
, cuCount(cuCountForTest)
, processor(AMDGPU::Processor::gfx942)
{
}

Expand All @@ -3113,7 +3090,7 @@ namespace TensileLite
// We overwrite the XCC to 1 to make sure this can pass.
// But we also have to notice we are passing the correct XCC to kernel.
// (i.e. Remember to do param.setWGMXCC(1) when running the kernel)
size_t XCC = fallbackXCC(problem) ? 1 : value[0];
size_t XCC = (problem.getParams().fallbackStatus()) ? 1 : value[0];
size_t XCCG = (value[1] == -1) ? cuCount : value[1];
return ((XCC & (XCC - 1)) == 0) && XCCG % XCC == 0;
}
Expand All @@ -3123,7 +3100,7 @@ namespace TensileLite
{
if(value[0] == -1)
return true;
size_t XCC = fallbackXCC(problem) ? 1 : value[0];
size_t XCC = (problem.getParams().fallbackStatus()) ? 1 : value[0];
size_t XCCG = (value[1] == -1) ? cuCount : value[1];
return debugEvalCmp(problem,
stream,
Expand Down
38 changes: 4 additions & 34 deletions projects/hipblaslt/tensilelite/include/Tensile/PropertyMatching.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1075,45 +1075,15 @@ namespace TensileLite
iter->second.begin(), iter->second.end(), bval, lowerB);
bEnd++;
}
// The 2-D index carries no batch dimension, so a cell is reached on
// (M, N) alone and may hold no row for this problem's batch count --
// b == 1 against a cell whose rows are all batched leaves bStart ==
// bEnd. There is nothing here to return: the cell's solutions are all
// for a different batch size. Skip it and let the remaining cells, or
// the binary search below, resolve the problem.
if(bStart == bEnd)
{
if(Debug)
std::cout << "No entry for batch " << b << " in cell "
<< result.node->pt.coord[0] << ", "
<< result.node->pt.coord[1] << std::endl;
continue;
}

if(Debug)
{
std::cout << "bStart " << bStart->b << ", bEnd ";
if(bEnd == iter->second.end())
std::cout << "<end>";
else
std::cout << bEnd->b;
std::cout << " k at bStart " << bStart->k << ", k at bEnd ";
if(bEnd == iter->second.end())
std::cout << "<end>";
else
std::cout << bEnd->k;
std::cout << std::endl;
}
std::cout << "bStart " << bStart->b << ", bEnd " << bEnd->b
<< " k at bStart " << bStart->k << ", k at bEnd "
<< bEnd->k << std::endl;

auto lower = std::lower_bound(bStart, bEnd, k, compK);
if(lower == bEnd)
lower--;
// Keep the clamped candidate as prev, step lower back one entry, and
// take whichever k is nearer. The step is guarded against bStart so
// it cannot walk before the first entry of the b-group.
auto prev = lower;
if(lower != bStart)
--lower;
auto prev = lower == bStart ? lower : lower--;
if(prev != lower)
{
if(std::abs(prev->k - k) < std::abs(lower->k - k))
Expand Down
Loading