NVFP4EffMA 0.6.0

NVFP4EffMA 0.6.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.6.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.6.0-linux-x86_64-cuda12.tar.xz (tar -xJf) and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.6.0-linux-x86_64-cuda12.tar.xz (12M)
• nvfp4effma-0.6.0-src.tar.gz (208K)
• nvfp4effma-0.6.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (75M)

What's in it

A transformer layer's epilogue in the call, and the calls' work as items.

• nvfp4a_gemm_ex, nvfp4a_gemv_tc_ex, nvfp4a_gemm_tc_ex extend the plain calls (B INT8 / BF16,
  A's scales linear E4M3) by any combination of: the deferred RMSNorm column scale (the parts of each
  column's sum of squares read through L2, one warp of the CTA reduces them and shares the scales), SwiGLU
  of interleaved gate / up rows, and a second weighted output with the parts of its sum of squares — the
  input of the next layer's norm. Fixed summation orders, no atomics, on the CUDA-core GEMV, the DP4A /
  FFMA tiles and the tensor-core paths alike.
• nvfp4effma_items.h: the same work as items for an executor of one's own — the number of items,
  their CTA shape and shared memory, per-item read / write ranges (with boundary fallbacks) and device
  entries; coherent intra-launch reads of mutable activations through L2, and the executor's publication
  protocol (waits / signals stay with the executor). Works from NVRTC too; the pairing of a compute_75
  library with a different-architecture executor is checked through cudaFuncAttributes::ptxVersion.
• Weight images on any GPU (nvfp4a_tc_image): the image — and the A16 ↔ A8 and image → tensor
  conversions — build wherever the weights are, not only where the tensor cores are.
• GEMV extended-epilogue performance: the NORM scales are computed by one warp while the others load
  the epilogue's inputs, and the INT8 kernels keep the plain call's k-loop schedule (a ptxas register
  allocation that had halved the loop's bandwidth is pinned by launch bounds): against the plain call,
  n = 8 tails cost ×1.0 … ×1.2 (was ×1.4 … ×1.9 on INT8), BF16 tiles ≈ ×1.0, tensor cores ×0.97 … ×1.23
  by mask (bench_epilogue, bench_gemm / bench_gemv with the new Extended epilogue sections;
  benches/1080.md and benches/l40s.md).

Tested by the author on a GTX 1080 (sm_61 and sm_60 builds: the full C suite, all architectures' ordinary
kernels byte-identical where nothing changed) and an L40S (sm_89 and compute_75: the full suite, items and
epilogue); the wheel installs and passes its pytest on both (the tensor-core tests skip below sm_75).

The wheel here is built for Python 3.12, PyTorch 2.7 and CUDA 12.6.

NVFP4EffMA 0.5.0

NVFP4EffMA 0.5.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.5.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.5.0-linux-x86_64-cuda12.tar.xz (tar -xJf) and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.5.0-linux-x86_64-cuda12.tar.xz (8,7M)
• nvfp4effma-0.5.0-src.tar.gz (168K)
• nvfp4effma-0.5.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (59M)

What's in it

Tensor cores for sm_75 … sm_89.

• nvfp4a_gemm_tc / nvfp4a_gemv_tc: W4A4 and W4A8 on INT8 mma with the same exact block sums as the DP4A kernel (one mma per block of 16, the 1.5 · 2²³ accumulator trick), W4A16 on BF16 mma with the weight decoded to a · scale exactly (sm_80+). FP32 accumulation throughout; the order is the tensor cores', so the values match nvfp4a_gemm within the rounding of an FP32 accumulation, not bit for bit.
• One image of the weight (nvfp4a_tc_image, kinds NVFP4A_TC_A16 for BF16 activations and NVFP4A_TC_A8 for INT8 / NVFP4), built once from any NVFP4 variant, serves both the GEMM and the GEMV; activations go through NVFP4A_TC_B per GEMM call, the GEMV takes them as they are.
• GEMM: persistent CTAs with stream-K for the last wave (nvfp4a_tc_workspace_bytes, zeroed once), rastered tiles; on sm_80+ a cp.async copier and an mbarrier pipeline, ldmatrix for BF16 activations; on sm_75 producer warps and sequence flags.
• GEMV (1 … 8 columns): the weight's image streamed from DRAM once in 16-byte loads (L1 bypassed, evict-first in L2), the k sums kept in registers to the end, one store per output.
• PyTorch: TCWeight(weight, act="exact" | "int8" | "nvfp4") holds the image; linear(x, tw) runs the tensor-core GEMV up to 8 rows and the GEMM beyond.
• NVFP4A_DP4A=1 in the environment turns the tensor-core paths off. nvfp4a_gemm and linear on a plain NVFP4Tensor are unchanged (DP4A / FFMA).

Tested by the author on a GTX 1080 (DP4A / FFMA paths, as before) and an L40S (sm_89): the C tests check every GEMM / GEMV case also through the tensor-core functions against a float64 reference, and make test-sm75 runs the Turing variant on the L40S from compute_75 PTX. The tensor-core GEMM's integer paths keep a few values in local memory at k-tile boundaries; every other kernel has none (Jetson TX1 / TX2 as before).

The wheel here is built for Python 3.12, PyTorch 2.7 and CUDA 12.6.

NVFP4EffMA 0.4.0

NVFP4EffMA 0.4.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.4.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.4.0-linux-x86_64-cuda12.tar.xz (tar -xJf) and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.4.0-linux-x86_64-cuda12.tar.xz (8,6M)
• nvfp4effma-0.4.0-src.tar.gz (136K)
• nvfp4effma-0.4.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (57M)

What's in it

For code that builds on NVFP4EffMA:

• Int8Tensor and quantize_int8(x): the int8 activations of linear(..., act="int8") (one FP32 scale per 16 values) as a public type. linear takes an Int8Tensor as x (bit for bit act="int8"), so quantized activations can be kept for a backward pass without reaching into the extension; act="int8" itself now goes through it.
• get_include(): the wheel ships nvfp4effma.h and the device-side codecs nvfp4effma.cuh (E2M1 / E4M3 decode and RNE encode, scale addressing), so CUDA extensions can build against the installed package. The C archive carries them too (src/).

Kernels are unchanged from 0.3.0.

NVFP4EffMA 0.3.0

NVFP4EffMA 0.3.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.3.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.3.0-linux-x86_64-cuda12.tar.xz (tar -xJf) and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.3.0-linux-x86_64-cuda12.tar.xz (8,5M)
• nvfp4effma-0.3.0-src.tar.gz (136K)
• nvfp4effma-0.3.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (57M)

What's new

• Every NVIDIA NVFP4 variant, read and produced as it is, bit for bit with Transformer Engine's arithmetic: ModelOpt checkpoints (1 × 16, decode factor), Transformer Engine's 1 × 16 and 16 × 16 (its default for weights) recipes, amax globals per tensor or per row (row-scaled activations), 4over6 (map-to-4 / map-to-6, MAE / MSE, E4M3 bound 448 or 256), UE5M3 block scales, rowwise and columnwise data, and the GEMM-swizzled scale layout of cuBLASLt — in GEMM, GEMV, quantization and dequantization, A and B in any combination. Not supported: the random Hadamard transform, stochastic rounding and 4over6's fast-math error mode.
• C API: nvfp4a_tensor gains flags (NVFP4A_SCALES_SWIZZLED, NVFP4A_SCALES_UE5M3, NVFP4A_GLOBAL_AMAX, NVFP4A_GLOBAL_PER_ROW) and scale_max; nvfp4a_quantize_ex (rowwise / columnwise, 1 × 16 / 16 × 16, 4over6, amax computed), nvfp4a_gemm_dual (forward and backward GEMMs from rowwise / columnwise halves), nvfp4a_swizzle_scales, nvfp4a_transpose (16 × 16 tensors).
• PyTorch: quantize(x, block=, four_over_six=, scale_max=, row_scaled=, columnwise=, swizzled=, scale_dtype=) with Transformer Engine's current scaling by default (the amax is computed and stored); NVFP4Tensor holds every layout, NVFP4Tensor.from_te wraps a Transformer Engine NVFP4 tensor without copying, .with_scale_layout(), .transpose_2d(); linear takes any formats on either side.
• No cost for the formats in the kernels: block scale formats are compile-time kernel instances; the GEMM rows of 0.2.0 are unchanged or slightly faster (NVFP4 × NVFP4 9.42 ms, W4A8 9.33 ms at 4096³), the Transformer Engine formats within 1.5 % of plain NVFP4 in GEMM and −6.5 … +3 % in GEMV (swizzled x scales are faster). Creating them costs more — quantize once, not per call (docs/performance.md).
• Quantization: 16 × 16 tiles at 0.74 ms (BF16, 8192²), 4over6 at 1.9 ms with the division by 6·M computed exactly through its reciprocal (exhaustively checked), the amax pass vectorized.
• Benchmarks: a Transformer Engine block in every benchmark (each row against the plain format of the same run), a warm-up before the first row, NVFP4A_BENCH_FILTER to run selected rows.

Changes in behaviour. Quantization now follows Transformer Engine's multiplications (x · inv, bmax · S_enc / 6) instead of 0.2.0's divisions, which rarely differ by one rounding step: some blocks get a different code than with 0.2.0. Python quantize(x) without a global_scale now stores an amax global (Transformer Engine's convention) instead of the decode factor amax / (6 · 448); values are the same. The C struct nvfp4a_tensor has two more trailing fields (source-compatible: initializers without them mean the ModelOpt format; not ABI-compatible).

The C tarball is now .tar.xz (8.6 MB; the format instances compress well, gzip would be 53 MB).

Tested and tuned by the author on a GTX 1080 (sm_61) only, plus the code path of GPUs without DP4A (SASS built for sm_60, run on the GTX 1080). Other architectures are compiled, not run. No stack or local memory in any kernel, except on Jetson TX1 / TX2 (sm_53 / sm_62, 32 K registers per block), where the heaviest kernels spill.

The wheel here is built for Python 3.12, PyTorch 2.7 and CUDA 12.6.

NVFP4EffMA 0.2.0

NVFP4EffMA 0.2.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.2.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.2.0-linux-x86_64-cuda12.tar.gz and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.2.0-linux-x86_64-cuda12.tar.gz (5.6M)
• nvfp4effma-0.2.0-src.tar.gz (96K)
• nvfp4effma-0.2.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (8.4M)

What's new

• W4A8: NVFP4 weights × int8 activations with an FP32 scale per 16 values (NVFP4A_INT8; nvfp4a_quantize produces them). Exact integer block sums on DP4A (|p| ≤ 24384), one rounding per block, the same bits on the FFMA path below sm_61. 9.3–9.4 ms at 4096³ (7.3 TMAC/s, 2.1× cuBLAS SGEMM) — the accuracy class of llama.cpp's int8 activations.
• Small and medium batches: the GEMV kernel now also runs beyond 8 columns, in groups of 8, up to where the tiled kernel takes over (48–56 columns, 32 for FP32); a 64-wide DP4A tile (two blocks per SM) covers 57–64 and 129–192 columns.
• PyTorch: linear(..., act=) with "auto" (default: exact for up to 2 rows, int8 activations beyond, about twice as fast), "exact", "int8", "nvfp4"; the same option on NVFP4Linear / quantize_linear_.
• Against llama.cpp's NVFP4 kernels on the same GTX 1080 (interleaved runs, int8 activations on both sides): faster on every shape measured — 1.47× at 4096³, 3.9–4.5× for 8-token decoding, 1.45–1.66× at 512 tokens. Table and method: docs/performance.md.
• GEMM candidates with 128-thread blocks no longer spill (register budget of 4 × 128 threads).
• The GEMV kernel's NVFP4-activation path computes the exact block sums of the tiled kernels (DP4A per block), 2–3× faster for 4–8 columns.

API change: nvfp4a_dtype gains NVFP4A_INT8 = 4; ld_scale of INT8 operands counts floats. Everything else is source-compatible with 0.1.0.

Tested and tuned by the author on a GTX 1080 (sm_61) only, plus the code path of GPUs without DP4A (SASS built for sm_60, run on the GTX 1080). Other architectures are compiled, not run. No stack or local memory in any kernel, except on Jetson TX1 / TX2 (sm_53 / sm_62, 32 K registers per block), where the heaviest kernels spill.

The wheel here is built for Python 3.12, PyTorch 2.7 and CUDA 12.6.

NVFP4EffMA 0.1.0

NVFP4EffMA 0.1.0

NVFP4 GEMM, GEMV and quantization for NVIDIA GPUs without native FP4 (sm_50 … sm_90).

Install

PyTorch: pick the wheel matching your Python (cpXY), PyTorch (torchA.B) and its CUDA (cuNNN):

pip install nvfp4effma-0.1.0+<cuNNNtorchA.B>-<cpXY>-<cpXY>-linux_x86_64.whl

Other combinations: pip install ./binds/python --no-build-isolation from the source archive.

C / C++: unpack nvfp4effma-0.1.0-linux-x86_64-cuda12.tar.gz and link lib/libnvfp4effma.a (see its README).

Every artifact carries SASS for sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75, sm_80, sm_86, sm_87, sm_89, sm_90 (CUDA 12.0).

Assets

• nvfp4effma-0.1.0-linux-x86_64-cuda12.tar.gz (3.9M)
• nvfp4effma-0.1.0-src.tar.gz (88K)
• nvfp4effma-0.1.0+cu126torch2.7-cp312-cp312-linux_x86_64.whl (6.7M)

What's in it

First release. The NVFP4 sibling of E4M3EffMA and
E5M2EffMA, with the same build, tuning, test and benchmark setup.

• GEMM D = α·A·Bᵀ + bias + β·C: A NVFP4 (E2M1 values, E4M3 scale per 16 values, FP32 global scale), B ∈ {NVFP4, BF16, FP16, FP32}, D ∈ {BF16, FP16, FP32}, accumulation into D (e.g. an FP32 main_grad), amax of the result.
• NVFP4 × NVFP4 (W4A4) on DP4A from sm_61: twice an E2M1 value is an exact int8, so block sums of 16 are exact integers and each block is rounded once. Below sm_61 (Maxwell, Tesla P100) an FFMA kernel produces the same bits.
• NVFP4 × BF16 / FP16 / FP32 (W4A16): the block scale is applied when the weights are decoded; exact products, one FFMA per MAC.
• GEMV path for up to 8 columns (decoding), picked automatically; DP4A per block for NVFP4 activations.
• Quantization / dequantization from and to BF16, FP16 and FP32, with the rounding of ModelOpt / Transformer Engine and the checkpoint layout (weight, weight_scale, weight_scale_2).
• PyTorch: quantize / dequantize / linear, NVFP4Tensor, and NVFP4Linear with quantize_linear_ to swap the nn.Linear layers of a model. Imports alongside e4m3effma / e5m2effma.

On a GTX 1080: NVFP4 × NVFP4 GEMM 2.01× cuBLAS SGEMM (9.72 ms at 4096³, 7.1 TMAC/s; 1.98× when accumulating into FP32), NVFP4 × BF16 at SGEMM speed with 7× fewer weight bytes, GEMV n = 1 at 84 % of DRAM bandwidth (7.4× cuBLAS SGEMV), quantization at 89 %. Details: docs/performance.md, benches/.

Tested and tuned by the author on a GTX 1080 (sm_61) only, plus the code path of GPUs without DP4A (SASS built for sm_60, run on the GTX 1080). Other architectures are compiled, not run. No stack or local memory in any kernel, except on Jetson TX1 / TX2 (sm_53 / sm_62, 32 K registers per block), where the heaviest kernels spill.

The wheel here is built for Python 3.12, PyTorch 2.7 and CUDA 12.6.