Using general batched GEMM

Using general batched GEMM#

General batched GEMM is a hipBLASLt feature that performs multiple independent GEMM operations in a single API call. Each batch references its own matrices through device-resident pointer arrays instead of a single contiguous buffer with uniform strides.

Operation: D = alpha * op(A) * op(B) + beta * C

Each batch can use separate memory allocations for A, B, C, and D. This guide covers bench usage, the API workflow, a complete standalone example, and common pitfalls.

Prerequisites: A ROCm installation with hipBLASLt built for your target GPU architecture. For build instructions, see Build from source.

What is general batched GEMM?#

In GEMM, general is standard BLAS terminology for the usual rectangular matrix multiply (as opposed to symmetric, triangular, or other specialized variants). In general batched GEMM, general refers to the batching layout: each batch matrix may live at a separate device address, referenced through a pointer array, rather than in one contiguous strided buffer. This is also known as pointer-array batched GEMM.

General batched GEMM runs multiple GEMM operations where:

  • Each batch has its own matrices in separate memory locations.

  • Matrices are referenced through device pointer arrays (A[], B[], C[], D[]).

  • Batch mode (hipblasLtBatchMode_t) is set to HIPBLASLT_BATCH_MODE_POINTER_ARRAY (or 1).

  • Strided-batch offset attributes are not used; each pointer in the array identifies the base address of one matrix.

This mode is useful when:

  • Matrices for different batches cannot be stored contiguously.

  • You are integrating with code that already owns separate allocations.

  • You need maximum flexibility in where each batch’s data lives in device memory.

General batched vs. strided batched GEMM#

hipBLASLt supports two batching modes. Choosing the right one affects both setup and performance.

Strided batched GEMM (HIPBLASLT_BATCH_MODE_STRIDED)#

Memory layout:

A: [A0 | A1 | A2 | A3]  <- single contiguous allocation
B: [B0 | B1 | B2 | B3]  <- single contiguous allocation
C: [C0 | C1 | C2 | C3]  <- single contiguous allocation
D: [D0 | D1 | D2 | D3]  <- single contiguous allocation

Characteristics:

  • All matrices live in one contiguous buffer per operand.

  • Batch mode (hipblasLtBatchMode_t) is set to HIPBLASLT_BATCH_MODE_STRIDED (or 0).

  • Uniform stride between batches set with HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET.

  • More memory-efficient and often slightly better cache locality.

  • Less flexible: requires contiguous allocation.

Use when:

  • All batches share identical dimensions and leading dimensions.

  • You can allocate one buffer per operand.

  • Memory is limited or you want the simplest batched setup.

General batched GEMM (HIPBLASLT_BATCH_MODE_POINTER_ARRAY)#

Memory layout:

A: [ptr->A0, ptr->A1, ptr->A2, ptr->A3]  <- array of pointers
B: [ptr->B0, ptr->B1, ptr->B2, ptr->B3]
C: [ptr->C0, ptr->C1, ptr->C2, ptr->C3]
D: [ptr->D0, ptr->D1, ptr->D2, ptr->D3]

A0, A1, A2, A3 may reside anywhere in device memory.

Characteristics:

  • Each batch matrix is independently allocated.

  • No strided-batch offset attributes; pointer arrays carry the per-batch bases.

  • Works with pre-existing, separately allocated matrices.

  • Requires extra setup to build and copy pointer arrays to the device.

Use when:

  • Matrices are already allocated separately.

  • Contiguous batch buffers are impractical (for example, due to fragmentation).

  • You need to integrate with legacy or modular code paths.

Quick comparison#

Feature

Strided batched

General batched

Batch mode value (hipblasLtBatchMode_t)

0 (HIPBLASLT_BATCH_MODE_STRIDED)

1 (HIPBLASLT_BATCH_MODE_POINTER_ARRAY)

Memory layout

Single contiguous buffer

Separate allocations

Stride attributes

Required (uniform stride)

Not used

Setup complexity

Simple

Moderate (pointer arrays on device)

Memory efficiency

Higher

Moderate

Flexibility

Limited

Maximum (non-contiguous bases)

Best for

Uniform batch problems

Pre-allocated or non-contiguous batches

Important

A single hipblasLtMatmul call uses one matrix layout descriptor per operand. All batches in that call must therefore share the same m, n, k, and leading dimensions (lda, ldb, ldc, and ldd). General batched mode lets each batch live at a different address; it does not let each batch use different problem sizes in one call. For variable problem sizes, use grouped GEMM instead (see below).

Not the same as grouped GEMM#

The grouped GEMM API runs multiple GEMMs with different m, n, and/or k values in one launch. General batched GEMM runs multiple GEMMs that share the same dimensions and layout metadata but use separate memory for each batch.

For strided batching examples in the hipBLASLt repository, see clients/samples/02_hipblaslt_gemm_batched/.

When to use general batched GEMM#

Choose general batched GEMM when you have:

  1. Pre-allocated matrices already living in separate device buffers.

  2. Dynamic batch management where batches are added or removed independently.

  3. Memory fragmentation that prevents large contiguous allocations.

  4. Legacy integration with code paths that already use per-batch pointers.

  5. Different data sources where matrices originate from separate modules or structures.

Choose strided batched GEMM when you:

  1. Are allocating memory specifically for batched GEMM.

  2. Have uniform batch dimensions and can use contiguous buffers.

  3. Want the simplest setup and typically the best performance for uniform problems.

Using hipblaslt-bench#

Basic command structure#

./hipblaslt-bench [options]

Key options for general batched GEMM#

Option

Description

Default

Example

--batch_mode

Batch mode: 0 = strided, 1 = general (pointer array)

0

--batch_mode 1

--batch_count

Number of batches

1

--batch_count 4

-m, --sizem

Rows of op(A) and D

128

-m 512

-n, --sizen

Columns of op(B) and D

128

-n 256

-k, --sizek

Columns of op(A) / rows of op(B)

128

-k 128

--alpha

Scalar alpha

1.0

--alpha 1.5

--beta

Scalar beta

0.0

--beta 0.5

--transA

Transpose A: N/T/C

N

--transA N

--transB

Transpose B: N/T/C

N

--transB T

-r, --precision

Data type

f16_r

--precision f32_r

--lda

Leading dimension of A

auto

--lda 512

--ldb

Leading dimension of B

auto

--ldb 128

--ldc

Leading dimension of C

auto

--ldc 512

--ldd

Leading dimension of D

auto

--ldd 512

-v, --verify

Enable CPU verification

disabled

-v

-i, --iters

Timing iterations

10

-i 100

Example: Basic general batched GEMM (FP32)#

./hipblaslt-bench --batch_mode 1 --batch_count 4 \
  -m 128 -n 128 -k 128 \
  --precision f32_r \
  --verify

This runs four independent 128 x 128 x 128 FP32 GEMMs in pointer-array mode and verifies the results against a CPU reference.

The first output line is a CSV header; subsequent lines report problem parameters and timing. The batch_count column reflects the number of pointer-array batches, and grouped_gemm remains 0 for this mode.

API overview#

Key steps#

The following outline shows the essential workflow. Steps omitted here (matmul preference, workspace, stream, and cleanup) are included in the complete example below.

// 1. Create hipBLASLt handle and stream
hipblasLtHandle_t handle;
hipblasLtCreate(&handle);

// 2. Allocate individual device matrices for each batch
for(int i = 0; i < batch_count; ++i) {
    hipMalloc(&d_a[i], size_a * sizeof(float));
    // ... B, C, D ...
}

// 3. Build device-resident pointer arrays
float** d_ptr_array_a;
hipMalloc(&d_ptr_array_a, sizeof(float*) * batch_count);
hipMemcpy(d_ptr_array_a, d_a.data(),
          sizeof(float*) * batch_count, hipMemcpyHostToDevice);
// Repeat for B, C, D.

// 4. Create matrix layouts (same m, n, k, lda for all batches)
hipblasLtMatrixLayoutCreate(&mat_a, HIP_R_32F, m, k, lda);

// 5. Set batch attributes on all four layouts
hipblasLtBatchMode_t batch_mode = HIPBLASLT_BATCH_MODE_POINTER_ARRAY;
hipblasLtMatrixLayoutSetAttribute(
    mat_a, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT,
    &batch_count, sizeof(batch_count));
hipblasLtMatrixLayoutSetAttribute(
    mat_a, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE,
    &batch_mode, sizeof(batch_mode));

// 6. Create matmul descriptor (set transposes and epilogue)
hipblasLtMatmulDescCreate(&matmul_desc, HIPBLAS_COMPUTE_32F, HIP_R_32F);

// 7. Query a heuristic and allocate workspace
hipblasLtMatmulAlgoGetHeuristic(
    handle, matmul_desc, mat_a, mat_b, mat_c, mat_d,
    pref, 1, &heuristic_result, &returned_count);

// 8. Execute with pointer arrays passed as the matrix pointers
hipblasLtMatmul(handle, matmul_desc,
    &alpha,
    d_ptr_array_a, mat_a,
    d_ptr_array_b, mat_b,
    &beta,
    d_ptr_array_c, mat_c,
    d_ptr_array_d, mat_d,
    &heuristic_result.algo,
    workspace, workspace_size,
    stream);

Required layout attributes#

For general batched GEMM, set both of the following on all four matrix layouts (A, B, C, D):

  1. HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, number of batches.

  2. HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, HIPBLASLT_BATCH_MODE_POINTER_ARRAY.

Do not set HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET in this mode.

Optional sub-matrix offset#

When a batch matrix is a sub-region of a larger allocation, set HIPBLASLT_MATRIX_LAYOUT_OFFSET (in elements from the base pointer) on the layout. Offsets are only valid when batch mode is pointer-array on all four operands. See hipBLASLt datatypes.

Known limitations#

  • Uniform dimensions per call: All batches in one hipblasLtMatmul must share the same m, n, k, and leading dimensions.

  • Scaling formats: When using hipblaslt-bench with --batch_mode 1, only tensor-wide scaling is supported for matrices A and B; vector and block scaling modes are rejected.

  • Pointer arrays must be on device: Pass the device pointer to the pointer array, not a host-side array of addresses.

Complete example#

The following standalone program demonstrates general batched GEMM from scratch: per-batch allocation, device pointer arrays, layout configuration, heuristic selection, execution, CPU verification, and cleanup.

Compile and run#

From the docs/data/how-to directory (adjust ROCM_PATH if hipBLASLt is installed elsewhere):

ROCM_PATH=$(hipconfig --rocm-path)
hipcc -o general_batched_gemm \
    example_general_batched_gemm_standalone.cpp \
    -I${ROCM_PATH}/include \
    -L${ROCM_PATH}/lib -lhipblaslt

./general_batched_gemm

Expected output:

=== general batched GEMM example ===
Problem size: M=128, N=64, K=96
Batch count: 4
Alpha=1.5, Beta=-0.5

Algorithm selected with workspace size: XXXXX bytes
Executing general batched GEMM...
GEMM execution completed.
Batch 0: PASSED
Batch 1: PASSED
Batch 2: PASSED
Batch 3: PASSED

SUCCESS: All batches passed verification!

Source code#

// Copyright Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

/*
 * Standalone General Batched GEMM Example using hipBLASLt
 *
 * This sample demonstrates how to use hipBLASLt to perform General Batched GEMM operations.
 * General Batched GEMM uses pointer arrays to reference individual matrices for each batch,
 * unlike Strided Batched GEMM which uses a single base pointer with uniform strides.
 *
 * Operation: D = alpha * op(A) * op(B) + beta * C
 *
 * Where:
 *   - Each batch has its own matrix data stored separately
 *   - Matrices are referenced via device pointer arrays
 *   - Batch mode is set to HIPBLASLT_BATCH_MODE_POINTER_ARRAY
 */

#include <hip/hip_runtime.h>
#include <hipblaslt/hipblaslt.h>

#include <cmath>
#include <cstddef>
#include <cstdlib>
#include <iostream>
#include <vector>

// Helper macro for checking HIP errors
#define HIP_CHECK(cmd)                                                                           \
    do {                                                                                         \
        hipError_t error = (cmd);                                                                \
        if(error != hipSuccess) {                                                                \
            std::cerr << "HIP error: " << hipGetErrorString(error) << " at " << __FILE__ << ":" \
                      << __LINE__ << std::endl;                                                  \
            exit(EXIT_FAILURE);                                                                  \
        }                                                                                        \
    } while(0)

// Helper macro for checking hipBLASLt errors
#define HIPBLASLT_CHECK(cmd)                                                                     \
    do {                                                                                         \
        hipblasStatus_t status = (cmd);                                                          \
        if(status != HIPBLAS_STATUS_SUCCESS) {                                                   \
            std::cerr << "hipBLASLt error: " << status << " at " << __FILE__ << ":" << __LINE__ \
                      << std::endl;                                                              \
            exit(EXIT_FAILURE);                                                                  \
        }                                                                                        \
    } while(0)

// Simple host GEMM for verification
void host_gemm(hipblasOperation_t        trans_a,
               hipblasOperation_t        trans_b,
               int                       m,
               int                       n,
               int                       k,
               float                     alpha,
               const std::vector<float>& a,
               int                       lda,
               const std::vector<float>& b,
               int                       ldb,
               float                     beta,
               const std::vector<float>& c,
               int                       ldc,
               std::vector<float>&       d,
               int                       ldd)
{
    for(int col = 0; col < n; ++col)
    {
        for(int row = 0; row < m; ++row)
        {
            float accum = 0.0f;
            for(int inner = 0; inner < k; ++inner)
            {
                float a_val, b_val;

                // Fetch A element based on transpose
                if(trans_a == HIPBLAS_OP_N)
                    a_val = a[row + inner * lda];
                else
                    a_val = a[inner + row * lda];

                // Fetch B element based on transpose
                if(trans_b == HIPBLAS_OP_N)
                    b_val = b[inner + col * ldb];
                else
                    b_val = b[col + inner * ldb];

                accum += a_val * b_val;
            }
            d[row + col * ldd] = alpha * accum + beta * c[row + col * ldc];
        }
    }
}

// Initialize matrix with pattern
void initialize_matrix(std::vector<float>& matrix, int rows, int cols, int ld, float seed)
{
    matrix.resize(ld * cols);
    for(int col = 0; col < cols; ++col)
    {
        for(int row = 0; row < rows; ++row)
        {
            matrix[row + col * ld] = seed + static_cast<float>(row + 1) * 0.25f
                                     + static_cast<float>(col + 1) * 0.5f;
        }
    }
}

// Compare two vectors with tolerance
bool verify_results(const std::vector<float>& result,
                    const std::vector<float>& reference,
                    float                     tolerance = 1.0e-4f)
{
    if(result.size() != reference.size())
        return false;

    for(size_t i = 0; i < result.size(); ++i)
    {
        if(std::fabs(result[i] - reference[i]) > tolerance)
        {
            std::cerr << "Mismatch at index " << i << ": result=" << result[i]
                      << ", reference=" << reference[i] << std::endl;
            return false;
        }
    }
    return true;
}

int main()
{
    // ============================================================================
    // Problem configuration
    // ============================================================================
    const hipblasOperation_t trans_a = HIPBLAS_OP_N;  // No transpose for A
    const hipblasOperation_t trans_b = HIPBLAS_OP_N;  // No transpose for B

    const int m = 128;                // Rows of op(A) and D
    const int n = 64;                 // Columns of op(B) and D
    const int k = 96;                 // Columns of op(A) and rows of op(B)

    const int lda = m;                // Leading dimension of A
    const int ldb = k;                // Leading dimension of B
    const int ldc = m;                // Leading dimension of C
    const int ldd = m;                // Leading dimension of D

    const int batch_count = 4;        // Number of batches
    const hipblasLtBatchMode_t batch_mode = HIPBLASLT_BATCH_MODE_POINTER_ARRAY;

    const float alpha = 1.5f;         // Scalar alpha
    const float beta  = -0.5f;        // Scalar beta

    const size_t max_workspace_bytes = 64 * 1024 * 1024;  // 64 MB workspace limit

    std::cout << "=== general batched GEMM example ===" << std::endl;
    std::cout << "Problem size: M=" << m << ", N=" << n << ", K=" << k << std::endl;
    std::cout << "Batch count: " << batch_count << std::endl;
    std::cout << "Alpha=" << alpha << ", Beta=" << beta << std::endl;
    std::cout << std::endl;

    // ============================================================================
    // Initialize hipBLASLt
    // ============================================================================
    hipblasLtHandle_t handle;
    HIPBLASLT_CHECK(hipblasLtCreate(&handle));

    hipStream_t stream;
    HIP_CHECK(hipStreamCreate(&stream));

    // ============================================================================
    // Allocate and initialize host matrices for each batch
    // ============================================================================
    std::vector<std::vector<float>> h_a(batch_count);
    std::vector<std::vector<float>> h_b(batch_count);
    std::vector<std::vector<float>> h_c(batch_count);
    std::vector<std::vector<float>> h_d(batch_count);
    std::vector<std::vector<float>> h_d_ref(batch_count);

    const size_t size_a = lda * k;
    const size_t size_b = ldb * n;
    const size_t size_c = ldc * n;
    const size_t size_d = ldd * n;

    for(int i = 0; i < batch_count; ++i)
    {
        initialize_matrix(h_a[i], m, k, lda, 10.0f * i + 1.0f);
        initialize_matrix(h_b[i], k, n, ldb, 20.0f * i + 2.0f);
        initialize_matrix(h_c[i], m, n, ldc, 30.0f * i + 3.0f);

        h_d[i].resize(size_d, 0.0f);
        h_d_ref[i].resize(size_d, 0.0f);

        // Compute reference on CPU
        host_gemm(trans_a, trans_b, m, n, k, alpha, h_a[i], lda, h_b[i], ldb,
                  beta, h_c[i], ldc, h_d_ref[i], ldd);
    }

    // ============================================================================
    // Allocate device memory for each batch
    // ============================================================================
    std::vector<float*> d_a(batch_count);
    std::vector<float*> d_b(batch_count);
    std::vector<float*> d_c(batch_count);
    std::vector<float*> d_d(batch_count);

    for(int i = 0; i < batch_count; ++i)
    {
        HIP_CHECK(hipMalloc(&d_a[i], sizeof(float) * size_a));
        HIP_CHECK(hipMalloc(&d_b[i], sizeof(float) * size_b));
        HIP_CHECK(hipMalloc(&d_c[i], sizeof(float) * size_c));
        HIP_CHECK(hipMalloc(&d_d[i], sizeof(float) * size_d));

        // Copy host data to device
        HIP_CHECK(hipMemcpy(d_a[i], h_a[i].data(), sizeof(float) * size_a, hipMemcpyHostToDevice));
        HIP_CHECK(hipMemcpy(d_b[i], h_b[i].data(), sizeof(float) * size_b, hipMemcpyHostToDevice));
        HIP_CHECK(hipMemcpy(d_c[i], h_c[i].data(), sizeof(float) * size_c, hipMemcpyHostToDevice));
    }

    // ============================================================================
    // Create device pointer arrays (key for General Batched GEMM)
    // ============================================================================
    float** d_ptr_array_a;
    float** d_ptr_array_b;
    float** d_ptr_array_c;
    float** d_ptr_array_d;

    HIP_CHECK(hipMalloc(&d_ptr_array_a, sizeof(float*) * batch_count));
    HIP_CHECK(hipMalloc(&d_ptr_array_b, sizeof(float*) * batch_count));
    HIP_CHECK(hipMalloc(&d_ptr_array_c, sizeof(float*) * batch_count));
    HIP_CHECK(hipMalloc(&d_ptr_array_d, sizeof(float*) * batch_count));

    // Copy pointer arrays to device
    HIP_CHECK(hipMemcpy(d_ptr_array_a, d_a.data(), sizeof(float*) * batch_count, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_ptr_array_b, d_b.data(), sizeof(float*) * batch_count, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_ptr_array_c, d_c.data(), sizeof(float*) * batch_count, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_ptr_array_d, d_d.data(), sizeof(float*) * batch_count, hipMemcpyHostToDevice));

    // ============================================================================
    // Create matrix layouts
    // ============================================================================
    hipblasLtMatrixLayout_t mat_a, mat_b, mat_c, mat_d;

    HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&mat_a, HIP_R_32F, m, k, lda));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&mat_b, HIP_R_32F, k, n, ldb));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&mat_c, HIP_R_32F, m, n, ldc));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&mat_d, HIP_R_32F, m, n, ldd));

    // Set batch count for all matrices
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_a, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_b, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_c, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_d, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count, sizeof(batch_count)));

    // Set batch mode to General Batched (pointer array mode)
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_a, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_b, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_c, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode)));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(
        mat_d, HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE, &batch_mode, sizeof(batch_mode)));

    // Note: Stride attributes are not used in General Batched mode,
    // as each batch has its own separate memory location

    // ============================================================================
    // Create matmul descriptor
    // ============================================================================
    hipblasLtMatmulDesc_t matmul_desc;

    HIPBLASLT_CHECK(hipblasLtMatmulDescCreate(&matmul_desc, HIPBLAS_COMPUTE_32F, HIP_R_32F));
    HIPBLASLT_CHECK(hipblasLtMatmulDescSetAttribute(
        matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSA, &trans_a, sizeof(trans_a)));
    HIPBLASLT_CHECK(hipblasLtMatmulDescSetAttribute(
        matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSB, &trans_b, sizeof(trans_b)));

    hipblasLtEpilogue_t epilogue = HIPBLASLT_EPILOGUE_DEFAULT;
    HIPBLASLT_CHECK(hipblasLtMatmulDescSetAttribute(
        matmul_desc, HIPBLASLT_MATMUL_DESC_EPILOGUE, &epilogue, sizeof(epilogue)));

    // ============================================================================
    // Create matmul preference
    // ============================================================================
    hipblasLtMatmulPreference_t pref;

    HIPBLASLT_CHECK(hipblasLtMatmulPreferenceCreate(&pref));
    HIPBLASLT_CHECK(hipblasLtMatmulPreferenceSetAttribute(
        pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &max_workspace_bytes, sizeof(max_workspace_bytes)));

    // ============================================================================
    // Get heuristic for algorithm selection
    // ============================================================================
    hipblasLtMatmulHeuristicResult_t heuristic_result;
    int returned_algo_count = 0;

    HIPBLASLT_CHECK(hipblasLtMatmulAlgoGetHeuristic(
        handle,
        matmul_desc,
        mat_a,
        mat_b,
        mat_c,
        mat_d,
        pref,
        1,  // Request one algorithm
        &heuristic_result,
        &returned_algo_count));

    if(returned_algo_count == 0)
    {
        std::cerr << "Error: No suitable algorithm found!" << std::endl;
        return EXIT_FAILURE;
    }

    std::cout << "Algorithm selected with workspace size: " << heuristic_result.workspaceSize
              << " bytes" << std::endl;

    // ============================================================================
    // Allocate workspace if needed
    // ============================================================================
    void* workspace = nullptr;
    if(heuristic_result.workspaceSize > 0)
    {
        HIP_CHECK(hipMalloc(&workspace, heuristic_result.workspaceSize));
    }

    // ============================================================================
    // Perform General Batched GEMM
    // ============================================================================
    std::cout << "Executing general batched GEMM..." << std::endl;

    HIPBLASLT_CHECK(hipblasLtMatmul(
        handle,
        matmul_desc,
        &alpha,
        d_ptr_array_a,  // Pointer array for A matrices
        mat_a,
        d_ptr_array_b,  // Pointer array for B matrices
        mat_b,
        &beta,
        d_ptr_array_c,  // Pointer array for C matrices
        mat_c,
        d_ptr_array_d,  // Pointer array for D matrices
        mat_d,
        &heuristic_result.algo,
        workspace,
        heuristic_result.workspaceSize,
        stream));

    HIP_CHECK(hipStreamSynchronize(stream));
    std::cout << "GEMM execution completed." << std::endl;

    // ============================================================================
    // Copy results back to host and verify
    // ============================================================================
    bool all_passed = true;

    for(int i = 0; i < batch_count; ++i)
    {
        HIP_CHECK(hipMemcpy(h_d[i].data(), d_d[i], sizeof(float) * size_d, hipMemcpyDeviceToHost));

        if(!verify_results(h_d[i], h_d_ref[i]))
        {
            std::cerr << "Verification FAILED for batch " << i << std::endl;
            all_passed = false;
        }
        else
        {
            std::cout << "Batch " << i << ": PASSED" << std::endl;
        }
    }

    // ============================================================================
    // Cleanup
    // ============================================================================
    if(workspace)
        HIP_CHECK(hipFree(workspace));

    for(int i = 0; i < batch_count; ++i)
    {
        HIP_CHECK(hipFree(d_a[i]));
        HIP_CHECK(hipFree(d_b[i]));
        HIP_CHECK(hipFree(d_c[i]));
        HIP_CHECK(hipFree(d_d[i]));
    }

    HIP_CHECK(hipFree(d_ptr_array_a));
    HIP_CHECK(hipFree(d_ptr_array_b));
    HIP_CHECK(hipFree(d_ptr_array_c));
    HIP_CHECK(hipFree(d_ptr_array_d));

    HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(mat_a));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(mat_b));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(mat_c));
    HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(mat_d));

    HIPBLASLT_CHECK(hipblasLtMatmulDescDestroy(matmul_desc));
    HIPBLASLT_CHECK(hipblasLtMatmulPreferenceDestroy(pref));

    HIP_CHECK(hipStreamDestroy(stream));
    HIPBLASLT_CHECK(hipblasLtDestroy(handle));

    // ============================================================================
    // Final result
    // ============================================================================
    std::cout << std::endl;
    if(all_passed)
    {
        std::cout << "SUCCESS: All batches passed verification!" << std::endl;
        return EXIT_SUCCESS;
    }
    else
    {
        std::cout << "FAILURE: Some batches failed verification." << std::endl;
        return EXIT_FAILURE;
    }
}

Performance considerations#

  • Strided vs. general: For uniform batches where you control allocation, strided batched GEMM is usually simpler and can be faster because the runtime avoids pointer-array indirection and extra device memory for the arrays.

  • Pointer-array overhead: general batched mode allocates four pointer arrays plus batch_count separate buffers. Factor that into memory planning.

  • Heuristic reuse: Query heuristics once with hipblasLtMatmulAlgoGetHeuristic() and reuse the selected algorithm for repeated calls with the same problem descriptor. See also Use logging and heuristics.

  • Workspace: Always honor workspaceSize returned by the heuristic; some algorithms require non-zero workspace for correct results.

Troubleshooting#

returned_algo_count == 0 (no suitable algorithm)

Widen HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, verify data types and transpose combinations are supported on your GPU, or try a different precision.

Verification failures

Confirm batch_count and HIPBLASLT_MATRIX_LAYOUT_BATCH_MODE match on all four layouts. Ensure pointer arrays were copied to the device with hipMemcpyHostToDevice.

Invalid value / status errors

Mixing batch modes across A/B/C/D layouts is invalid. Do not set strided-batch offsets when using pointer-array mode.

Segfaults or garbage results

The matrix pointer arguments to hipblasLtMatmul must be device pointers to pointer arrays (float** on device), not host arrays of float*.

Summary#

General batched GEMM is the right choice when you need separate device allocations per batch and cannot rely on contiguous strided buffers.

Key points:

  1. Set --batch_mode 1 (or HIPBLASLT_BATCH_MODE_POINTER_ARRAY) for pointer-array batching.

  2. Pass device-resident pointer arrays to hipblasLtMatmul.

  3. Do not use strided-batch offset attributes in this mode.

  4. All batches in one call share the same dimensions and leading dimensions.

  5. For variable m/n/k per problem, use grouped GEMM instead.

Quick-start bench command:

./hipblaslt-bench --batch_mode 1 --batch_count 4 \
  -m 512 -n 512 -k 512 --precision f32_r --verify
Contents