> ## Documentation Index
> Fetch the complete documentation index at: https://mintlify.com/NVIDIA/cutlass/llms.txt
> Use this file to discover all available pages before exploring further.

# Hopper Architecture (SM90)

> CUTLASS support for NVIDIA Hopper architecture GPUs with WGMMA and TMA

The NVIDIA Hopper architecture (compute capability 9.0) represents a major leap in GPU computing with Warpgroup Matrix Multiply-Accumulate (WGMMA), Tensor Memory Accelerator (TMA), and native FP8 support.

## Supported GPUs

| GPU Model                   | Compute Capability | Min CUDA Toolkit |
| --------------------------- | ------------------ | ---------------- |
| NVIDIA H100 Tensor Core GPU | 9.0                | 11.8             |
| NVIDIA H200 Tensor Core GPU | 9.0                | 11.8             |

<Warning>
  Hopper architecture-accelerated features require compiling with the `sm_90a` target (note the "a" suffix) to enable WGMMA and TMA instructions:

  ```bash theme={null}
  cmake .. -DCUTLASS_NVCC_ARCHS="90a"
  ```
</Warning>

## Key Features

### 1. Warpgroup Matrix Multiply-Accumulate (WGMMA)

WGMMA operates on 4 warps (128 threads) simultaneously, dramatically improving instruction throughput:

```cpp theme={null}
// From include/cute/arch/mma_sm90_gmma.hpp
namespace cute::SM90::GMMA {

enum class Major {
  K  = 0,  // K-major layout
  MN = 1   // MN-major layout
};

enum class ScaleOut {
  Zero = 0,  // Zero out accumulator
  One  = 1   // Accumulate
};

// GMMA 64x8x16 F16+=F16*F16
template <
  GMMA::Major tnspA,
  GMMA::Major tnspB,
  GMMA::ScaleIn  scaleA = GMMA::ScaleIn::One,
  GMMA::ScaleIn  scaleB = GMMA::ScaleIn::One
>
struct MMA_64x8x16_F16F16F16_SS {
  using DRegisters = void;
  using ARegisters = uint64_t[1];
  using BRegisters = uint64_t[1];
  using CRegisters = uint32_t[2];
  
  CUTE_HOST_DEVICE static void
  fma(uint64_t const& desc_a,
      uint64_t const& desc_b,
      uint32_t      & d0, uint32_t      & d1,
      GMMA::ScaleOut const scale_D = GMMA::ScaleOut::One) {
#if defined(CUTE_ARCH_MMA_SM90A_ENABLED)
    asm volatile(
      "wgmma.mma_async.sync.aligned.m64n8k16.f16.f16.f16 "
      "{%0, %1}, %2, %3, %4, %5, %6;\n"
      : "+r"(d0), "+r"(d1)
      : "l"(desc_a), "l"(desc_b),
        "n"(int32_t(scaleA)), "n"(int32_t(scaleB)),
        "n"(int32_t(scale_D)));
#endif
  }
};

} // namespace cute::SM90::GMMA
```

**WGMMA Synchronization Primitives:**

```cpp theme={null}
// From include/cute/arch/mma_sm90_gmma.hpp
CUTE_HOST_DEVICE void warpgroup_arrive() {
#if defined(CUTE_ARCH_MMA_SM90A_ENABLED)
  asm volatile ("wgmma.fence.sync.aligned;\n" ::: "memory");
#endif
}

template <int N>
CUTE_HOST_DEVICE void warpgroup_wait() {
  static_assert(N >= 0 && N <= 7, "WGMMA wait: N must be in range [0, 7]");
#if defined(CUTE_ARCH_MMA_SM90A_ENABLED)
  asm volatile("wgmma.wait_group.sync.aligned %0;\n" :: "n"(N) : "memory");
#endif
}

CUTE_HOST_DEVICE void warpgroup_commit_batch() {
#if defined(CUTE_ARCH_MMA_SM90A_ENABLED)
  asm volatile("wgmma.commit_group.sync.aligned;\n" ::: "memory");
#endif
}
```

**Usage Pattern:**

```cpp theme={null}
// Issue WGMMA operations
warpgroup_arrive();
mma_operation(desc_a, desc_b, accum);
warpgroup_commit_batch();

// Continue with other work
// ...

// Wait for completion (0 = wait for all)
warpgroup_wait<0>();
```

### 2. Tensor Memory Accelerator (TMA)

TMA provides hardware-accelerated bulk data transfer between global and shared memory:

```cpp theme={null}
// From include/cute/arch/copy_sm90_tma.hpp
struct SM90_TMA_LOAD_2D {
  CUTE_HOST_DEVICE static void
  copy(void const* desc_ptr, 
       uint64_t* mbar_ptr, 
       uint64_t cache_hint,
       void* smem_ptr,
       int32_t const& crd0, 
       int32_t const& crd1) {
#if defined(CUTE_ARCH_TMA_SM90_ENABLED)
    uint64_t gmem_int_desc = reinterpret_cast<uint64_t>(desc_ptr);
    uint32_t smem_int_mbar = cast_smem_ptr_to_uint(mbar_ptr);
    uint32_t smem_int_ptr  = cast_smem_ptr_to_uint(smem_ptr);
    
    asm volatile (
      "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes"
      ".L2::cache_hint [%0], [%1, {%3, %4}], [%2], %5;\n"
      :
      : "r"(smem_int_ptr), "l"(gmem_int_desc), "r"(smem_int_mbar),
        "r"(crd0), "r"(crd1), "l"(cache_hint)
      : "memory");
#endif
  }
};
```

**TMA Benefits:**

* Hardware-managed address generation
* Automatic boundary checking
* Integrated with async barriers
* Support for up to 5D tensors
* Prefetch capabilities

### 3. Enhanced FP64 Tensor Cores

Hopper provides improved FP64 tensor core shapes:

```cpp theme={null}
// From include/cutlass/arch/mma_sm90.h
template <>
struct Mma<
  gemm::GemmShape<16,8,4>,
  32,
  double,
  layout::RowMajor,
  double,
  layout::ColumnMajor,
  double,
  layout::RowMajor,
  OpMultiplyAdd> {
  
  using Shape = gemm::GemmShape<16,8,4>;
  using ArchTag = arch::Sm90;
  
  CUTLASS_HOST_DEVICE
  void operator()(FragmentC &d, FragmentA const &a, 
                  FragmentB const &b, FragmentC const &c) const {
#if defined(CUTLASS_ARCH_MMA_SM90_F64_MMA_ENABLED)
    asm volatile(
      "mma.sync.aligned.m16n8k4.row.col.f64.f64.f64.f64.rn "
      "{%0, %1, %2, %3}, {%4, %5}, {%6}, {%7, %8, %9, %10};\n"
      : "=d"(D[0]), "=d"(D[1]), "=d"(D[2]), "=d"(D[3])
      : "d"(A[0]), "d"(A[1]), "d"(B[0]),
        "d"(C[0]), "d"(C[1]), "d"(C[2]), "d"(C[3]));
#endif
  }
};
```

**Supported FP64 Shapes:**

* 16x8x4 (new in Hopper)
* 16x8x8 (new in Hopper)
* 16x8x16 (new in Hopper)

### 4. FP8 Tensor Core Support

Hopper introduces native FP8 support with two formats:

* **E4M3**: 4 exponent bits, 3 mantissa bits (better for forward pass)
* **E5M2**: 5 exponent bits, 2 mantissa bits (better for gradients)

```cpp theme={null}
// FP8 E4M3 x E4M3 -> FP32
using ElementA = cutlass::float_e4m3_t;
using ElementB = cutlass::float_e4m3_t;
using ElementAccumulator = float;

// GMMA instruction shape for FP8
// 64x8x32 (much larger K dimension than FP16)
```

## Instruction Shapes

| Data Type       | WGMMA Shape (MxNxK)         | Accumulator |
| --------------- | --------------------------- | ----------- |
| FP8 (E4M3/E5M2) | 64x8x32, 64x16x32, 64x32x32 | FP32        |
| FP16            | 64x8x16, 64x16x16, 64x32x16 | FP16/FP32   |
| BF16            | 64x8x16, 64x16x16, 64x32x16 | FP32        |
| TF32            | 64x8x8, 64x16x8, 64x32x8    | FP32        |
| INT8            | 64x8x32, 64x16x32, 64x32x32 | INT32       |
| FP64            | 16x8x4, 16x8x8, 16x8x16     | FP64        |

<Note>
  Notice the significantly larger M and K dimensions compared to Ampere's 16x8xK shapes. This enables more efficient computation per instruction.
</Note>

## Thread Block Clusters

Hopper introduces thread block clusters for multi-CTA cooperation:

```cpp theme={null}
// From include/cute/arch/cluster_sm90.hpp
namespace cute {

// Get cluster dimensions
CUTE_HOST_DEVICE dim3 cluster_shape() {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900
  dim3 shape;
  asm volatile("mov.u32 %0, %%cluster_nctaid.x;\n" : "=r"(shape.x));
  asm volatile("mov.u32 %0, %%cluster_nctaid.y;\n" : "=r"(shape.y));
  asm volatile("mov.u32 %0, %%cluster_nctaid.z;\n" : "=r"(shape.z));
  return shape;
#else
  return {1, 1, 1};
#endif
}

// Synchronize cluster
CUTE_HOST_DEVICE void cluster_arrive() {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900
  asm volatile("barrier.cluster.arrive;\n" ::: "memory");
#endif
}

CUTE_HOST_DEVICE void cluster_wait() {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900
  asm volatile("barrier.cluster.wait;\n" ::: "memory");
#endif
}

} // namespace cute
```

**Cluster Configuration:**

```cpp theme={null}
cudaLaunchConfig_t config;
config.gridDim = dim3(num_blocks_x, num_blocks_y, num_blocks_z);
config.blockDim = dim3(128, 1, 1);  // Must be multiple of 128 for WGMMA

// Configure cluster
cudaLaunchAttribute cluster_attr;
cluster_attr.id = cudaLaunchAttributeClusterDimension;
cluster_attr.val.clusterDim.x = 2;  // 2x1x1 cluster
cluster_attr.val.clusterDim.y = 1;
cluster_attr.val.clusterDim.z = 1;

config.attrs = &cluster_attr;
config.numAttrs = 1;

cudaLaunchKernelEx(&config, kernel, args...);
```

## Asynchronous Barriers

Hopper enhances barrier functionality for TMA integration:

```cpp theme={null}
// Barrier initialization for TMA
__shared__ __align__(16) uint64_t barrier[1];

// Initialize barrier
if (threadIdx.x == 0) {
  uint32_t expected_bytes = tile_size_bytes;
  asm volatile(
    "mbarrier.init.shared.b64 [%0], %1;\n"
    :: "r"(cast_smem_ptr_to_uint(barrier)),
       "r"(expected_bytes));
}
__syncthreads();

// Issue TMA load with barrier
if (threadIdx.x == 0) {
  SM90_TMA_LOAD_2D::copy(
    desc_ptr, barrier, cache_hint,
    smem_ptr, coord0, coord1
  );
}

// Wait for TMA completion
asm volatile(
  "mbarrier.arrive_expect_tx.shared.b64 _, [%0], %1;\n"
  :: "r"(cast_smem_ptr_to_uint(barrier)),
     "r"(tile_size_bytes));
```

## Complete WGMMA+TMA Example

```cpp theme={null}
#include "cutlass/cutlass.h"
#include "cutlass/gemm/device/gemm_universal.h"

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
  cute::Shape<cute::_128, cute::_128, cute::_64>,  // TileShape MxNxK
  cutlass::gemm::collective::CollectiveMma<
    cutlass::gemm::TileScheduler,
    cute::Shape<cute::_128, cute::_128, cute::_64>,
    cutlass::half_t,    // ElementA
    cutlass::layout::RowMajor,
    cutlass::half_t,    // ElementB  
    cutlass::layout::ColumnMajor,
    float,              // ElementAccumulator
    cute::Shape<cute::_2, cute::_1, cute::_1>,  // ClusterShape
    cutlass::gemm::collective::StageCount<3>,
    cutlass::gemm::collective::KernelTmaWarpSpecialized  // Use WGMMA+TMA
  >,
  cutlass::epilogue::collective::DefaultEpilogue<...>
>;

using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

int main() {
  Gemm gemm_op;
  
  typename Gemm::Arguments args{
    cutlass::gemm::GemmUniversalMode::kGemm,
    {M, N, K},
    {A, lda, B, ldb},
    {{alpha, beta}, C, ldc, D, ldd}
  };
  
  // Kernel uses WGMMA instructions automatically
  cutlass::Status status = gemm_op(args);
  
  return status == cutlass::Status::kSuccess ? 0 : -1;
}
```

## Performance Optimization

### Optimal Tile Sizes for WGMMA

```cpp theme={null}
// Recommended threadblock shapes for Hopper
using TileShape_128x128x64  = cute::Shape<cute::_128, cute::_128, cute::_64>;
using TileShape_128x256x64  = cute::Shape<cute::_128, cute::_256, cute::_64>;
using TileShape_256x128x64  = cute::Shape<cute::_256, cute::_128, cute::_64>;
```

### Cluster Configuration

* **Small problems**: 1x1x1 cluster (single CTA)
* **Medium problems**: 2x1x1 or 1x2x1 cluster
* **Large problems**: 2x2x1 cluster (maximum practical)

### Pipeline Stages

* **WGMMA kernels**: 3-4 stages typical
* **Balance**: TMA latency vs shared memory capacity
* **Hopper H100**: Up to 227 KB shared memory per SM

## Compilation

```bash theme={null}
# Required: sm_90a for WGMMA and TMA
nvcc -arch=sm_90a -std=c++17 example.cu -o example

# CMake
cmake .. -DCUTLASS_NVCC_ARCHS="90a"

# Verify architecture-accelerated features
nvcc --ptx -arch=sm_90a example.cu -o example.ptx
grep "wgmma" example.ptx  # Should find wgmma instructions
```

<Warning>
  Using `sm_90` (without "a") will **not** enable WGMMA or TMA. Always use `sm_90a` for Hopper architecture-accelerated features.
</Warning>

## Examples

Hopper-specific examples in CUTLASS:

* `examples/48_hopper_warp_specialized_gemm/` - WGMMA-based GEMM
* `examples/57_hopper_warp_specialized_gemm_with_epilogue_visitor/` - Epilogue fusion
* `examples/58_hopper_persistent_gemm/` - Persistent kernel design
* `examples/111_hopper_ssd/` - State Space Decomposition

## See Also

* [Ampere Architecture](/architectures/ampere) - Previous generation
* [Blackwell Architecture](/architectures/blackwell) - Next generation
* [CuTe Documentation](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/00_quickstart.html)
* [NVIDIA Hopper Architecture Whitepaper](https://resources.nvidia.com/en-us-tensor-core)