Using hipBLASLt offline tuning

Using hipBLASLt offline tuning#

hipblaslt-bench can help find the best-performing GEMM kernel for a given set of GEMM problems and provide the best solution index for a given problem size. This index can be used directly in future GEMM calls through the User Offline Tuning mechanism. However, these indices cannot be reused across library releases or across different device architectures.

Use the command line interface to access this functionality. See hipBLASLt clients for more details.

Using hipblaslt-bench to run the tuning with the best GEMM kernel#

To find and use the best GEMM kernel for a problem, follow these steps:

  1. Generate the tuning command line by setting the environment variable HIPBLASLT_LOG_MASK=32 before calling any hipBLASLt APIs. For more details on how to use hipblaslt-bench, see Logging and heuristics. In the Bash shell, set the following environment variable:

    export HIPBLASLT_LOG_MASK=32

  2. The following command uses sample_hipblaslt_gemm.cpp as an example:

    ./sample_hipblaslt_gemm

    The tuning command displays the following log entry:

    hipblaslt-bench --api_method c -m 1024 -n 512 -k 1024 --lda 1024 --ldb 1024 --ldc 1024 --ldd 1024  --stride_a 0 --stride_b 0 --stride_c 0 --stride_d 0  --alpha 1.000000 --beta 1.000000 --transA N --transB N --batch_count 1  --a_type f16_r --b_type f16_r --c_type f16_r --d_type f16_r --scale_type f32_r --bias_type f32_r   --compute_type f32_r --algo_method index --solution_index 56073

  3. Set the environment variable HIPBLASLT_TUNING_FILE=<file_name> to tune and store the tuning result, which indicates the best solution indices for the GEMM problems. The <file_name> points to the tuning file.

    In the Bash shell, set the following environment variable:

    export HIPBLASLT_TUNING_FILE=tuning.txt

    Additionally, you can set the environment variable to specify that the solution found in the tuning stage is under the constraint of the max workspace size setting:

    export HIPBLASLT_TUNING_USER_MAX_WORKSPACE=<value> (Default value is: 128 * 1024 * 1024)

    The default settings for the following parameters in hipblaslt-bench are changed in the tuning environment.

    --iters |-i <value>             (Default value is: 1000)
--cold_iters |-j <value>        (Default value is: 1000)
--requested_solution <value>    (Default value is: -1)
--rotating <value>              (Default value is: 512)

    After the tuning completes, the expected output is displayed as follows:

    ./hipblaslt-bench --api_method c -m 1024 -n 512 -k 1024 --lda 1024 --ldb 1024 --ldc 1024 --ldd 1024  --stride_a 0 --stride_b 0 --stride_c 0 --stride_d 0  --alpha 1.000000 --beta 1.000000 --transA N --transB N --batch_count 1  --a_type f16_r --b_type f16_r --c_type f16_r --d_type f16_r --scale_type f32_r --bias_type f32_r   --compute_type f32_r --algo_method index --solution_index 56073

Winner:
transA,transB,grouped_gemm,batch_count,m,n,k,alpha,lda,stride_a,beta,ldb,stride_b,ldc,stride_c,ldd,stride_d,a_type,b_type,c_type,d_type,compute_type,scaleA,scaleB,scaleC,scaleD,amaxD,activation_type,bias_vector,bias_type,rotating_buffer,hipblaslt-Gflops,hipblaslt-GB/s,us,solution_index,kernel_name
N,N,0,1,1024,512,1024,1,1024,1048576,1,1024,524288,1024,524288,1024,524288,f16_r,f16_r,f16_r,f16_r,f32_r,0,0,0,0,0,none,0,f32_r,512,66613.8,363.509,16.1189,56537,<kernel_name>

  4. Set the environment variable HIPBLASLT_TUNING_OVERRIDE_FILE=<file_name> to load the tuning file and override the default kernel selection with the optimal kernel choices, where <file_name> points to the tuning file.

    In the Bash shell, set the following environment variable:

    export HIPBLASLT_TUNING_OVERRIDE_FILE=tuning.txt

    For example, you can use hipblaslt-bench with algo_method set to heuristic to obtain the solutions for a problem, which include the best tuning solution index.

    ./hipblaslt-bench --api_method c -m 1024 -n 512 -k 1024 --lda 1024 --ldb 1024 --ldc 1024 --ldd 1024  --stride_a 0 --stride_b 0 --stride_c 0 --stride_d 0  --alpha 1.000000 --beta 1.000000 --transA N --transB N --batch_count 1  --a_type f16_r --b_type f16_r --c_type f16_r --d_type f16_r --scale_type f32_r --bias_type f32_r   --compute_type f32_r --algo_method heuristic --requested_solution 1 --print_kernel_info

transA,transB,grouped_gemm,batch_count,m,n,k,alpha,lda,stride_a,beta,ldb,stride_b,ldc,stride_c,ldd,stride_d,a_type,b_type,c_type,d_type,compute_type,scaleA,scaleB,scaleC,scaleD,amaxD,activation_type,bias_vector,bias_type,rotating_buffer,hipblaslt-Gflops,hipblaslt-GB/s,us,solution_index,kernel_name
[0]:
N,N,0,1,1024,512,1024,1,1024,1048576,1,1024,524288,1024,524288,1024,524288,f16_r,f16_r,f16_r,f16_r,f32_r,0,0,0,0,0,none,0,f32_r,512,37575.2,205.047,28.5758,56537,<kernel_name>

How tuned entries are validated#

A tuning file records a kernel_name next to each solution_index. When the file is replayed, the index is treated as a lookup hint: hipBLASLt resolves it in the current library and uses it only if it still names the same kernel. If it does not, that entry is rejected and the problem falls back to normal heuristic selection.

Validation is per entry, so upgrading hipBLASLt thins a tuning file rather than discarding it, and only the shapes that failed need re-tuning. A row that carries a solution_name instead is validated the same way. Files written by an older hipBLASLt carry no name at all, and those are used only when the file’s Git Version line matches the running build. A build made outside a git checkout has no version to match, so it uses none of them.

kernel_name names the compiled kernel, not the whole solution. The index still selects the solution, including its GSU, WGM and StaggerU defaults, and several solutions can share one compiled kernel, which is why the name is paired with the index rather than used on its own. What validation guarantees is that the index resolves to the same compiled kernel as when the row was written. It does not detect a rebuild that changes only that solution’s defaults, whose row is used with the new defaults, and it cannot find a kernel that moved to a different index: that row is rejected and the shape falls back to heuristic selection.

Replaying a tuning cache#

Instead of HIPBLASLT_TUNING_OVERRIDE_FILE, a tuning file can be replayed as a runtime cache:

export HIPBLASLT_TUNING_MODE=cache
export HIPBLASLT_TUNING_CACHE_PATH=tuning.txt

cache replays what the file contains, validating each entry as described above. off, the default, changes nothing. Cache mode only reads the file, so any number of processes can share one, for example the ranks of an MPI job.

Which kernel a matmul launches depends on whether it passes an algorithm. A call to hipblasLtMatmul with an explicit algo always launches that algorithm; such a caller gets the cached entry from hipblasLtMatmulAlgoGetHeuristic, which returns it as its first result. A call with algo=nullptr leaves the choice to the library and launches the cached entry when there is a usable one, which the override file never did.

Rows written by hipblaslt-bench match on transpose, shape and the principal datatypes only, exactly as they do through HIPBLASLT_TUNING_OVERRIDE_FILE: they do not distinguish leading dimensions, batch strides, epilogue details or the device. Rows that carry a schema_version column are keyed on the complete problem, and every column of such a row is required, so a row with a missing or malformed value, including one cut short by an interrupted write, is ignored.

HIPBLASLT_TUNING_CACHE_PATH and HIPBLASLT_TUNING_OVERRIDE_FILE are mutually exclusive. With a tuning mode set, only the cache is consulted and the override file is ignored.

A process running in a secure execution context (set-user-ID, set-group-ID, or another credential-changing exec such as file capabilities) ignores both tuning variables, so the cache stays off in it.

Reading the cache output#

cache mode reports what it loaded without needing any logging variable:

tuning-cache: mode=cache path=tuning.txt load=ok loaded=12
tuning-cache: summary shapes=14 matched=12 fellback=2 invalidated=0

The first line is written when the file is first read. load is ok, not-found, read-error or no-path, the last meaning HIPBLASLT_TUNING_CACHE_PATH is unset and the cache does nothing.

The summary, written as the process exits normally, is how you confirm a deployed cache is being used. shapes counts distinct problems the process looked up, not calls, so it stays comparable with loaded. matched is how many of them a cache entry served and fellback how many fell through to default selection; those are the shapes still worth tuning. invalidated counts entries rejected because a rebuild moved the kernel they named. A process that is killed or aborts writes no summary, and one that loaded nothing and looked nothing up leaves it out.

Setting HIPBLASLT_LOG_LEVEL=4 adds a cache hit, miss or invalidation line once per problem. See hipBLASLt environment variables for where each notice is written.

Limitations#

  • Grouped GEMM is not served from the cache, and neither is a pointer-array batch called with algo=nullptr.

  • alpha, beta and whether C and D alias are not part of the lookup key, so an entry can serve a caller whose values differ from those it was recorded with.

Contents