### Build and install hipBLASLt package with client Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds all dependencies, the library, and the client, then installs the hipBLASLt package. Requires sudo access. Use `-i` to install for all users; omit it to keep the build local. ```bash ./install.sh -idc ``` -------------------------------- ### Grouped GEMM Pseudo-code Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Demonstrates the workflow for setting up and executing a grouped GEMM operation using hipBLASLt extension APIs, including algorithm selection, problem setup, and default UserArguments retrieval. ```c++ // Pseudo code // Step 1: Get all algorithms std::vector heuristicResult; CHECK_HIPBLASLT_ERROR(hipblaslt_ext::getAllAlgos(handle, HIPBLASLT_GEMM, HIPBLAS_OP_N, HIPBLAS_OP_N, in_out_datatype, in_out_datatype, in_out_datatype, in_out_datatype, HIPBLAS_COMPUTE_32F, heuristicResult)); hipblaslt_ext::GemmPreference pref; pref.setMaxWorkspaceBytes(1000000); // Step 2: Setup problem std::vector m(gemm_count); std::vector n(gemm_count); std::vector k(gemm_count); std::vector batch_count(gemm_count); std::vector epilogue(gemm_count); std::vector inputs(gemm_count); for(int i = 0; i < gemm_count; i++) { m[i] = 1; n[i] = 1; k[i] = 1; batch_count[i] = 1; epilogue[i].setMode(HIPBLASLT_EPILOGUE_GELU); inputs[i].setA(d_a[i]); inputs[i].setB(d_b[i]); inputs[i].setC(d_c[i]); inputs[i].setD(d_d[i]); inputs[i].setAlpha(&alpha[i]); inputs[i].setBeta(&beta[i]); } // Step 3: Create grouped gemm instance hipblaslt_ext::GroupedGemm groupedGemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIPBLAS_COMPUTE_32F); // Step 4: Set problem groupedGemm.setProblem(m, n, k, batch_count, epilogue, inputs); // m, n, k, batch // Step 5: Get default value from the instance hipblaslt_ext::UserArguments* dUAFloat = new hipblaslt_ext::UserArguments[gemm_count]; groupedGemm.getDefaultValueForDeviceUserArguments((void*)dUAFloat); // Once you get the default value here, you can make several copies and change the values // from the host // Next copy them to the device memory hipblaslt_ext::UserArguments* d_dUAFloat = nullptr; hipMalloc(&d_dUAFloat, sizeof(hipblaslt_ext::UserArguments) * gemm_count); ``` -------------------------------- ### Build and install hipBLASLt package and client Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds and installs the hipBLASLt package and builds the client, requiring sudo access. Use `-i` for system-wide installation; omit it to keep the build local. ```bash ./install.sh -ic ``` -------------------------------- ### Install hipBLASLt Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-tuning-utility.html Run the install script to set up hipBLASLt before using the tuning utility. Ensure `MatchTable.yaml` exists in the build directory. ```bash ./install.sh ``` -------------------------------- ### Build and install hipBLASLt package Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds the hipBLASLt library and installs the package to `/opt/rocm/hipblaslt`. This command requires sudo access and installs for all users. Omit `-i` to keep the library in the local directory. ```bash ./install.sh -i ``` -------------------------------- ### hipblaslt-bench Tuning Command Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-offline-tuning.html This is an example of a tuning command generated by hipBLASLt, used to find the best GEMM kernel for specific problem dimensions and data types. It specifies the algorithm method as 'index'. ```bash ./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 ``` -------------------------------- ### hipBLASLt install.sh script help Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Displays help information for the `install.sh` script, which is used to build and install hipBLASLt and its dependencies. This is useful for understanding available options. ```bash ./install.sh -h ``` -------------------------------- ### Install hipBLASLt using apt Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Installs the hipBLASLt package using the apt package manager. Ensure your system's package list is up-to-date before running this command. ```bash sudo apt update && sudo apt install hipblaslt ``` -------------------------------- ### Run hipblaslt-bench Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-offline-tuning.html Execute the sample hipBLASLt GEMM application to generate tuning-related log entries. This command is typically run after setting the logging mask. ```bash ./sample_hipblaslt_gemm ``` -------------------------------- ### hipblaslt-bench with Heuristic Algorithm Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-offline-tuning.html This example demonstrates using hipblaslt-bench with the 'heuristic' algorithm method to obtain kernel solutions, including the best tuning solution index. The 'requested_solution' parameter is used to specify which solution to retrieve. ```bash ./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 ``` -------------------------------- ### Grouped GEMM Setup Pseudo-code Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Illustrates the typical setup for a grouped GEMM operation using `hipblasLtExt`. It shows how to populate vectors for problem dimensions, strides, epilogues, inputs, and problem types before calling the `setProblem` API. Note that only a single `GemmProblemType` is supported for all problems in the group. ```c++ // Pseudo code std::vector m, n, k; // ... for(size_t i = 0; i < problem_size, i++) { // ... } std::vector problemtypes; problemtypes.push_back(problemtype); groupedgemm.setProblem(m, n, k, batch_count, lda, ldb, ldc, ldd, strideA, strideB, strideC, strideD, epilogue, inputs, problemtypes); ``` -------------------------------- ### Get Algorithm Index and Heuristic Results Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Demonstrates how to retrieve an algorithm index from test results and then use that index to get heuristic results. Handles cases where the index might be out of bounds. ```cpp int index = hipblaslt_ext::getIndexFromAlgo(testResults[i].algo); // Save the index to disk or somewhere else for later use. // Get the index from previous state. std::vector algoIndex{index}; std::vector heuristicResults; // If the index is out of the bound of solutions, getAlgosFromIndex will return HIPBLAS_STATUS_INVALID_VALUE if(HIPBLAS_STATUS_INVALID_VALUE == hipblaslt_ext::getAlgosFromIndex(handle, algoIndex, heuristicResults)) { std::cout << "Indexes are all out of bound." << std::endl; break; } ``` -------------------------------- ### run (stream, start, stop) Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Execute the kernel arguments stored inside the hipblaslt_ext::GemmInstance. This function executes the kernel with optional HIP events for timing. ```APIDOC ## run (stream, start, stop) ### Description Execute the kernel arguments stored inside the hipblaslt_ext::GemmInstance. ### Parameters * **stream** - The HIP stream where all the GPU work will be submitted. * **start** - The HIP event which will record the start of the kernel. * **stop** - The HIP event which will record the end of the kernel. ### Return Values * **HIPBLAS_STATUS_SUCCESS** - If the operation completed successfully. ``` -------------------------------- ### Build hipBLASLt Clients with CMake Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/conceptual/hipblaslt-clients.html Use these CMake options to build only the clients, samples, tests, and benchmarks without building Tensile libraries. This is an alternative to using the install script with the `-n` flag. ```bash mkdir build && cd build cmake [other options] -DBUILD_CLIENTS_SAMPLES=ON -DBUILD_CLIENTS_TESTS=ON -DBUILD_CLIENTS_BENCHMARKS=ON -DTensile_SKIP_BUILD=ON .. make ``` -------------------------------- ### Build hipBLASLt Dependencies Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds optional dependencies for hipBLASLt clients, such as GoogleTest and LAPACK, from source. Assumes Boost is installed via a package manager. Installs to the default CMake directory (/usr/local). ```bash mkdir -p [HIPBLASLT_BUILD_DIR]/release/deps cd [HIPBLASLT_BUILD_DIR]/release/deps ccmake -DBUILD_BOOST=OFF [HIPBLASLT_SOURCE]/deps # assuming boost is installed through package manager as above make -j$(nproc) install ``` -------------------------------- ### GemmInputs Class Methods Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Methods for setting and getting input matrices, scaling factors, and bias for GEMM operations. ```APIDOC ## GemmInputs ### Description Manages the input data for GEMM operations, including matrices A, B, C, D, scaling factors (alpha, beta), bias, and auxiliary data. ### Methods - `setA()` - `setB()` - `setC()` - `setD()` - `setAlpha()` - `setBeta()` - `setBias()` - `setScaleA()` - `setScaleB()` - `setScaleC()` - `setScaleD()` - `setScaleAux()` - `setScaleAlphaVec()` - `setAux()` - `setAmaxD()` - `getA()` - `getB()` - `getC()` - `getD()` - `getAlpha()` - `getBeta()` - `getBias()` - `getScaleA()` - `getScaleB()` - `getScaleC()` - `getScaleD()` - `getScaleAux()` - `getScaleAlphaVec()` - `getAux()` - `getAmaxD()` ``` -------------------------------- ### Build hipBLASLt Library Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds the hipBLASLt library. Specify CMAKE_INSTALL_PREFIX for custom install locations and CMAKE_BUILD_TYPE for build configurations. Uses amdclang++ as the C++ compiler. ```bash mkdir -p [HIPBLASLT_BUILD_DIR]/release cd [HIPBLASLT_BUILD_DIR]/release # Default install location is in /opt/rocm, define -DCMAKE_INSTALL_PREFIX= to specify other # Default build config is 'Release', define -DCMAKE_BUILD_TYPE= to specify other CXX=/opt/rocm/bin/amdclang++ ccmake [HIPBLASLT_SOURCE] make -j$(nproc) sudo make install # sudo required if installing into system directory such as /opt/rocm ``` -------------------------------- ### Grouped GEMM Pseudo Code Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Illustrates how to set up and call the `setProblem` function for grouped GEMM operations, emphasizing the use of vectors for problem parameters and the `GemmProblemType`. ```APIDOC // Pseudo code std::vector m, n, k; // ... populate m, n, k vectors ... std::vector batch_count; // ... populate batch_count vector ... std::vector lda, ldb, ldc, ldd; // ... populate stride vectors ... std::vector strideA, strideB, strideC, strideD; // ... populate stride vectors ... std::vector epilogue; // ... populate epilogue vector ... std::vector inputs; // ... populate inputs vector ... std::vector problemtypes; GemmProblemType problemtype; // ... configure problemtype ... problemtypes.push_back(problemtype); // Assuming 'groupedgemm' is an instance of the class that has the setProblem method groupedgemm.setProblem(m, n, k, batch_count, lda, ldb, ldc, ldd, strideA, strideB, strideC, strideD, epilogue, inputs, problemtypes); ``` -------------------------------- ### Initialize and Run Grouped GEMM Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Demonstrates the process of setting up a grouped GEMM problem, including algorithm selection, problem definition, and execution. It shows how to retrieve default user arguments and update them before running the kernel. ```c++ // Pseudo code // Step 1: Get all algorithms std::vector heuristicResult; CHECK_HIPBLASLT_ERROR(hipblaslt_ext::getAllAlgos(handle, HIPBLASLT_GEMM, HIPBLAS_OP_N, HIPBLAS_OP_N, in_out_datatype, in_out_datatype, in_out_datatype, in_out_datatype, HIPBLAS_COMPUTE_32F, heuristicResult)); hipblaslt_ext::GemmPreference pref; pref.setMaxWorkspaceBytes(1000000); // Step 2: Setup problem std::vector m(gemm_count); std::vector n(gemm_count); std::vector k(gemm_count); std::vector batch_count(gemm_count); std::vector epilogue(gemm_count); std::vector inputs(gemm_count); // Step 2.1: Calculate sum of n int64_t sum_of_n = 0; for(int i = 0; i < gemm_count; i++) { sum_of_n += n_arr[i]; } // {sum_of_n, 1, 1, 1, ...}; // The array of N, the first element is the sum of N for(int i = 0; i < gemm_count; i++) { m[i] = m_arr[i]; if(i == 0) n[i] = sum_of_n; else n[i] = 1; k[i] = k_arr[i]; batch_count[i] = 1; inputs[i].setA(d_a[i]); inputs[i].setB(d_b[i]); inputs[i].setC(d_c[i]); inputs[i].setD(d_d[i]); inputs[i].setAlpha(&alpha[i]); inputs[i].setBeta(&beta[i]); } // Step 3: Create grouped gemm instance hipblaslt_ext::GroupedGemm groupedGemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIPBLAS_COMPUTE_32F); // Step 4: Set problem groupedGemm.setProblem(m, n, k, batch_count, epilogue, inputs); // m, n, k, batch // Step 5: Get default value from the instance hipblaslt_ext::UserArguments* dUAFloat = new hipblaslt_ext::UserArguments[gemm_count]; groupedGemm.getDefaultValueForDeviceUserArguments((void*)dUAFloat); // Once you get the default value here, you can make several copies and change the values // from the host // Next Copy them to the device memory hipblaslt_ext::UserArguments* d_dUAFloat = nullptr; hipMalloc(&d_dUAFloat, sizeof(hipblaslt_ext::UserArguments) * gemm_count); hipMemcpy(d_dUAFloat, dUAFloat, sizeof(hipblaslt_ext::UserArguments) * gemm_count, hipMemcpyHostToDevice); validIdx.clear(); for(int j = 0; j < heuristicResult.size(); j++) { size_t workspace_size = 0; if(groupedGemm.isAlgoSupported(heuristicResult[j].algo, workspace_size) == HIPBLAS_STATUS_SUCCESS) { validIdx.push_back(j); } } int threads = 256; int blocks = ceil((double)gemm_count / threads); // Step 6: Initialize and run if(validIdx.size() > 1) { groupedGemm.initialize(heuristicResult[validIdx[0]].algo, d_workspace); for(int i = 0; i < 10; i++) { hipLaunchKernelGGL(kernelUpdateN, dim3(blocks), dim3(threads), 0, stream, gemm_count, d_dUAFloat, d_n_vec); // d_n_vec is a device pointer with Ns groupedGemm.run(userArgs, stream); } } // ..... __global__ void kernelUpdateN(uint32_t gemm_count, void* userArgs, int32_t* sizes_n) { uint64_t id = hipBlockIdx_x * 256 + hipThreadIdx_x; if(id >= gemm_count) return; hipblaslt_ext::UserArguments* dUAFloat = static_cast(userArgs); dUAFloat[id].n = sizes_n[id]; } ``` -------------------------------- ### hipBLASLt Bench Help Options Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/conceptual/hipblaslt-clients.html Display the help message for the hipblaslt-bench command-line tool. This output lists all available options for configuring matrix sizes, data types, precision, and execution parameters. ```bash ./hipblaslt-bench --help ``` -------------------------------- ### hipBLASLt AMax Operation Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/samples/sample_hipblaslt_ext_op_amax.html This C++ code demonstrates the usage of hipblasltExtAMax for finding the absolute maximum value of a 2D tensor. It includes examples for both float and half-precision float data types. Ensure hipBLASLt and HIP runtime are correctly set up. ```cpp /******************************************************************************* * * MIT License * * Copyright (C) 2022-2024 Advanced Micro Devices, Inc. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal * in the Software without restriction, including without limitation the rights * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell * copies of the Software, and to permit persons to whom the Software is * furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice shall be included in * all copies or substantial portions of the Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. * *******************************************************************************/ #include #include #include #include #include "helper.h" void simpleAMax(hipDataType type, hipDataType dtype, void* d_out, void* d_in, int64_t m, int64_t n, hipStream_t stream); int main() { /** This is a amax example * in = (m, n). lda = m * out = (1). ldb = 1 */ OptAMaxRunner runnerF32(135, 345); runnerF32.run([&runnerF32] { simpleAMax(HIP_R_32F, HIP_R_32F, runnerF32.d_out, runnerF32.d_in, runnerF32.m, runnerF32.n, runnerF32.stream); }); OptAMaxRunner runnerF16(135, 345); runnerF16.run([&runnerF16] { simpleAMax(HIP_R_16F, HIP_R_16F, runnerF16.d_out, runnerF16.d_in, runnerF16.m, runnerF16.n, runnerF16.stream); }); return 0; } void simpleAMax(hipDataType type, hipDataType dtype, void* d_out, void* d_in, int64_t m, int64_t n, hipStream_t stream) { CHECK_HIPBLASLT_ERROR(hipblasltExtAMax(type, dtype, d_out, d_in, m, n, stream)); } ``` -------------------------------- ### LayerNorm Extension Operation Example Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/samples/sample_hipblaslt_ext_op_layernorm.html This C++ code demonstrates the usage of the hipblasltExtLayerNorm function to perform Layer Normalization on a 2D tensor. It includes necessary headers, a helper function for the operation, and a main function to run the example with float data type. ```cpp /******************************************************************************* * * MIT License * * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal * in the Software without restriction, including without limitation the rights * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell * copies of the Software, and to permit persons to whom the Software is * furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice shall be included in * all copies or substantial portions of the Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. * *******************************************************************************/ #include #include #include #include #include #include "helper.h" void simpleLayerNorm(hipDataType type, void* d_out, void* d_mean, void* d_invvar, void* d_in, int64_t m, int64_t n, float eps, void* d_gamma, void* d_beta, hipStream_t stream); int main() { LayerNormRunner runnerF32(135, 345); runnerF32.run([&runnerF32] { simpleLayerNorm(HIP_R_32F, runnerF32.d_out, runnerF32.d_mean, runnerF32.d_invvar, runnerF32.d_in, runnerF32.m, runnerF32.n, 1e-5, runnerF32.d_gamma, runnerF32.d_beta, runnerF32.stream); }); return 0; } void simpleLayerNorm(hipDataType type, void* d_out, void* d_mean, void* d_invvar, void* d_in, int64_t m, int64_t n, float eps, void* d_gamma, void* d_beta, hipStream_t stream) { CHECK_HIPBLASLT_ERROR(hipblasltExtLayerNorm( type, d_out, d_mean, d_invvar, d_in, m, n, eps, d_gamma, d_beta, stream)); } ``` -------------------------------- ### Initialize and Run Grouped GEMM Operations Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html This snippet demonstrates initializing grouped GEMM operations with a selected algorithm and then running them multiple times. It requires a valid list of supported algorithms and user arguments. ```cpp hipMemcpy(d_dUAFloat, dUAFloat, sizeof(hipblaslt_ext::UserArguments) * gemm_count, hipMemcpyHostToDevice); validIdx.clear(); for(int j = 0; j < heuristicResult.size(); j++) { size_t workspace_size = 0; if(groupedGemm.isAlgoSupported(heuristicResult[j].algo, workspace_size) == HIPBLAS_STATUS_SUCCESS) { validIdx.push_back(j); } } // Step 6: Initialize and run if(validIdx.size() > 1) { groupedGemm.initialize(heuristicResult[validIdx[0]].algo, d_workspace, stream); for(int i = 0; i < 10; i++) { groupedGemm.run(userArgs, stream); } } ``` -------------------------------- ### Tuning Utility Logic Creation Progress Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-tuning-utility.html This message shows the progress of the tuning utility creating logic YAML files after reading benchmark results and matching tables. ```text Creating exact logic --Reading matching table: /build/release/library/MatchTable.yaml --Reading bench files --Found file /0_Bench/result_NN_SSS_88x12x664.txt Writing logic yaml files: 100%| | 1/1 [00:05<00:00, 5.69s/it] ``` -------------------------------- ### hipblaslt-bench Output with Tuning Results Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-offline-tuning.html This output shows the hipblaslt-bench command used after tuning and the 'Winner' section, which details the optimal kernel configuration found, including the solution index. ```text ./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,soulution_index 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 ``` -------------------------------- ### hipblasLtMatmulPreference_t Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/datatypes.html A descriptor for matrix multiplication preferences, used to guide the heuristic function. ```APIDOC ## hipblasLtMatmulPreference_t ### Description Descriptor of the matrix multiplication preference. This is a pointer to an opaque structure holding the description of the preferences for the hipblasLtMatmulAlgoGetHeuristic() configuration. Use hipblasLtMatmulPreferenceCreate() to create an instance and hipblasLtMatmulPreferenceDestroy() to destroy it. ### Type typedef hipblasLtMatmulPreferenceOpaque_t *hipblasLtMatmulPreference_t ``` -------------------------------- ### hipblasLtMatrixTransformDescGetAttribute Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/api-reference.html Gets an attribute from a matrix transform operation descriptor. This retrieves the current configuration of a specific attribute. ```APIDOC ## hipblasLtMatrixTransformDescGetAttribute() ### Description Matrix transform operation getter. Get matrix transform operation descriptor attribute. ### Signature ```c hipblasStatus_t hipblasLtMatrixTransformDescGetAttribute(hipblasLtMatrixTransformDesc_t transformDesc, hipblasLtMatrixTransformDescAttributes_t attr, void *buf, size_t sizeInBytes, size_t *sizeWritten) ``` ### Parameters * **transformDesc** (*hipblasLtMatrixTransformDesc_t*) - [in] The descriptor. * **attr** (*hipblasLtMatrixTransformDescAttributes_t*) - [in] The attribute to get. * **buf** (*void* *) - [out] Memory address to store the attribute value. * **sizeInBytes** (*size_t*) - [in] Size of buf buffer for verification (in bytes). * **sizeWritten** (*size_t* *) - [out] Number of bytes actually written or needed. ### Return values * **HIPBLAS_STATUS_INVALID_VALUE** – if sizeInBytes is 0 and sizeWritten is NULL, or if sizeInBytes is non-zero and buf is NULL or sizeInBytes doesn’t match size of internal storage for selected attribute. * **HIPBLAS_STATUS_SUCCESS** – if attribute’s value was successfully written to user memory. ``` -------------------------------- ### Build hipBLASLt dependencies, library, and client locally Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Builds the library dependencies, client dependencies, library, and client in the local directory. The `-d` flag is for initial dependency building. ```bash ./install.sh -dc ``` -------------------------------- ### GemmInstance::setMaxWorkspaceBytes and GemmInstance::getMaxWorkspaceBytes Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Functions to set and get the maximum workspace size in bytes for GEMM operations. ```APIDOC ## GemmInstance::setMaxWorkspaceBytes ### Description This function sets the max workspace size. ### Method `void setMaxWorkspaceBytes(size_t workspaceBytes)` ### Parameters #### Path Parameters - **workspaceBytes** (size_t) - [in] Set the max workspace size in bytes. ## GemmInstance::getMaxWorkspaceBytes ### Description This function returns the set max workspace size. ### Method `const size_t getMaxWorkspaceBytes() const` ### Return values - **size_t**: Returns the set max workspace size. ``` -------------------------------- ### hipblasltExtAMax Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-ops.html Determines the absolute maximum value of a 2D tensor. This is an independent function used to call and get the result. ```APIDOC ## hipblasltExtAMax() ### Description Determines the absolute maximum value of a 2D tensor. This is an independent function used to call and get the result. ### Signature (Signature not provided in source text) ### Parameters (Parameters not provided in source text) ### Return Values (Return values not provided in source text) ``` -------------------------------- ### Configure and Retrieve GEMM Heuristics Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Sets GEMM problem details, configures preferences, retrieves heuristic solutions, and prepares for kernel execution. ```cpp // Pseudo code ipblaslt_ext::GemmPreference pref; pref.setMaxWorkspaceBytes(1000000); // Default epilogue mode is HIPBLASLT_EPILOGUE_DEFAULT ipblaslt_ext::GemmEpilogue epilogue; ipblaslt_ext::GemmInputs inputs; inputs.setA(d_a); inputs.setB(d_b); inputs.setC(d_c); inputs.setD(d_d); inputs.setAlpha(&alpha); inputs.setBeta(&beta); ipblaslt_ext::Gemm gemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIPBLAS_COMPUTE_32F); std::vector heuristic; gemm.setProblem(1, 1, 1, 1, epilogue, inputs); // m, n, k, batch gemm.algoGetHeuristic(gemm, pref, heuristic); ``` -------------------------------- ### Algorithm Indexing Functions Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Provides functions to get an algorithm's index from its `hipblasLtMatmulAlgo_t` object and to retrieve heuristic results using a vector of algorithm indices. ```APIDOC ## Algorithm Indexing Functions ### Description Extension APIs to get the algorithm index using `hipblasLtMatmulAlgo_t` and to retrieve heuristic results using a vector of algorithm indices. ### Methods - **getIndexFromAlgo** - **Description**: Gets the index from a given algorithm. - **Signature**: `HIPBLASLT_EXPORT int getIndexFromAlgo(hipblasLtMatmulAlgo_t& algo)` - **getAlgosFromIndex** - **Description**: Retrieves heuristic results based on a provided vector of algorithm indices. - **Signature**: ```cpp HIPBLASLT_EXPORT hipblasStatus_t getAlgosFromIndex( hipblasLtHandle_t handle, std::vector& algoIndex, std::vector& heuristicResults ); ``` ### Usage Example ```cpp // You can use an index vector to retrieve the heuristic results. // Example: Assuming 'algoIndex' is populated with valid indices. // HIPBLASLT_EXPORT hipblasStatus_t getAlgosFromIndex(handle, algoIndex, heuristicResults); ``` ``` -------------------------------- ### Configure and Build hipBLASLt Clients Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/install/building-installing-hipblaslt.html Configures and builds hipBLASLt clients, tests, and benchmarks. Use -DCMAKE_PREFIX_PATH to specify locations of custom dependency builds if not in default system directories. Uses amdclang++ as the C++ compiler. ```bash -DCMAKE_PREFIX_PATH="" # Default install location is in /opt/rocm, use -DCMAKE_INSTALL_PREFIX= to specify other CXX=/opt/rocm/bin/amdclang++ ccmake -DBUILD_CLIENTS_TESTS=ON -DBUILD_CLIENTS_BENCHMARKS=ON [HIPBLASLT_SOURCE] make -j$(nproc) sudo make install # sudo required if installing into system directory such as /opt/rocm ``` -------------------------------- ### Get Algorithm Index from Algo Object Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html This function retrieves an integer index for a given hipblasLtMatmulAlgo_t object. This index can be used with other APIs to refer to specific algorithms. ```cpp HIPBLASLT_EXPORT int getIndexFromAlgo(hipblasLtMatmulAlgo_t& algo); ``` -------------------------------- ### Run Quick hipBLASLt Demo Tests Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/conceptual/hipblaslt-clients.html Execute a quick set of demo tests for hipBLASLt by filtering tests that contain 'quick' in their name. This provides a fast check of basic functionality. ```bash ./hipblaslt-test --gtest_filter=*quick* ``` -------------------------------- ### GEMM Heuristics and Execution Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Code snippet demonstrating how to retrieve heuristic kernel configurations and prepare for GEMM execution. ```APIDOC ## GEMM Heuristics and Execution ### Description This section shows how to retrieve heuristic kernel configurations and set up inputs for GEMM execution. ### Pseudo Code Example ```cpp // Pseudo code hipblaslt_ext::GemmPreference pref; pref.setMaxWorkspaceBytes(1000000); // Default epilogue mode is HIPBLASLT_EPILOGUE_DEFAULT hipblaslt_ext::GemmEpilogue epilogue; hipblaslt_ext::GemmInputs inputs; inputs.setA(d_a); inputs.setB(d_b); inputs.setC(d_c); inputs.setD(d_d); inputs.setAlpha(&alpha); inputs.setBeta(&beta); hipblaslt_ext::Gemm gemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIPBLAS_COMPUTE_32F); std::vector heuristic; gemm.setProblem(1, 1, 1, 1, epilogue, inputs); // m, n, k, batch gemm.algoGetHeuristic(gemm, pref, heuristic); ``` ``` -------------------------------- ### hipblasLtCreate Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/api-reference.html Creates a hipblaslt handle, initializing the library and allocating resources. This must be called before any other hipBLASLt functions. ```APIDOC ## hipblasLtCreate() ### Description Create a hipblaslt handle. This function initializes the hipBLASLt library and creates a handle to an opaque structure holding the hipBLASLt library context. It allocates light hardware resources on the host and device, and must be called prior to making any other hipBLASLt library calls. The hipBLASLt library context is tied to the current ROCm device. To use the library on multiple devices, one hipBLASLt handle should be created for each device. ### Parameters #### Path Parameters - **handle** (hipblasLtHandle_t *) - [out] Pointer to the allocated hipBLASLt handle for the created hipBLASLt context. ### Return values - **HIPBLAS_STATUS_SUCCESS** – The allocation completed successfully. - **HIPBLAS_STATUS_INVALID_VALUE** – `handle` == NULL. ``` -------------------------------- ### GEMM Operation with hipblaslt Extension Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Demonstrates how to perform a General Matrix Multiply (GEMM) operation using the hipblaslt extension. It covers algorithm discovery, preference setting, epilogue configuration, and running the GEMM kernel. ```cpp // Pseudo code for gemm problem // Get all algorithms std::vector heuristicResult; CHECK_HIPBLASLT_ERROR(hipblaslt_ext::getAllAlgos(handle, HIPBLASLT_GEMM, HIPBLAS_OP_N, HIPBLAS_OP_N, in_out_datatype, in_out_datatype, in_out_datatype, in_out_datatype, HIPBLAS_COMPUTE_32F, heuristicResult)); ipblaslt_ext::GemmPreference pref; pref.setMaxWorkspaceBytes(1000000); ipblaslt_ext::GemmEpilogue epilogue; epilogue.setMode(HIPBLASLT_EPILOGUE_GELU); ipblaslt_ext::GemmInputs inputs; inputs.setA(d_a); inputs.setB(d_b); inputs.setC(d_c); inputs.setD(d_d); inputs.setAlpha(&alpha); inputs.setBeta(&beta); ipblaslt_ext::Gemm gemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIP_R_16F, HIPBLAS_COMPUTE_32F); gemm.setProblem(1, 1, 1, 1, epilogue, inputs); // m, n, k, batch validIdx.clear(); for(int j = 0; j < heuristicResult.size(); j++) { size_t workspace_size = 0; if(gemm.isAlgoSupported(heuristicResult[j].algo, workspace_size) == HIPBLAS_STATUS_SUCCESS) { validIdx.push_back(j); heuristicResult[j].workspaceSize = workspace_size; } else { heuristicResult[j].workspaceSize = 0; } } if(validIdx.size() > 1) { gemm.initialize(heuristicResult[validIdx[0]].algo, d_workspace, stream); for(int i = 0; i < 10; i++) { gemm.run(stream); } } ``` -------------------------------- ### Limit Stream-K Compute Units Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-streamk.html Use TENSILE_STREAMK_MAX_CUS to restrict the number of compute units that Stream-K kernels can utilize. This example limits GEMM kernels to 32 compute units. ```bash export TENSILE_STREAMK_MAX_CUS=32 ``` -------------------------------- ### GemmInstance::initialize Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html Creates kernel arguments from a GemmInstance, optionally performing an implicit heuristics query. ```APIDOC ## GemmInstance::initialize (overload 1) ### Description Create kernel arguments from a given hipblaslt_ext::GemmInstance. This function creates kernel arguments from a given hipblaslt_ext::GemmInstance then saves the arguments inside the instance. ### Method `hipblasStatus_t initialize(const hipblasLtMatmulAlgo_t &algo, void *workspace, bool useUserArgs = true, hipStream_t stream = 0)` ### Parameters #### Path Parameters - **algo** (const hipblasLtMatmulAlgo_t&) - [in] Handle for matrix multiplication algorithm to be used. When NULL, an implicit heuristics query with default search preferences will be performed to determine actual algorithm to use. - **workspace** (void*) - [in] Pointer to the workspace buffer allocated in the GPU memory. Pointer must be 16B aligned. - **useUserArgs** (bool) - [in] Use user args, this does not affect vanilla gemm. (May be deprecated in the future) - **stream** (hipStream_t) - [in] The HIP stream where all the GPU work will be submitted. (May be deprecated in the future) ### Return values - **HIPBLAS_STATUS_SUCCESS**: If the operation completed successfully. - **HIPBLAS_STATUS_INVALID_VALUE**: If the gemm_count = 0 or workspace is null but workspaceBytes is greater than zero. Note that workspaceBytes should be set with setMaxWorkspaceBytes. ``` ```APIDOC ## GemmInstance::initialize (overload 2) ### Description Create kernel arguments from a given hipblaslt_ext::GemmInstance with specific tuning parameters. This function creates kernel arguments from a given hipblaslt_ext::GemmInstance then saves the arguments inside the instance. ### Method `hipblasStatus_t initialize(const hipblasLtMatmulAlgo_t &algo, GemmTuning &tuning, void *workspace, bool useUserArgs = true, hipStream_t stream = 0)` ### Parameters #### Path Parameters - **algo** (const hipblasLtMatmulAlgo_t&) - [in] Handle for matrix multiplication algorithm to be used. When NULL, an implicit heuristics query with default search preferences will be performed to determine actual algorithm to use. - **tuning** (GemmTuning&) - [in] The tuning parameters. - **workspace** (void*) - [in] Pointer to the workspace buffer allocated in the GPU memory. Pointer must be 16B aligned. - **useUserArgs** (bool) - [in] Use user args, this does not affect vanilla gemm. (May be deprecated in the future) - **stream** (hipStream_t) - [in] The HIP stream where all the GPU work will be submitted. (May be deprecated in the future) ### Return values - **HIPBLAS_STATUS_SUCCESS**: If the operation completed successfully. - **HIPBLAS_STATUS_INVALID_VALUE**: If the gemm_count = 0 or workspace is null but workspaceBytes is greater than zero. Note that workspaceBytes should be set with setMaxWorkspaceBytes. ``` -------------------------------- ### Get Heuristic Results from Algorithm Indices Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html This API retrieves heuristic results for a list of algorithms specified by their integer indices. It requires a handle to the hipBLASLt library and a vector of algorithm indices. ```cpp HIPBLASLT_EXPORT hipblasStatus_t getAlgosFromIndex(hipblasLtHandle_t handle, std::vector& algoIndex, std::vector& heuristicResults); ``` -------------------------------- ### Run hipBLASLt Tuning Utility Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/how-to/how-to-use-hipblaslt-tuning-utility.html Execute the `find_exact.py` script with the path to your YAML configuration file, the hipBLASLt build directory, and the desired output folder. ```bash python3 find_exact.py /build/release ``` -------------------------------- ### Get All Algorithms for a Problem Type Source: https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/reference/ext-reference.html This API retrieves all available algorithms for a specific matrix multiplication problem defined by transpose operations, data types, and compute type. It does not require problem size or epilogue information. ```cpp HIPBLASLT_EXPORT hipblasStatus_t hipblaslt_ext::getAllAlgos(hipblasLtHandle_t handle, hipblasLtExtGemmTypeEnum_t typeGemm, hipblasOperation_t opA, hipblasOperation_t opB, hipDataType typeA, hipDataType typeB, hipDataType typeC, hipDataType typeD, hipblasComputeType_t typeCompute, std::vector& heuristicResults); ```