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 fromkernels/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
{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
- Downloads last month
- -
cuda
kernel
gpu-optimization
hpc
apache-2.0