hipBLASLt

https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/
Admin
hipBLASLt is a library that provides General Matrix-Matrix (GEMM) operations with flexible APIs and...
Tokens:56,945
Snippets:432
Trust Score:10
Update:5 months ago
Show doc for...
Tokens:
Raw
### 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<hipblasLtMatmulHeuristicResult_t> 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<int64_t> m(gemm_count);
std::vector<int64_t> n(gemm_count);
std::vector<int64_t> k(gemm_count);
std::vector<int64_t> batch_count(gemm_count);
std::vector<hipblaslt_ext::GemmEpilogue> epilogue(gemm_count);
std::vector<hipblaslt_ext::GemmInputs> 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<int64_t> m, n, k;
// ...
for(size_t i = 0; i < problem_size, i++)
{
    // ...
}
std::vector<GemmProblemType> 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<int> algoIndex{index};
std::vector<hipblasLtMatmulHeuristicResult_t> 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=<path> to specify other
# Default build config is 'Release', define -DCMAKE_BUILD_TYPE=<config> 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<int64_t> m, n, k;
// ... populate m, n, k vectors ...
std::vector<int64_t> batch_count;
// ... populate batch_count vector ...
std::vector<int64_t> lda, ldb, ldc, ldd;
// ... populate stride vectors ...
std::vector<int64_t> strideA, strideB, strideC, strideD;
// ... populate stride vectors ...
std::vector<GemmEpilogue> epilogue;
// ... populate epilogue vector ...
std::vector<GemmInputs> inputs;
// ... populate inputs vector ...

std::vector<GemmProblemType> 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<hipblasLtMatmulHeuristicResult_t> 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<int64_t> m(gemm_count);
std::vector<int64_t> n(gemm_count);
std::vector<int64_t> k(gemm_count);
std::vector<int64_t> batch_count(gemm_count);
std::vector<hipblaslt_ext::GemmEpilogue> epilogue(gemm_count);
std::vector<hipblaslt_ext::GemmInputs> 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<hipblaslt_ext::UserArguments*>(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 <hip/hip_runtime.h>
#include <hipblaslt/hipblaslt-ext-op.h>
#include <hipblaslt/hipblaslt.h>
#include <iostream>

#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<float> 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<hipblasLtHalf> 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 <hip/hip_runtime.h>
#include <hipblaslt/hipblaslt-ext-op.h>
#include <hipblaslt/hipblaslt.h>
#include <hipblaslt_arguments.hpp>
#include <iostream>

#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<float> 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: <hipblaslt_root_folder>/build/release/library/MatchTable.yaml
--Reading bench files
--Found file <output folder>/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<hipblasLtMatmulHeuristicResult_t> 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<int>&                              algoIndex,
        std::vector<hipblasLtMatmulHeuristicResult_t>& 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="<semicolon separated paths>"
# Default install location is in /opt/rocm, use -DCMAKE_INSTALL_PREFIX=<path> 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<hipblasLtMatmulHeuristicResult_t> 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<hipblasLtMatmulHeuristicResult_t> 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<int>&                              algoIndex,
                      std::vector<hipblasLtMatmulHeuristicResult_t>& 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 <your yaml file> <hipblaslt_root_folder>/build/release <output folder>

```

--------------------------------

### 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<hipblasLtMatmulHeuristicResult_t>& heuristicResults);
```
Tokens:
Raw