English | 简体中文
HALOps (High-performance Accelerated LLM Operators) is a high-performance operator library for large language model inference, developed by the ERNIE team. The operators are built to exploit NVIDIA GPU hardware features, and use PDL, multi-stream overlap, operator fusion and MegaKernels to remove bubbles between kernels, keeping compute, memory and communication as close to saturation as possible for high-throughput, low-latency inference.
Highlights | Performance | Hardware and Software Requirements | Quick Start | Documentation | Build Guide | Benchmark Guide | Contributing
HALOps is designed from the ground up for NVIDIA Blackwell (SM100). Built on CUTLASS / CuTe, it uses tcgen05, TMEM and TMA directly and covers the Attention, MoE and MLA hot paths of LLM inference.
- GQA attention: variable-length prefill and paged-KV decode across BF16 / FP8 KV cache, GQA and speculative decoding; split-KV parallelism, adaptive reduction and fused pre/post-processing make full use of SM100 compute and bandwidth in long-context and small-batch decode.
- Multi-precision GEMM: FP8 (MXFP8) and FP4 (NVFP4 / MXFP4) block-scaled GEMM for dense, batched and MoE grouped workloads; with ultra-lightweight flexible tile configurations and pre-mapped tile scheduling, it delivers industry-leading performance in most scenarios on SM100.
- MegaKernel sparse MLA: high-performance MegaKernel sparse attention for HCA and SWA, removing kernel scheduling and intermediate memory round trips, with tile-level pipelining that hides compute and memory latency to push GPU utilization further.
All results are measured on SM100: MXFP8 dense GEMM outperforms DeepGEMM and FlashInfer in some cases, and MegaMLA consistently outperforms the multi-stream + PDL implementation at small concurrency.
 |
 |
| Figure 1: Dense Mxfp8 W8A8 GEMM performance comparison | Figure 2: HCA attention layer performance comparison |
| Component | Requirement |
|---|---|
| GPU | NVIDIA Blackwell SM100 GPU |
| CUDA Toolkit | 13.0 or later, hardware architecture compute_100f/sm_100f |
| NVIDIA Driver | R580 or later, compatible with CUDA 13.x |
| Python / PyTorch | Python 3.10+; PyTorch 2.13+ |
See the NVIDIA CUDA compatibility matrix for driver requirements, and the Build Guide for the full source-build requirements.
Prebuilt wheels are available for Python 3.12 with CUDA 13.2 on Linux x86_64 and aarch64. With a CUDA 13.2 build of PyTorch installed, install the prebuilt HALOps package directly:
python -m pip install halops==0.1.0 -f https://www.paddlepaddle.org.cn/packages/stable/cu132/halops/For other Python or CUDA versions, build from source.
To build it yourself, run from the repository root:
git submodule update --init --recursive
python -m pip install -r requirements.txt
python -m pip install --no-build-isolation --no-deps -v .See the Build Guide for development installs, wheel builds, CUDA path configuration and build parallelism.
After installation, the W8A8-MXFP8 dense GEMM example below checks the quantization and GEMM path. It should print torch.Size([128, 256]) torch.bfloat16.
import torch
from halops import dense_gemm, group_wise_quant
M, N, K = 128, 256, 512
a = torch.randn((M, K), device="cuda", dtype=torch.bfloat16)
b = torch.randn((N, K), device="cuda", dtype=torch.bfloat16)
a_fp8, scale_a = group_wise_quant(a, mode="w8a8-mxfp8", use_swizzle=True)
b_fp8, scale_b = group_wise_quant(b, mode="w8a8-mxfp8", use_swizzle=True)
output = torch.empty((M, N), device="cuda", dtype=torch.bfloat16)
dense_gemm(a_fp8, b_fp8, scale_a, scale_b, output, mode="w8a8-mxfp8")
print(output.shape, output.dtype)- Complete GEMM coverage: add W4A8, BF16 and other GEMM types to build a full quantization-type support matrix, so users can pick the trade-off between inference speed and accuracy they need.
- Higher-performance MegaKernels: cover more compute-communication fusion scenarios and further overlap compute with memory traffic.
- High-performance kernels for small workloads: kernels tuned for small problem sizes that deliver the lowest latency at low concurrency.
See CONTRIBUTING.md for the development setup, code quality and testing requirements, and the pull request checklist.
HALOps is built on CUTLASS and CuTe; we thank the CUTLASS developers for their work. We are also grateful to FlashInfer, FlashMLA and DeepGEMM, whose prior work on LLM inference kernels inspired many of our designs and served as valuable references for validating them.
HALOps is released under the Apache License 2.0.