Skip to content

fix(hipBLASLt) : mockup of int32 overflow test and fix #12212

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

Closed
TorreZuk wants to merge 3 commits into from
Closed

fix(hipBLASLt) : mockup of int32 overflow test and fix #12212

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
ae07ecb
add test for int32 overflow seen on gfx12
TorreZuk Sep 16, 2026
100445e
reduce host memory requirements
TorreZuk Sep 16, 2026
2431ae5
fix kernel instructions that fail
TorreZuk Sep 16, 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
10 changes: 10 additions & 0 deletions projects/hipblaslt/clients/common/include/hipblaslt_test.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,16 @@
} \
} while(0)

#define CHECK_HOST_ALLOCATION(BUF, NUM_ELEMENTS) \
do \
{ \
if((NUM_ELEMENTS) > 0 && (BUF).buf() == nullptr) \
{ \
SUCCEED() << LIMITED_MEMORY_STRING; \
return; \
} \
} while(0)

// This wraps the hipBLASLt call with catch_signals_and_exceptions_as_failures().
// By placing it at the hipBLASLt call site, memory resources are less likely to
// be leaked in the event of a caught signal.
Expand Down
29 changes: 23 additions & 6 deletions projects/hipblaslt/clients/common/include/testing_matmul.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2002,6 +2002,11 @@ void testing_matmul_with_bias(const Arguments& arg,
bool do_swizzle_a = arg.swizzle_a && isSwizzleSupported(TiA);
bool do_swizzle_b = arg.swizzle_b && isSwizzleSupported(TiB);
bool mx_use_rocroller = MXUseRocroller();
// Host mirrors of A/B/C are only needed for verification, swizzle, MX init, or dumps.
// Bench runs without --unit_check should not pin multi-GB host buffers.
const bool need_host_ABC = arg.unit_check || arg.norm_check || arg.allclose_check
|| do_swizzle_a || do_swizzle_b || arg.dump_matrix
|| isBlockScaling(arg.scaleA) || isBlockScaling(arg.scaleB);

// Need to split into two for loop to calculate the rotating buffer
int64_t totalRotatingSizeNeeded = 0;
Expand Down Expand Up @@ -2645,9 +2650,15 @@ void testing_matmul_with_bias(const Arguments& arg,
}

// Naming: dX is in GPU (device) memory. hK is in CPU (host) memory
hA.emplace_back(TiA, size_A[i]);
hB.emplace_back(TiB, size_B[i]);
hC.emplace_back(To, size_C[i]);
if(need_host_ABC)
{
hA.emplace_back(TiA, size_A[i]);
CHECK_HOST_ALLOCATION(hA.back(), size_A[i]);
hB.emplace_back(TiB, size_B[i]);
CHECK_HOST_ALLOCATION(hB.back(), size_B[i]);
hC.emplace_back(To, size_C[i]);
CHECK_HOST_ALLOCATION(hC.back(), size_C[i]);
}
hD_gold.emplace_back(To, size_D_copy[i]);
hD_1.emplace_back(To, size_D_copy[i]);
if(size_bias[i] * block_count != 0)
Expand Down Expand Up @@ -2743,9 +2754,15 @@ void testing_matmul_with_bias(const Arguments& arg,
}

// Naming: dX is in GPU (device) memory. hK is in CPU (host) memory
hA.emplace_back(TiA, size_A[i]);
hB.emplace_back(TiB, size_B[i]);
hC.emplace_back(To, size_C[i]);
if(need_host_ABC)
{
hA.emplace_back(TiA, size_A[i]);
CHECK_HOST_ALLOCATION(hA.back(), size_A[i]);
hB.emplace_back(TiB, size_B[i]);
CHECK_HOST_ALLOCATION(hB.back(), size_B[i]);
hC.emplace_back(To, size_C[i]);
CHECK_HOST_ALLOCATION(hC.back(), size_C[i]);
}
hD_gold.emplace_back(To, size_D_copy[i]);
hD_1.emplace_back(To, size_D_copy[i]);
if(size_bias[i] * block_count != 0)
Expand Down
14 changes: 14 additions & 0 deletions projects/hipblaslt/clients/common/include/utility.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -81,6 +81,20 @@
#endif

#define LIMITED_MEMORY_STRING "Error: Attempting to allocate more memory than available."

// Check HipHostBuffer after allocation when host mirrors are required (bench).
// Gtest builds define CHECK_HOST_ALLOCATION in hipblaslt_test.hpp.
#ifndef GOOGLE_TEST
#define CHECK_HOST_ALLOCATION(BUF, NUM_ELEMENTS) \
do \
{ \
if((NUM_ELEMENTS) > 0 && (BUF).buf() == nullptr) \
{ \
hipblaslt_cerr << LIMITED_MEMORY_STRING << std::endl; \
exit(EXIT_FAILURE); \
} \
} while(0)
#endif
#define TOO_MANY_DEVICES_STRING "Error: Too many devices requested."
#define HMM_NOT_SUPPORTED "Error: HMM not supported."
#define KNOWN_BUG_STRING "Known bug for current GPU platform (see known_bugs.yaml)."
Expand Down
3 changes: 3 additions & 0 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 @@ -97,6 +97,9 @@ Definitions:
- { M: 2048, N: 2048, K: 2048, lda: 1049100, ldb: 2048, ldc: 2048, ldd: 2048 }
- { M: 2048, N: 2048, K: 2048, lda: 2048, ldb: 1049100, ldc: 2048, ldd: 2048 }
- { M: 2048, N: 2048, K: 2048, lda: 2048, ldb: 2048, ldc: 1049100, ldd: 1049100 }
# Conv3d-shaped boundary: lda*K = 2^31 elements, max A index = INT32_MAX.
# Requires ~8 GiB device memory for f32 A; host mirrors skipped unless unit_check.
- { M: 67108864, N: 1, K: 32, lda: 67108864, ldb: 32, ldc: 67108864, ldd: 67108864 }

# Output-only oversized leading dimension (C and D): the ROCM-26455 write
# path. K is kept small so the CPU reference / unit_check stays inexpensive.
Expand Down
28 changes: 17 additions & 11 deletions projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@
VAddCCOU32, VAddCOU32, VAddI32, VAddU32, VAndB32, VBfiB32, \
VCmpEQU32, VCmpGtU32, VCmpLtU32, VCmpXNeU32, VCndMaskB32, VLShiftLeftB32, \
VMadI32I24, VMovB32, VMulLOU32, VSubI32, VSubU32
from rocisa.functions import vectorAddMultiplyBpe, vectorMultiply64Bpe
from rocisa.functions import vectorAddMultiplyBpe, vectorAddMultiply64Bpe, vectorMultiply64Bpe
from .Common import INDEX_CHARS, DataDirection, log2

##############################################################################
Expand Down Expand Up @@ -333,11 +333,17 @@ def emitExtractAndScalePackedDims(self, kernel, ss, tmpVgpr, storeChar):
module.add(VAddU32(dst=vgpr(addrVgpr), src0=vgpr(addrVgpr), \
src1=vgpr(tmpVgpr+2), comment="addrCalc += scaled extracted dim "))

module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, addrVgpr, bpe, \
comment="packed: add rowPtr and scaleToBpe"))
self.vectorAddMultiplyStoreBpe(module, kw, addrVgpr, rowPtr, addrVgpr, bpe, tmpVgpr, \
comment="packed: add rowPtr and scaleToBpe")

return module

def vectorAddMultiplyStoreBpe(self, module, kw, addrVgpr, rowPtr, elementVgpr, bpe, tmpVgpr, comment):
if kw.states.use64bShadowLimitStore:
module.add(vectorAddMultiply64Bpe(addrVgpr, rowPtr, elementVgpr, bpe, tmpVgpr + 2, comment=comment))
else:
module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, elementVgpr, bpe, comment=comment))

def emitScaleToBpe(self, kernel, ss, tmpVgpr, tmpSgpr, singleUpdate, tc, dim):
"""
Needs 3 temporary VGPRs
Expand Down Expand Up @@ -473,8 +479,8 @@ def emitScaleToBpe(self, kernel, ss, tmpVgpr, tmpSgpr, singleUpdate, tc, dim):
ss.singleColGateAddrUpdated = True
else:
ss.singleColDAddrUpdated = True
module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, elementVgpr, bpe, \
comment="optSingleColVgpr scaleToBpe: sharedAddrVgpr <- cinRowPtr + coord0, scaled by BPE. BSHERE:coord0=%d, coord0Vgpr=%d"%(kw.vgprs.coord0, self.coord0Vgpr)))
self.vectorAddMultiplyStoreBpe(module, kw, addrVgpr, rowPtr, elementVgpr, bpe, tmpVgpr, \
comment="optSingleColVgpr scaleToBpe: sharedAddrVgpr <- cinRowPtr + coord0, scaled by BPE. BSHERE:coord0=%d, coord0Vgpr=%d"%(kw.vgprs.coord0, self.coord0Vgpr))
elif ss.optSharedColVgpr:
# Need an address calculation for the first address in each row:
if d1==0 and vc1==0:
Expand Down Expand Up @@ -541,8 +547,8 @@ def emitScaleToBpe(self, kernel, ss, tmpVgpr, tmpSgpr, singleUpdate, tc, dim):
module.add(self.emitExtractAndScalePackedDims(kernel, ss, tmpVgpr, tc))
else:
updatedAddr = True
module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, elementVgpr, bpe,
comment="optSharedColVgpr scaleToBpe for first row: col addr <- cinRowPtr + coord0, scaled by BPE"))
self.vectorAddMultiplyStoreBpe(module, kw, addrVgpr, rowPtr, elementVgpr, bpe, tmpVgpr,
comment="optSharedColVgpr scaleToBpe for first row: col addr <- cinRowPtr + coord0, scaled by BPE")
else:
# Generate final address calculation (to bytes) for each element
# The unpacking takes 8-10 instructions so could be worth optimizing someday :
Expand Down Expand Up @@ -611,15 +617,15 @@ def emitScaleToBpe(self, kernel, ss, tmpVgpr, tmpSgpr, singleUpdate, tc, dim):
module.add(self.emitExtractAndScalePackedDims(kernel, ss, tmpVgpr, tc))
else:
updatedAddr = True
module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, elementVgpr, bpe, \
comment="scaleToBpe: accumulate d0 lower and *= bpe into Cin addr"))
self.vectorAddMultiplyStoreBpe(module, kw, addrVgpr, rowPtr, elementVgpr, bpe, tmpVgpr, \
comment="scaleToBpe: accumulate d0 lower and *= bpe into Cin addr")

# if not optSrdIncForRow then we may have moved the row pointer
# and depending on paths above may not have refreshed addrVgpr already.
# if so - do it here:
if self.rowIncDirtyRowPtr and not updatedAddr:
module.add(vectorAddMultiplyBpe(addrVgpr, rowPtr, kw.vgprs.coord0, bpe, \
comment="scaleToBpe: Update address with new rowPtr"))
self.vectorAddMultiplyStoreBpe(module, kw, addrVgpr, rowPtr, kw.vgprs.coord0, bpe, tmpVgpr, \
comment="scaleToBpe: Update address with new rowPtr")

return module

Expand Down
14 changes: 10 additions & 4 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -270,6 +270,7 @@ class StateValues:
groOffsetInMacroTile: int = 0
use64bShadowLimit: bool = True
use64bShadowLimitMX: bool = False
use64bShadowLimitStore: bool = False
preventVgprOverflowDuringNewTile: int = -1
interleaveStoreVmcnt: bool = False
srdShiftLeft:dict = field(init=False)
Expand Down Expand Up @@ -7619,6 +7620,7 @@ def readWriteVectors(mat, vw, kernel):
# use 64-bit buffer limit shadow register
# but not implemented or tested
self.states.use64bShadowLimit = kernel["Use64bShadowLimit"] and kernel["BufferLoad"]
self.states.use64bShadowLimitStore = kernel["Use64bShadowLimit"] and kernel["BufferStore"]
# Keep gfx950's dedicated MX shadow-limit switch, but match legacy gfx1250
# behavior where MX tensors share the same shadow-limit mode as A/B.
if isgfx950:
Expand Down Expand Up @@ -8411,7 +8413,8 @@ def _mxsTileSpanActive(tc):
if kernel["enableTDMA"]:
self.states.a.numVgprGlobalReadOffsets = 0
elif kernel["BufferLoad"]:
self.states.a.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsA * self.states.rpgo)
groVgprMul = 2 if self.states.use64bShadowLimit else 1
self.states.a.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsA * self.states.rpgo) * groVgprMul
else:
numVgprGlobalReadAddressesA = numGlobalReadInstructionsA * self.states.rpga

Expand All @@ -8431,7 +8434,8 @@ def _mxsTileSpanActive(tc):
if kernel["enableTDMA"]:
self.states.mxsa.numVgprGlobalReadOffsets = 0
elif kernel["BufferLoad"]:
self.states.mxsa.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsMXSA * self.states.rpgo)
groVgprMulMX = 2 if self.states.use64bShadowLimitMX else 1
self.states.mxsa.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsMXSA * self.states.rpgo) * groVgprMulMX
else:
numVgprGlobalReadAddressesMXSA = numGlobalReadInstructionsMXSA * self.states.rpga

Expand All @@ -8450,7 +8454,8 @@ def _mxsTileSpanActive(tc):
if kernel["enableTDMB"]:
self.states.b.numVgprGlobalReadOffsets = 0
elif kernel["BufferLoad"]:
self.states.b.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsB * self.states.rpgo)
groVgprMul = 2 if self.states.use64bShadowLimit else 1
self.states.b.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsB * self.states.rpgo) * groVgprMul
else:
numVgprGlobalReadAddressesB = numGlobalReadInstructionsB * self.states.rpga

Expand All @@ -8469,7 +8474,8 @@ def _mxsTileSpanActive(tc):
if kernel["enableTDMB"]:
self.states.mxsb.numVgprGlobalReadOffsets = 0
elif kernel["BufferLoad"]:
self.states.mxsb.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsMXSB * self.states.rpgo)
groVgprMulMX = 2 if self.states.use64bShadowLimitMX else 1
self.states.mxsb.numVgprGlobalReadOffsets = roundUp(numGlobalReadInstructionsMXSB * self.states.rpgo) * groVgprMulMX
else:
numVgprGlobalReadAddressesMXSB = numGlobalReadInstructionsMXSB * self.states.rpga

Expand Down
Loading
Loading