Skip to content
Merged
11 changes: 11 additions & 0 deletions projects/hipblaslt/CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip
* Stream-K flags are now per-stream: a handle reserves an extra fixed 8 MiB at creation and serves at most 64 distinct streams for Stream-K matmuls (claimed on a stream's first use, held until the handle is destroyed); beyond that the matmul returns `HIPBLAS_STATUS_INTERNAL_ERROR`.
* Stream-K workspace size reported by the heuristic APIs is now smaller, and the SK grid is bounded, so `TENSILE_STREAMK_GRID_MULTIPLIER` values past that bound no longer take effect.
* Solution cache key now includes `HIPBLASLT_MATMUL_DESC_SM_COUNT_TARGET` and the StreamK tile scheduling mode, so the same problem can select a different kernel than before.
* A tuning file is now trusted one entry at a time instead of all or nothing. `HIPBLASLT_TUNING_OVERRIDE_FILE` records solution indices, which are positions in one build's kernel library. Previously the C API ignored the whole file when its build-version line did not match the running build, while the C++ API applied it regardless and could run kernels it was never tuned on. `hipblaslt-bench` now records a `kernel_name` beside each `solution_index` in `HIPBLASLT_TUNING_FILE`, and on both APIs a row that records a name is checked at replay and dropped only if its index no longer names that kernel. A row without a name is used only when the file's `Git Version` line matches the running build; a build made outside a git checkout has no version, so it uses no such rows. Problems whose rows are dropped fall back to normal kernel selection.

### Removed

Expand Down Expand Up @@ -65,6 +66,16 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip
* Corrected `PhysicalMaxVgprCU` for gfx1102 and gfx1103, where gfx1103 was reported with a 1536-VGPR per-SIMD file instead of 1024 and gfx1102 omitted the two-SIMDs-per-CU factor, both affecting occupancy and kernel selection.
* Fixed a StreamK per-XCD work-queue counter that failed to reset between launches when the StreamK grid size was not a multiple of the XCD count, causing progressively slower execution on repeated GEMM launches (`WorkGroupMappingXCC == -1` and `StreamKXCCMapping` chiplet-remap paths).
* Fixed out-of-bounds stores in subtile GEMM kernels on gfx950 and gfx1250 when the `M` dimension does not evenly fill the macro tile (for example `M=8` with a 32-row tile).
* When a tuning file held several entries for one problem and the first could not be used, the later entries were never used either. The replay loop reused a single vector across entries and always read element `[0]`, so each later entry was checked against the first entry's solution instead of its own.
* Solution index `0` in a tuning file was rejected outright, despite being a valid index that is present in the shipped logic.
* A grouped GEMM heuristic query crashed when `HIPBLASLT_TUNING_OVERRIDE_FILE` held any entry, because the override lookup read the grouped problem as a single GEMM. Grouped GEMMs now use default selection.
* `hipblaslt-bench` recorded grouped GEMM winners in `HIPBLASLT_TUNING_FILE`, where a single-group winner read back as a single GEMM of the same shape. Grouped runs no longer write to the file.
* An XF32 problem on the C++ extension API could fall back to FP32 kernels. When a tuning file entry failed both the XF32 check and its FP32 fallback check, the problem was left in FP32 mode for the remaining entries and for default kernel selection.
* `*returnAlgoCount` was read uninitialised when an override satisfied a single-algo request, and then used to scan one element past the end of the caller's array.
* A tuning file that yielded no usable rows was re-read and re-parsed on every heuristic query instead of once.
* A tuning file row cut short, for example when the disk fills while `hipblaslt-bench` appends it, could be replayed with a truncated `solution_index`. A row whose cells do not line up with its header is now ignored.
* `GemmInstance::getSolutionName()` in `hipblaslt_ext` crashed for a RocRoller solution. It looked the encoded RocRoller index up in the Tensile library and dereferenced the null result. It now returns the RocRoller short name, as `hipblaslt_ext::getSolutionNameFromAlgo()` already did.
* The logger's destructor called `close()` on the log file, which throws out of a destructor and terminates the process at exit when the log file has failed, for example on a full disk. Reproducible with `HIPBLASLT_LOG_LEVEL=4` and a failing `HIPBLASLT_LOG_FILE`.

## hipBLASLt 1.4.1 for ROCm 7.14

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -313,6 +313,11 @@ class ArgumentModel
name_list << delim << "solution_index";
value_list << delim << solution_index;

// Replay resolves solution_index and uses it only if it still
// names this kernel.
name_list << delim << "kernel_name";
value_list << delim << kernel_name;

const char* tuningEnv = getenv("HIPBLASLT_TUNING_FILE");
std::string tuningPath = tuningEnv;
std::ofstream file(tuningPath, std::ios::app);
Expand Down
22 changes: 17 additions & 5 deletions projects/hipblaslt/clients/common/include/testing_matmul.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -6155,15 +6155,19 @@ void testing_matmul_with_bias(const Arguments& arg,
std::string archName = "";
std::string cuNum = "";

if(tuningEnv && heuristicResult.size() == 1)
// The tuning file has no grouped GEMM key, so a grouped winner would
// be read back as a single GEMM of the same shape.
if(tuningEnv && !do_grouped_gemm && heuristicResult.size() == 1)
{
archName = deviceProps.gcnArchName;
cuNum = std::to_string(deviceProps.multiProcessorCount);
}

if(arg.print_solution_found)
if(arg.print_solution_found || tuningEnv)
{
if(arg.print_kernel_info)
// A tuning run records the winner's kernel name, which comes
// from best_k_name below.
if(arg.print_kernel_info || tuningEnv)
Comment thread
jinchen62 marked this conversation as resolved.
Resolved
{
if(arg.use_ext && batchMode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY)
{
Expand All @@ -6177,6 +6181,13 @@ void testing_matmul_with_bias(const Arguments& arg,
solutionName = groupedGemmVec[0].getSolutionName();
kernelName = groupedGemmVec[0].getKernelName();
}

// The ext accessor joins every kernel the solution
// launches with "; ". Replay compares against the main
// kernel's name alone, which the algo accessor returns.
if(tuningEnv)
kernelName = hipblaslt_ext::getKernelNameFromAlgo(
handle, heuristicResult[sol].algo);
}
else
{
Expand Down Expand Up @@ -6236,13 +6247,14 @@ void testing_matmul_with_bias(const Arguments& arg,
std::string kernelName = "";
std::string archName = "";
std::string cuNum = "";
if(tuningEnv)
if(tuningEnv && !do_grouped_gemm)
{
archName = deviceProps.gcnArchName;
cuNum = std::to_string(deviceProps.multiProcessorCount);
}

if(arg.print_kernel_info)
// Same reason as the per-candidate loop above.
if(arg.print_kernel_info || tuningEnv)
{
solutionName = best_s_name;
kernelName = best_k_name;
Expand Down
1 change: 1 addition & 0 deletions projects/hipblaslt/clients/tests/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,7 @@ target_sources(hipblaslt-test
${CMAKE_CURRENT_SOURCE_DIR}/streamk_null_stream_test.cpp
${CMAKE_CURRENT_SOURCE_DIR}/synchronizer_size_check_test.cpp
${CMAKE_CURRENT_SOURCE_DIR}/row_uniformity_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/tuning_cache_test.cpp
${CMAKE_CURRENT_SOURCE_DIR}/ulp_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/arguments_a2a_gtest.cpp
${CMAKE_CURRENT_SOURCE_DIR}/unit_gtest.cpp
Expand Down
Loading
Loading