---
license: apache-2.0
tags:
- cuda
- kernel
- gpu-optimization
- hpc
---

# hopper_tensor_core_bf16_async

Hopper BF16 async WMMA lane with cp.async double-buffered staging, 64x64x32 tiles, and four fragments per warp to cut exposed feed bubbles.

This repository contains the standalone CUDA source for the `hopper_tensor_core_bf16_async` lane from
 the PyC kernel lab. It is a source artifact for inspection and benchmarking;
it is not a precompiled binary and the result below is not a universal ranking.

## Performance

| Kernel | GPU / architecture | Shape | Best recorded result | Evidence |
|---|---|---|---|---|
| `hopper_tensor_core_bf16_async` | sm90 | 4096x4096x4096 | 1.059 ms / 129.824 TFLOPS | Measured on sm90, shape 4096x4096x4096; evidence `hopper-sm90-80-120-async-candidates-20260421T201129Z.json`. |

![Performance plot](performance.svg)

The result is reported with the original campaign's timing and correctness
context. Compare kernels only when GPU, CUDA version, matrix shape, warmup,
repeats, and reference/correctness mode match.

## Source

- `kernel.cu` — copied from `kernels/prototypes/hopper/tensor_core_async/kernel.cu`.
- Original lane tags: `cuda, matmul, hopper, sm90, prototype, tensor-core, bf16, async, cpasync, double-buffered`.

## Build/run contract

```text
{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}
{build_dir}/{name} 4096 4096 4096 3 30 1
```