# I2_S, TL1, and TL2 Kernel Implementations in BitNet: GPU vs CPU Architecture Deep Dive

> Explore BitNet's I2_S, TL1, and TL2 kernel implementations. Discover GPU vs CPU architecture differences for optimized performance and hardware utilization.

- Repository: [Microsoft/BitNet](https://github.com/microsoft/BitNet)
- Tags: deep-dive
- Published: 2026-03-13

---

**BitNet employs three distinct quantization kernels—**I2_S TLI** for NVIDIA GPUs using inline PTX bit-manipulation, **TL2** for x86 CPUs with AVX2 lookup-table acceleration, and **TL1** for ARM GPUs and CPU fallback scenarios—each optimized for specific hardware constraints and memory layouts.**

The microsoft/BitNet repository implements ultra-low-bit inference through these specialized compute paths. Understanding the architectural differences between **I2_S**, **TL1**, and **TL2** kernel implementations in BitNet is essential for deploying 2-bit quantized models across heterogeneous hardware environments, from datacenter GPUs to edge CPUs.

## I2_S TLI Kernel: NVIDIA GPU Implementation

The **I2_S TLI** (int-2-bit-s transpose-like interleaved) kernel is the primary compute path for NVIDIA GPUs, designed to unpack and multiply 2-bit weights directly within CUDA warps.

### Target Hardware and Activation

This kernel targets NVIDIA GPUs utilizing CUDA compute capabilities. It is distinct from the ARM path, which uses **TL1** instead. The implementation resides in [`gpu/bitnet_kernels.h`](https://github.com/microsoft/BitNet/blob/main/gpu/bitnet_kernels.h), where two critical functions handle the quantization pipeline.

### Core Functions and Data Layout

The kernel operates on the **I2S** layout, which stores 2-bit values in a transpose-like order: `{e0, e4, e8, …}`. Two functions manage the decoding and multiplication:

- **`decode_i2s_to_i8s`** (lines 23-44 in [`gpu/bitnet_kernels.h`](https://github.com/microsoft/BitNet/blob/main/gpu/bitnet_kernels.h)): Unpacks 16 packed 2-bit values into 16 signed 8-bit integers using CUDA inline PTX assembly.
- **`ladder_int8xint2_kernel`** (lines 46-73 in [`gpu/bitnet_kernels.h`](https://github.com/microsoft/BitNet/blob/main/gpu/bitnet_kernels.h)): Executes the general matrix multiplication (GEMM) between int8 activations and unpacked int2 weights.

### Bit-Level Optimization

The implementation leverages **inline PTX assembly** for maximum throughput. The `decode_i2s_to_i8s` function uses `lop3.b32` logic operations and `__vsubss4` vector subtraction to perform bit extraction and sign extension without branching, enabling single-pass decoding directly in registers.

```cpp
// CUDA kernel invocation for I2_S TLI path
ladder_int8xint2_kernel<<<grid, block>>>(
    A,               // int8_t* input activations
    B,               // int8_t* packed I2S weights
    dtype_transform, // output buffer (typically bfloat16)
    s,               // scaling factors
    ws);             // workspace pointer

```

## TL2 Kernel: x86 AVX2 Lookup Table Implementation

The **TL2** kernel provides a high-throughput path for x86 CPUs with **AVX2** or higher instruction sets, utilizing precomputed lookup tables (LUTs) and tiled matrix multiplication strategies.

### Files and API Surface

The TL2 implementation spans multiple files:
- **[`include/ggml-bitnet.h`](https://github.com/microsoft/BitNet/blob/main/include/ggml-bitnet.h)** (lines 42-45): Declares the public API including `ggml_qgemm_lut` and `ggml_preprocessor`
- **[`preset_kernels/bitnet_b1_58-large/bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/preset_kernels/bitnet_b1_58-large/bitnet-lut-kernels-tl2.h)**: Contains AVX2 intrinsic implementations for LUT construction

### Block Partitioning and Constraints

TL2 imposes a strict block size constraint: **BK must satisfy `BK % 6 == 0`** (documented in [`docs/codegen.md`](https://github.com/microsoft/BitNet/blob/main/docs/codegen.md) lines 43-45). To handle standard dimensions, the kernel splits the **K** dimension into two partitions:
- **threeK**: Blocks where `BK = 32` and divisible by 6, processed by TL2
- **twoK**: Remainder blocks that do not satisfy the modulo constraint, falling back to **TL1**

### LUT Generation with AVX2

The kernel constructs three separate lookup tables using AVX2 intrinsics (`_mm256_*`):
- **`three_lut_ctor`** (lines 71-89 in [`bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/bitnet-lut-kernels-tl2.h)): Generates the primary LUT for threeK blocks
- **`two_lut_ctor`** (lines 98-115 in [`bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/bitnet-lut-kernels-tl2.h)): Generates the secondary LUT for twoK blocks (though these typically execute via TL1 fallback)

```cpp
// TL2 LUT construction for threeK blocks (BK % 6 == 0)
three_lut_ctor<ACT_K>(qlut, b, lut_scales);

// Fallback to TL1 for twoK remainder blocks
two_lut_ctor<ACT_K>(qlut, b, lut_scales); // Typically routes to TL1 execution

```

### Compilation Control

Enable TL2 via the CMake option **`BITNET_X86_TL2`** (defined in [`CMakeLists.txt`](https://github.com/microsoft/BitNet/blob/main/CMakeLists.txt) lines 16-31), which sets `-DGGML_BITNET_X86_TL2` during compilation.

## TL1 Kernel: ARM GPU and CPU Fallback

The **TL1** kernel serves dual roles: as the primary compute path for **ARM-based GPUs** and as the fallback mechanism for x86 CPUs when TL2 constraints cannot be satisfied.

### ARM GPU Activation

When compiling BitNet for ARM GPU architectures (as opposed to NVIDIA), enable TL1 using the **`-DGGML_BITNET_ARM_TL1`** flag. This replaces the I2_S TLI path entirely for non-NVIDIA hardware.

### CPU Fallback Mechanism

On x86 systems, TL1 handles the **twoK** partition when the K dimension does not divide evenly into TL2-compatible blocks. When `BK % 6 != 0`, the `ggml_bitnet_mul_mat_task_compute` function routes the remainder tiles to TL1 instead of the AVX2-optimized TL2 path.

## Architectural Comparison

| Feature | I2_S TLI | TL2 | TL1 |
|---------|----------|-----|-----|
| **Primary Target** | NVIDIA GPUs (CUDA) | x86 CPUs (AVX2) | ARM GPUs / CPU fallback |
| **Key Files** | [`gpu/bitnet_kernels.h`](https://github.com/microsoft/BitNet/blob/main/gpu/bitnet_kernels.h) | [`include/ggml-bitnet.h`](https://github.com/microsoft/BitNet/blob/main/include/ggml-bitnet.h), [`preset_kernels/.../bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/preset_kernels/.../bitnet-lut-kernels-tl2.h) | ARM: CUDA compilation path; CPU: generic fallback |
| **Data Strategy** | On-the-fly unpacking with PTX | Precomputed LUTs with tiling | Standard quantization paths |
| **Block Handling** | Single-pass, no splitting | Splits K into threeK/twoK | Handles twoK remainders |
| **Instruction Set** | `lop3.b32`, `__vsubss4` | `_mm256_*` intrinsics | Architecture-dependent |
| **CMake Flag** | Default for CUDA | `BITNET_X86_TL2` | `GGML_BITNET_ARM_TL1` |

## Practical Implementation Examples

### GPU Inference with I2_S TLI

For NVIDIA hardware, directly invoke the CUDA kernels after allocating device memory for the packed I2S weights:

```cpp
// Device pointers must be allocated with cudaMalloc
int8_t* d_A; // int8 activations
int8_t* d_B; // packed I2S weights
void* d_output; // bfloat16 or fp16 output

// Launch configuration depends on matrix dimensions
ladder_int8xint2_kernel<<<grid_dim, block_dim>>>(
    d_A, d_B, d_output, scale_factors, workspace);

```

### CPU Inference with TL2

On x86 systems with AVX2, the high-level `ggml` API automatically selects TL2 when the compile flag is enabled:

```cpp
// Prepare ggml tensors
ggml_tensor* src0 = ggml_new_tensor_2d(ctx, GGML_TYPE_Q4_0, K, M); // Activations
ggml_tensor* src1 = ggml_new_tensor_2d(ctx, GGML_TYPE_I2_S, N, K); // BitNet weights
ggml_tensor* dst  = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, N, M);  // Output

// When GGML_BITNET_X86_TL2 is defined, this calls ggml_qgemm_lut internally
ggml_bitnet_mul_mat_task_compute(src0->data, src1->data, dst->data, ...);

```

### Handling Mixed Block Sizes

When implementing custom kernels that split K dimensions, explicitly handle the TL2 constraint:

```cpp
const int BK = 32;
const int threeK = (K / BK) * BK; // TL2 handles this portion
const int twoK = K - threeK;      // Falls back to TL1

// Process threeK with TL2 AVX2 kernels
three_lut_ctor<32>(qlut_three, weights, scales);

// Process twoK with TL1 generic path
if (twoK > 0) {
    // TL1 fallback execution
    process_twoK_remainder(weights + threeK, twoK);
}

```

## Summary

- **I2_S TLI** targets NVIDIA GPUs exclusively, using inline PTX assembly (`lop3.b32`) to decode 2-bit weights on-the-fly without lookup tables.
- **TL2** accelerates x86 CPUs via AVX2 intrinsics and precomputed LUTs, requiring block sizes where `BK % 6 == 0` and splitting K into threeK/twoK partitions.
- **TL1** serves ARM GPUs when compiled with `-DGGML_BITNET_ARM_TL1` and handles CPU fallback for remainder blocks that violate TL2 alignment constraints.
- **File locations** determine the implementation path: [`gpu/bitnet_kernels.h`](https://github.com/microsoft/BitNet/blob/main/gpu/bitnet_kernels.h) for I2_S, [`preset_kernels/.../bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/preset_kernels/.../bitnet-lut-kernels-tl2.h) for TL2, and ARM-specific paths for TL1.
- **Compilation flags** (`BITNET_X86_TL2` vs `GGML_BITNET_ARM_TL1`) are mutually exclusive indicators of the target hardware platform.

## Frequently Asked Questions

### What hardware requires I2_S TLI versus TL1?

**I2_S TLI** is required for NVIDIA GPUs to achieve optimal bit-manipulation performance using CUDA PTX instructions. **TL1** is mandatory for ARM-based GPU architectures (such as Mali or Adreno) when running BitNet inference, controlled by the `-DGGML_BITNET_ARM_TL1` compilation flag. On x86 CPUs, neither is primary; instead, **TL2** provides the optimized path with AVX2 acceleration.

### Why does TL2 split the K dimension into threeK and twoK?

The TL2 kernel in [`preset_kernels/bitnet_b1_58-large/bitnet-lut-kernels-tl2.h`](https://github.com/microsoft/BitNet/blob/main/preset_kernels/bitnet_b1_58-large/bitnet-lut-kernels-tl2.h) requires block sizes satisfying `BK % 6 == 0` to align with its AVX2 lookup table strategy. When processing matrices where **K** is not evenly divisible by 6, BitNet splits the dimension into **threeK** (the portion handled by TL2) and **twoK** (the remainder handled by TL1 fallback), as documented in [`docs/codegen.md`](https://github.com/microsoft/BitNet/blob/main/docs/codegen.md) lines 43-45.

### When should I enable TL1 compilation flags on x86 systems?

You do not explicitly enable TL1 for x86 via compile flags; TL1 serves automatically as the **fallback path** for TL2 when block constraints fail. However, if compiling for **ARM GPUs**, you must explicitly define `-DGGML_BITNET_ARM_TL1` to replace the default NVIDIA-centric I2_S TLI implementation with the ARM-compatible TL1 kernels.

### Which kernel delivers the best performance for BitNet inference on modern CPUs?

**TL2** delivers superior throughput on x86 CPUs with AVX2 support due to its lookup-table approach and tiled matrix multiplication using `_mm256_*` intrinsics. According to the source in [`include/ggml-bitnet.h`](https://github.com/microsoft/BitNet/blob/main/include/ggml-bitnet.h), the `ggml_qgemm_lut` function bypasses per-element decoding overhead by precomputing multiplication results. TL1 should only activate for edge-case remainders where `BK % 6 != 0`, making TL2 the primary performance path for production x86 deployments.