Shared Memory Bank Conflicts: Why Your Quantized Kernel Is Slow

Pack eight INT4 weights into a 32-bit word, write a straightforward unpack-and-store loop to stage them in shared memory for the matmul that follows, and profile it. If the unpack order lines every thread in a warp up on the same bank, Nsight Compute will show a kernel burning up to 32x the shared memory traffic it should, and nothing in the source code looks obviously wrong. That's what shared memory bank conflicts look like from the inside: correct output, quietly serialized hardware cycles, and packing low-bit weights into wider words is one of the easiest ways to trigger one by accident in a kernel that was supposed to make inference cheaper, not slower.

TL;DR: Why Shared Memory Bank Conflicts Slow Down Quantized Kernels

• Quantization link: packing INT4/INT2 weights changes the per-element stride, reconstructing that pattern in a dequant kernel.
• Fix: pad a tile by one element, or XOR-swizzle it (MARLIN's approach); both cut a measured RTX 3090 transpose from 1.10ms to 0.92ms.
• Detection: Nsight Compute's l1tex bank-conflict counters (_ld/_st) report conflicts directly.
• Spheron: confirming this needs real hardware counter access, which Spheron's bare-metal instances expose; see our FlashAttention-4 Blackwell kernel guide for the same requirement.

We've covered what happens when a warp's threads disagree about which branch to take in our warp divergence guide; a bank conflict is the memory-side version of the same lesson, a pattern that looks fine in source and costs you serialized hardware cycles anyway.

What a Shared Memory Bank Conflict Actually Is

A shared memory bank conflict happens when two or more threads in the same warp issue addresses, in the same instruction, that fall into the same memory bank but point at different locations. Shared memory can service only one address per bank per clock cycle, so instead of completing the warp's request in one pass, the hardware splits it into as many separate, sequential requests as there are distinct addresses competing for that bank. A load that should take one cycle's worth of bank bandwidth ends up taking several, and the warp just waits.

A few properties worth holding onto before going further:

• Shared memory is split into 32 banks, one per thread in a warp.
• Each bank delivers 4 bytes (one 32-bit word) of bandwidth per cycle.
• Conflicts are evaluated per warp, per instruction, not across the whole kernel's lifetime.
• Multiple threads reading the same address is a broadcast, not a conflict.
• The worst case, 32 threads hitting one bank at 32 different addresses, runs up to 32x slower than a conflict-free access.

Shared Memory's 32 Banks and the 4-Byte Mapping

Per NVIDIA's own guidance, "successive 32-bit words are assigned to successive banks," and on compute capability 5.x and later, each of the 32 banks has 32 bits of bandwidth per clock cycle. That means the address-to-bank mapping is just the word index modulo 32: word 0 sits in bank 0, word 1 in bank 1, word 32 wraps back to bank 0, word 33 to bank 1, and so on. This is the entire mechanism. There's no cache, no associativity, no eviction policy to reason about, only a modulo-32 lookup, which is exactly why the pattern that triggers a conflict is so easy to write without noticing: any access with a stride that's a multiple of 32 words lands every thread on the same bank.

This 32-bank SRAM is the same tier covered at a higher level in our GPU memory hierarchy guide, which looks at shared memory's capacity and bandwidth relative to registers and HBM. This post goes one level further down, into how that SRAM is actually organized internally and what breaks when you cross a bank boundary wrong.

What an N-Way Bank Conflict Costs: Serialization, Not Just Slowness

The cost model here is exact, not a vague "it gets slower." As NVIDIA's CUDA C++ Best Practices Guide puts it: "If multiple addresses of a memory request map to the same memory bank, the accesses are serialized. The hardware splits a memory request that has bank conflicts into as many separate conflict-free requests as necessary." An n-way conflict, meaning n distinct addresses from the warp landing on one bank, decreases effective bandwidth for that instruction by a factor of n, because the hardware now needs n passes instead of one to retire the same amount of work.

The Broadcast Exception Everyone Confuses With a Conflict

Here's the detail that trips people up when they first learn the rule: if every thread in a warp reads the exact same address, that's not a conflict at all. The hardware detects that the request is identical for every thread and services it with a single read, broadcasting the result to every thread that asked, rather than treating it as 32 competing requests for one bank. This is why a kernel that loads a shared scalar, a bias term, a scale factor, a reduction result, into every thread in a warp pays nothing extra for doing so, even though on paper 32 threads are "hitting the same bank." The rule that matters is whether the addresses are identical, not whether they share a bank; identical addresses broadcast for free, differing addresses on the same bank serialize.

Where Shared Memory Bank Conflicts Come From in Ordinary Kernels

You don't need quantization to hit this. The textbook trigger is a 2D shared memory tile accessed two different ways: row-wise in one pass, column-wise in the next. A tile-based matrix transpose is the canonical example:

__shared__ float tile[32][32];

// Load: contiguous threads write contiguous addresses. No conflict.
tile[threadIdx.y][threadIdx.x] = input[row * width + col];
__syncthreads();

// Store: contiguous threads now read down a COLUMN of the tile.
output[col * width + row] = tile[threadIdx.x][threadIdx.y];

The load is fine: for a fixed threadIdx.y, the 32 threads in a warp write 32 contiguous addresses, one per bank. The store is where it breaks. For a fixed threadIdx.y, reading tile[threadIdx.x][threadIdx.y] across a warp means threadIdx.x varies while the row stride is 32 floats, so every thread's address differs from the next by exactly 32 words, a full lap around the bank count. Every thread in the warp lands on the same bank, at a different offset, which is the 32-way worst case described above.

Vectorized shared memory accesses can produce the same outcome through a less obvious door. Loading or storing 128-bit chunks (four packed floats, or int4, at a time) is normally a win because it cuts the number of memory transactions per warp. But a 128-bit access from one thread spans four consecutive banks, not one, so if the per-thread stride between those four-bank spans lines up wrong, threads can still collide on the same subset of banks even though each individual thread's own access is internally conflict-free. Widening the access reduces instruction count; it doesn't automatically fix the underlying stride-to-bank arithmetic, and in a kernel that was conflict-free at 32-bit granularity, switching to a wider vectorized store is a real way to introduce a conflict that wasn't there before.

The general rule behind both of these: any shared memory access pattern where the stride between consecutive threads' addresses is a multiple of the bank count (32 for 4-byte elements, fewer effective threads per lap for wider ones) will put multiple threads on the same bank. Row-major 2D tiles accessed column-wise are just the most common way a kernel ends up with that stride by accident.

The Quantization Connection: How Packing Low-Bit Weights Triggers Conflicts

This is where the mechanism above stops being a CUDA-101 exercise and starts costing real inference throughput. Quantizing a model's weights to INT4 is the whole reason you'd bother, and our AWQ quantization guide covers why teams accept the accuracy tradeoff: it's roughly a 50% VRAM reduction, often the difference between needing one GPU and needing two. But that same packing, cramming more values into fewer bytes, is also the single easiest way to reconstruct the exact stride-32 pattern above inside the kernel that has to unpack those weights before using them.

Why INT4/INT2 Packing Changes Your Effective Stride

An unquantized FP16 weight tensor has a simple, well-behaved stride: one element is two bytes, consecutive elements are two bytes apart, and a transpose or tiling kernel written against that layout rarely surprises you. Pack eight INT4 values into a single 32-bit word, or sixteen INT2 values into the same word, and that relationship breaks. The element you actually want for thread i is now a sub-word slice (a 4-bit or 2-bit nibble) at a bit offset determined by i, not a standalone addressable unit, and the word that holds it is shared across several logical weight indices.

A naive dequant-and-rearrange kernel, the kind you'd write before thinking hard about it, unpacks that word into a shared memory scratch tile using the original weight indexing to decide where in the tile each unpacked value lands. If that scratch tile is then read back in a transposed or strided order for the matmul that follows, as it typically is, the effective stride between threads is no longer the simple one-FP16-element-at-a-time pattern a non-quantized kernel would have had. It's a function of the packing width and the tile's layout, and it's easy for that function to land on exactly the multiple-of-32 stride that produces the worst-case conflict, especially once you add a group-wise scale factor lookup on top, which introduces its own stride into the same loop. The packing itself didn't cause the conflict directly; it changed the addressing math enough that a kernel which would have been conflict-free at full precision isn't anymore.

How Marlin and Other Production INT4 Kernels Design Around It (XOR Swizzle)

Production INT4 GEMM kernels design around this instead of hoping it doesn't happen. MARLIN, a kernel built for mixed-precision auto-regressive inference on large language models, documents this directly in its own implementation notes: the kernel "performs several layout transformations to guarantee that all shared memory reads and writes are conflict-free, in particular for matrix loading instructions." Rather than relying on the natural row-major address-to-bank mapping and hoping the stride works out, it deliberately reorders where packed weight and activation tiles physically live in shared memory, so the loads feeding the tensor cores stay conflict-free regardless of which row or column a warp is reading.

The general technique costs no extra shared memory, unlike padding, at the price of a layout that's harder to reason about by eye. If you're packing low-bit weights or KV cache entries into dense storage for inference, as covered in a different context in our Google TurboQuant guide, this is the pattern to reach for once padding's wasted capacity actually matters.

Reproducing It: A Triton Kernel, Before and After Padding

Here's the same mechanism in a form you can write and profile yourself, using the pattern a quantized-weight unpack kernel actually has: load a tile of packed INT4 weights, unpack it, and transpose it into the layout the following matmul expects. Our Triton kernel development guide covers everything else Triton takes off a kernel author's plate; the shared memory layout decision below is the one piece it still leaves to you.

The Naive Kernel and Its ncu Trace

import triton
import triton.language as tl

@triton.jit
def dequant_transpose_naive(
    w_packed_ptr, scale_ptr, out_ptr,
    N: tl.constexpr, BLOCK: tl.constexpr,
):
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    rows = pid_m * BLOCK + tl.arange(0, BLOCK)
    cols = pid_n * BLOCK + tl.arange(0, BLOCK)

    # unpack 8x INT4 values per 32-bit word
    word = tl.load(w_packed_ptr + rows[:, None] * (N // 8) + cols[None, :] // 8)
    shift = (cols[None, :] % 8) * 4
    nibble = (word >> shift) & 0xF
    scale = tl.load(scale_ptr + rows[:, None])
    w = (nibble.to(tl.float16) - 8.0) * scale

    # transpose the unpacked BLOCK x BLOCK tile for the matmul layout
    tile = tl.trans(w)
    out_rows = pid_n * BLOCK + tl.arange(0, BLOCK)
    out_cols = pid_m * BLOCK + tl.arange(0, BLOCK)
    tl.store(out_ptr + out_rows[:, None] * N + out_cols[None, :], tile)

With BLOCK set to 32, tl.trans has to route a 32x32 tile through a shared memory round trip to flip rows and columns, which is exactly the matrix-transpose pattern from the CUDA example earlier, now sitting downstream of an INT4 unpack instead of a plain FP32 load. Profile this with ncu and the bank-conflict counters described below should show the same magnitude of regression the published stride-32 microbenchmark measured: efficiency collapsing from the mid-90s percent down into the low single digits.

The One-Line Padding Fix and the Trace After

The fix mirrors the CUDA fix exactly: give the intermediate tile one extra column so its row stride stops being an exact multiple of 32 words. In practice that means routing the transpose through an explicitly padded scratch buffer instead of a bare BLOCK x BLOCK tile:

# Allocate the scratch tensor from the host with one padding column:
# scratch = torch.empty((BLOCK, BLOCK + 1), device="cuda", dtype=torch.float16)

tl.store(scratch_ptr + rows_local[:, None] * (BLOCK + 1) + cols_local[None, :], w)
tl.debug_barrier()
tile = tl.load(scratch_ptr + cols_local[:, None] * (BLOCK + 1) + rows_local[None, :])

That BLOCK + 1 stride is the entire fix. It shifts every row of the scratch tile over by one extra element relative to the row before it, so a column-wise read across a warp no longer lands every thread on the same bank. The cost is one wasted column of shared memory per tile, typically a rounding error against the conflict it removes, which is exactly the tradeoff the padding fix for 2D tiles describes in the general CUDA case. A matrix-transpose benchmark measured on an RTX 3090 found both the padded version and an XOR-swizzled version landing at the same 0.92ms, down from 1.10ms for the conflicted version, about a 20% improvement from either fix.

Fixing It: Padding vs. Swizzling, and When Each Applies

Padding is the right default when shared memory capacity isn't the binding constraint; it's a one-line change and it's easy to verify. It stops being free once your tile size is already pushing the occupancy ceiling, since shared memory per block is one of the three resources that cap how many warps an SM can run concurrently, alongside registers and the thread-slot limit. Adding a padding column to a tile that's already near the shared-memory budget can drop occupancy enough to offset the gain from removing the conflict, which is exactly the scenario XOR swizzling exists for: the same conflict-free guarantee, with zero extra bytes, at the cost of a layout a reader can no longer eyeball from the index math alone.

Measuring Shared Memory Bank Conflicts With Nsight Compute

Nsight Compute exposes the mechanism directly rather than making you infer it from wall-clock time. Two metrics count bank conflicts on loads and stores separately: l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld for loads and the corresponding _st metric for stores. Run them against a kernel with:

ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st \
  python dequant_bench.py

A healthy, conflict-free kernel reports these counters at or near zero. A kernel with the stride-32 pattern from the transpose example will show a nonzero count scaling with the number of warps that hit the conflicted path.

Before spending an afternoon chasing a fix, confirm the conflict is actually the bottleneck. For the broader ncu workflow this fits into, replay modes, roofline charts, remote capture without a display, our Nsight Compute and PyTorch Profiler guide covers the full setup these two metrics slot into.

What This Costs You in Rented GPU-Hours When It Goes Unnoticed

A dequant kernel with an unnoticed 32-way bank conflict on its hottest loop isn't failing. It's running, producing correct output, and quietly burning multiples of the GPU-hours it should need, which on rented hardware is a direct, recurring line item rather than a one-time inefficiency. A kernel stuck at a fraction of its achievable throughput because of a conflicted shared memory access pattern means more wall-clock time per batch, which means more billed hours for the same workload, indefinitely, until someone profiles it.

Checking for this requires ncu's hardware performance counters, and that's a structural requirement most serverless and shared GPU platforms don't meet: tenant isolation on those platforms typically blocks the counter access ncu needs, independent of which provider is involved. Spheron rents GPUs as either a fully provisioned VM or a bare-metal instance billed per minute with no monthly minimum, and bare-metal allocation is what passes NVIDIA driver capabilities through to the container so ncu can read the l1tex counters directly instead of failing on a permissions error. A short profiling session, compile the kernel, run ncu, read the counters, costs a fraction of an hour at the $2.98/hr on-demand rate for an H100. It's worth being clear about what that buys you and what it doesn't: renting bare-metal access makes the conflict visible, it doesn't fix it. The fix is still a code change, padding or swizzling the shared memory layout, and no amount of better hardware underneath an unfixed kernel changes that.

Pricing fluctuates based on GPU availability. The Spheron rate above is live as of 02 Oct 2026. Check current GPU pricing → for live rates.

A shared memory bank conflict is one of the few GPU performance problems with a clean, mechanical cost model: count the distinct addresses landing on one bank, that's your serialization factor. Quantized kernels run into it more than full-precision ones because packing weights tighter changes the stride math a kernel author has to get right, not because low-bit formats are inherently conflict-prone. Pad the tile, swizzle the layout, or let Triton and CUTLASS make that decision for you, and the conflict disappears as completely as it arrived.

Profiling a dequant kernel for bank conflicts needs real ncu counter access, not a monthly commitment.

Check H100 availability →