Skip to content
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
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
9 changes: 9 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 was written by the running build. Problems whose rows are dropped fall back to normal kernel selection.

### Removed

Expand Down Expand Up @@ -65,6 +66,14 @@ 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.
* 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.
* `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
5 changes: 5 additions & 0 deletions projects/hipblaslt/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -389,6 +389,11 @@ if(HIPBLASLT_ENABLE_HOST)
if(HIPBLASLT_ENABLE_ROCROLLER)
target_compile_definitions(hipblaslt PUBLIC HIPBLASLT_USE_ROCROLLER)
endif()
if(HIPBLASLT_ENABLE_CLIENT AND (HIPBLASLT_BUILD_TESTING OR BUILD_TESTING))
# Tuning file test hooks, exported only from builds that include the
# client tests that call them.
target_compile_definitions(hipblaslt PRIVATE HIPBLASLT_ENABLE_TUNING_TEST_HOOKS)
endif()
if(NOT HIPBLASLT_ENABLE_YAML)
target_link_libraries(hipblaslt PRIVATE ${hipblaslt_msgpack_target})
endif()
Expand Down
16 changes: 16 additions & 0 deletions projects/hipblaslt/clients/common/include/argument_model.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,17 @@ void ArgumentModel_log_efficiency(hipblaslt_internal_ostream& name_line,
const Arguments& arg,
const double hipblaslt_gflops);

// The ext API's name accessors append " (Custom tuning: GSU: x, WGM: y)" when
// GemmTuning overrides differ from the solution's own values. Replay compares
// against getKernelNameFromAlgoIndex, which never appends it, so a persisted
// kernel_name must not carry it.
inline std::string hipblaslt_strip_custom_tuning_suffix(const std::string& name)
{
static const std::string marker = " (Custom tuning: ";
const size_t pos = name.find(marker);
return pos == std::string::npos ? name : name.substr(0, pos);
}

// ArgumentModel template has a variadic list of argument enums
template <hipblaslt_argument... Args>
class ArgumentModel
Expand Down Expand Up @@ -313,6 +324,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 << hipblaslt_strip_custom_tuning_suffix(kernel_name);

const char* tuningEnv = getenv("HIPBLASLT_TUNING_FILE");
std::string tuningPath = tuningEnv;
std::ofstream file(tuningPath, std::ios::app);
Expand Down
14 changes: 12 additions & 2 deletions projects/hipblaslt/clients/common/include/testing_matmul.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -6163,7 +6163,9 @@ void testing_matmul_with_bias(const Arguments& arg,

if(arg.print_solution_found)
{
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 +6179,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 @@ -6242,7 +6251,8 @@ void testing_matmul_with_bias(const Arguments& arg,
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 @@ -18,6 +18,7 @@ target_sources(hipblaslt-test
${CMAKE_CURRENT_SOURCE_DIR}/streamk_multistream_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
)

Expand Down
Loading
Loading