# CUDA Batched Matrix Multiplication Guide

December 25, 2025 11 min By RightNow AI Team

## Introduction

Batched GEMM executes multiple small matrix multiplications in parallel, critical for multi-head attention, batch processing, and grouped convolutions. Understanding when to use batched GEMM vs. reshaping to single large GEMM is key to transformer optimization.

This guide covers cuBLAS batched operations, memory layouts, and custom kernels for small matrices where cuBLAS overhead dominates.

## Common Performance Issues

- cuBLAS launch overhead dominates for small matrices
- Memory layout mismatches requiring transposes
- Not leveraging strided batching
- Suboptimal batch dimension ordering

## Optimization Techniques

### 1. Strided Batched GEMM

Use `cublasGemmStridedBatched` for contiguous batch layout without pointer arrays.

### 2. Persistent Kernels

For many small GEMMs, custom persistent kernels avoid launch overhead.

### 3. Memory Layout Optimization

Store Q,K,V as [batch, head, seq, dim] for strided access.

## Implementation Comparison

### Before (Naive Implementation)

Launching separate GEMMs has high overhead for many small matrices.

```c
// Naive: launch separate GEMM for each batch
void batched_gemm_naive(cublasHandle_t handle,
                        float* A, float* B, float* C,
                        int M, int N, int K, int batch) {
    float alpha = 1.0f, beta = 0.0f;
    for (int i = 0; i < batch; i++) {
        cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
                    N, M, K, &alpha,
                    B + i * K * N, N,
                    A + i * M * K, K,
                    &beta, C + i * M * N, N);
    }
}
```

### After (Optimized Implementation)

Strided batched GEMM eliminates loop overhead and enables parallelism.

```c
// Optimized: single batched GEMM launch
void batched_gemm_strided(cublasHandle_t handle,
                          float* A, float* B, float* C,
                          int M, int N, int K, int batch) {
    float alpha = 1.0f, beta = 0.0f;

// Strided batched - no pointer array needed
    cublasSgemmStridedBatched(handle, CUBLAS_OP_N, CUBLAS_OP_N,
        N, M, K,
        &alpha,
        B, N, K * N,      // B stride
        A, K, M * K,      // A stride
        &beta,
        C, N, M * N,      // C stride
        batch);
}

// For multi-head attention: reshape Q,K,V for efficient batching
// Input: [batch, seq, heads * dim]
// Reshape to: [batch * heads, seq, dim] for strided batched GEMM
```

## Performance Results

| Metric | Naive | Optimized | Improvement |
| --- | --- | --- | --- |
| Throughput (small matrices) | 1x | 8-15x | Reduced launch overhead |
| Multi-head Attention | 1x | 2x | Proper batching |

## Frequently Asked Questions

### When should I use batched GEMM vs single large GEMM?

Use batched GEMM when batch elements have different sizes or when reshape/concat would require memory copies. For same-size batches with contiguous memory, reshaping to large GEMM may be faster.