HPC-Ops: Tencent's CUDA Kernel Library for LLM Inference Hot Paths

The serving bottleneck HPC-Ops was built to attack

LLM serving latency is not spread evenly across a model. A small set of operations repeats on every decode step and every prefill batch: attention, the GEMMs inside MoE layers, sampling over the vocabulary, normalization, and the AllReduce that tensor parallelism inserts between layers. The README states that HPC-Ops focuses on exactly these hot paths, and lists Attention, MoE, GEMM, sampling, normalization, and communication-compute fusion as its scope. The intended user is an inference infrastructure engineer who already has a serving stack and wants to replace or supplement individual kernels inside it, not someone assembling a model runner from scratch. The README also positions the code as a CUDA tutorial: it says the repository contains examples of building kernels with CUDA, CuTe, CUTLASS, cp.async, TMA, PDL, and multicast. That dual purpose matters when you decide whether to read it or depend on it. A library written partly as a teaching artifact tends to expose its internals, which helps if you plan to fork a kernel, and hurts if you want a stable ABI.

How the kernels are organized around specific workload shapes

The June 2026 update notes describe five mechanisms in enough detail to reason about. Dynamic decode attention addresses mixed-length batches: the README states that static split-k scheduling cannot adapt when requests in one batch have very different KV-cache lengths, so the library splits requests into uniform KV tiles, assigns tiles before each decode step, and balances them across CTAs with a greedy bin-packing strategy. A combine kernel merges the split-k results. Sparse attention takes a precomputed block mask, skips masked-out KV tiles, and applies per-tile FP8 scaling. Route GEMM handles the case where activations are BF16 but weights need FP32 sensitivity: the implementation decomposes each FP32 weight into a high BF16 component and a low BF16 residual with a fixed scale of 1 / 256, then computes two BF16 Tensor Core GEMMs inside one kernel and writes the result once. Fused MoE reads original tokens directly through routing indices instead of gathering them first, and the README says the cp.async GEMM path removes Warp Specialization in the low-latency regime to raise CTA residency. Fused AllReduce plus RMSNorm has two modes: CUDA multicast for large-token prefill shapes, and a Lamport P2P two-kernel design with PDL overlap for small-token decode. The common thread is that each kernel is designed around a shape regime, not a single general case.

What the README does not tell you about getting it running

This is the weakest part of the material. The README has a Quick Start section in its table of contents, but the supplied text is truncated before that section, so no build command, no pip install line, no CMake invocation, and no environment variable is visible. The badges state CUDA 12.9 or newer, Python 3.8 or newer, and a GPU target of H20 or SM90. Those three constraints are the only concrete setup facts available. There is no retrieved release, so there is no version tag to pin. The licence badge says MIT while the repository metadata reports NOASSERTION, which means GitHub could not identify a licence file. Until you open the repository and read the actual LICENSE file, you cannot treat MIT as confirmed. The README does say the project exposes a Python API intended for integration into frameworks like vLLM and SGLang, and that tests and benchmarks are included, but it does not name the import path, the module name, or the function signatures. If you are evaluating this for production, the first hour of work is reading the repository tree, not reading the README.

H20 and SM90 tuning is a constraint, not a footnote

The GPU badge names H20 and SM90, and the README repeats that the kernels are tailored for NVIDIA H20 GPUs. That is narrower than it first appears. H20 is a China-market part with different compute-to-bandwidth ratios than the H100 or H200, and a kernel tuned for one balance of memory bandwidth against Tensor Core throughput does not automatically transfer to the other. The fused AllReduce plus RMSNorm path depends on NVLink-native communication and CUDA multicast, which assumes a topology where multicast is available. The README does not state which GPU generations outside SM90 are supported, nor whether the kernels degrade gracefully or fail to compile elsewhere. If your fleet is mixed, or if you run on A100, L40S, or AMD hardware, this library is not the right tool and the README gives you no path to make it one. Even within NVIDIA hardware, the absence of any stated fallback means the safe assumption is that the kernels are validated only on the documented target.

Where HPC-Ops overlaps with FlashInfer and vLLM's kernel layer

The README names its own comparison set: vLLM, SGLang, FlashInfer, NCCL, cuBLAS, and TensorRT-LLM. The most direct alternative is FlashInfer, which is also a standalone kernel library for LLM serving rather than a full inference engine. The difference in approach is visible in the update notes. FlashInfer's public material centers on a catalog of attention and sampling kernels with a JIT compilation path that generates variants for the shapes you actually run. HPC-Ops instead ships hand-written kernels with workload-specific scheduling built in, such as the greedy bin-packing task map for mixed-length decode and the two-shot communication schedule in the fused AllReduce path. That design can win on the shapes it was tuned for and lose flexibility everywhere else, because adapting it means editing CUDA rather than regenerating a variant. Against vLLM and SGLang, the relationship is different: those are full serving engines, and the README positions HPC-Ops as something you integrate into them. Against NCCL and cuBLAS, the pitch is fusion, since HPC-Ops merges AllReduce with RMSNorm and residual add where NCCL would leave them as separate stages.

Maintenance cost and the licence question

There are no retrieved releases, so there is no upgrade cadence to plan around. The README is organized as dated update notes, with a June 2026 entry covering five kernel families, which suggests development happens in the open repository rather than through tagged versions. For a library that lives inside your serving path, that means tracking main and rebuilding against CUDA 12.9 or newer on your own schedule. The kernels use CuTe, CUTLASS, cp.async, TMA, PDL, and multicast, so a CUDA toolkit upgrade is a real maintenance event, not a version bump. On licensing, the badge claims MIT but the repository metadata says NOASSERTION. MIT would permit commercial use and modification with attribution, but you should not rely on a badge. Read the LICENSE file in the repository root and confirm it before you copy any kernel into a product. This is not legal advice, and if the licence file is missing or ambiguous, that is a question for whoever handles licensing at your organization.

Who should read this repository first

The README describes the code as a modern CUDA tutorial alongside a production library, and that framing is the most useful thing in it. If you are writing your own attention or MoE kernels and want to see how PDL chaining, cp.async pipelines, or per-tile FP8 scaling are handled in compact implementations, the repository is worth reading even if you never depend on it. If you are running vLLM or SGLang on H20 hardware and your profiling points at attention tail latency in mixed-length decode batches or at the AllReduce and RMSNorm sequence, the mechanisms described here map directly onto those problems. If you are on any other accelerator, or if your procurement process requires a tagged release and a confirmed licence before evaluation, this is not ready for you yet. The gap between the README's claims and what it documents is the thing to weigh: five kernel families with real scheduling detail, and no install command, no release, and an unresolved licence identifier.

Editorial conclusion

HPC-Ops is worth evaluating if you serve LLMs on NVIDIA H20 or SM90 hardware, run vLLM or SGLang, and need FP8 or BF16 kernels for attention, MoE, GEMM or fused AllReduce plus RMSNorm that you can read and modify. It is the wrong starting point if you are on non-NVIDIA accelerators, on GPUs older than the documented SM90 target, or if you need a tagged release and a clear licence before you can ship anything. Before adopting, verify three things in the repository itself: the actual licence file, whether a build or install path exists beyond the Python API description, and whether the operator catalog covers the specific kernel your serving stack spends its time in.

Official sources

1. Issues
2. README
3. Tencent/hpc-ops on GitHub