File size: 1,527 Bytes
7d86ec7 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 | ---
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`. |
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
```
|