Amplifying Latency and Cost of LLM Inference via Rowhammer Attack

Mrityunjay Shukla; Shubham Roy; Biswabandan Panda

cs.CROct 22, 2026

Amplifying Latency and Cost of LLM Inference via Rowhammer Attack

Authors: Mrityunjay Shukla, Shubham Roy, Biswabandan Panda

Organizations: Indian Institute of Technology Bombay, Mumbai, India

Abstract

Large language models (LLMs) achieve high inference efficiency using optimized GPU kernels generated by just-in-time (JIT) compilation frameworks such as Triton. Triton compiles GPU kernels through an MLIR/LLVM-based pipeline. We investigate whether a Rowhammer-induced bit-flip in the CPU's DRAM can degrade GPU-based LLM inference performance while preserving output correctness. We present LLVMhammer, the first Rowhammer attack that causes LLM inference slowdown by inducing a bit-flip in Triton's compilation pipeline. A bit-flip in a performance-critical optimization pass causes Triton to generate a functionally identical but suboptimal GPU kernel. Since Triton caches the generated kernel on disk, a single successful bit-flip causes persistent slowdown across subsequent inference executions. Triton's target-independent stages are shared across GPU vendors, which makes our attack GPU vendor-agnostic. To identify exploitable bit-flip locations without manual inspection, we build Karnage, a tool that isolates useful bit-flip locations in Triton's target-independent compilation pipeline. We evaluate LLVMHammer across a diverse set of LLMs, including Qwen2.5-7B-Instruct, Llama-3.1-8B-Instruct, Mistral-7B-Instruct-v0.3, and Deepseek-llm-7B-chat, across NVIDIA RTX 4090 and AMD RX 7900 XTX GPUs. LLVMhammer increases inference latency of an LLM inference server running on a single GPU by 4-16%, which, in turn, can translate into additional annual operating costs at the scale of multi-GPU deployments. We further show that LLVMhammer can also corrupt LLM outputs across different GPU vendors, reducing BERTScore to as low as 25%, illustrating the broader security implications of bit-flips in Triton's compilation pipeline.

Figures & tables

Fig. 5: LLVMhammer on the target-independent stages of the Triton compilation pipeline.

Fig. 6: Overview of Karnage’s four-stage pruning of useful bit-flip locations.

Function count	NVIDIA	AMD
Total functions in libtriton.so	121,735	122,148
Target-independent functions (namespace filter)	7,963	8,219
Target-independent functions (Intel PT filter)	254	254

TABLE I: Number of functions in libtriton.so tailored for NVIDIA and AMD GPUs

Instruction. A	Opcode A	Instruction. B	Opcode B
je	0x74	jne	0x75
jl	0x7C	jge	0x7D
jle	0x7E	jg	0x7F
jb	0x72	jae	0x73
jbe	0x76	ja	0x77
js	0x78	jns	0x79

TABLE II: Control-flow altering instruction pairs at 1-bit Hamming distance in the x86-64-based CPU

	Potentially useful bit-flip locations		Time per trial (s)		Total time (hrs)
Stage	NVIDIA	AMD	NVIDIA	AMD	NVIDIA	AMD
Stage 1: Namespace filter	1,183,008	1,326,122	30	30	9,858	11,051
Stage 1: Intel PT filter	38,769	42,996	30	30	323	358
Stage 4: GPU profiling	1,176	810	30	30	10	7

TABLE III: Search space reduction and estimated profiling time with subsequent Karnage stages

IR level	Baseline (unmodified)	Flipped
Vectorized vs. Scalar inline-asm emission
LLVM IR	⬇ ; single vectorized load (4x32-bit) % val = tail call { i32 , i32 , i32 , i32 } asm sideeffect "ld.global.v4.b32 { $0,$ 1, $2,$ 3 }, [ $4 ];" , "=r,=r,=r,=r,l" ( ptr addrspace ( 1 ) % ptr ) # 2 One vectorized inline-asm call loads 4$ \times$ 32-bit values.	⬇ ; four separate scalar loads % v0 = tail call i32 asm sideeffect "ld.global.b32 { $0 }, [$ 1 ];" , "=r,l" ( ptr addrspace ( 1 ) % ptr0 ) #6 % v1 = tail call i32 asm sideeffect "ld.global.b32 { $0 }, [$ 1 ];" , "=r,l" ( ptr addrspace ( 1 ) % ptr1 ) #6 ; ... two more scalar calls A single four-element vector load is scalarized into four scalar loads.
128-bit coalesced load vs. scalar 32-bit loads
PTX	⬇ ; 128-bit vectorized load @ % p0 ld . global . v4 . b32 { % r0 , % r1 , % r2 , % r3 }, [ % rd0 + 0 ] ; ; 128-bit vectorized store @ % p1 st . global . v4 . b32 [ % rd1 + 0 ], { % r4 , % r5 , % r6 , % r7 } ; 128-bit coalesced memory transactions; four values per instruction.	⬇ ; four scalar 32-bit loads @ % p0 ld . global . b32 { % r0 }, [ % rd0 + 0 ] ; @ % p0 ld . global . b32 { % r1 }, [ % rd0 + 4 ] ; @ % p0 ld . global . b32 { % r2 }, [ % rd0 + 8 ] ; @ % p0 ld . global . b32 { % r3 }, [ % rd0 + 12 ] ; ; four scalar 32-bit stores @ % p1 st . global . b32 [ % rd1 + 0 ], % r4 ; @ % p1 st . global . b32 [ % rd1 + 4 ], % r5 ; One 32-bit scalar load per value, resulting in 4 $\times$ more memory transactions.
128-bit LDG/STG vs. 16-bit scalar LDG
SASS	⬇ ; 128-bit vectorized global loads @P 0 LDG . E .128 R4 , [ R2 .64 ] ; @P 0 LDG . E .128 R8 , [ R6 .64 ] ; ; 128-bit vectorized global stores @P 1 STG . E .128 [ R1 0.64 ] , R4 ; @P 1 STG . E .128 [ R1 2.64 ] , R8 ; Each LDG/STG.E.128 loads/stores 128 bits in one transaction.	⬇ ; 16-bit scalar loads replace 128-bit loads @P 0 LDG . E . U16 R4 , [ R2 .64 ] ; @P 0 LDG . E . U16 R5 , [ R2 .64+ 0x 10 ] ; @P 0 LDG . E . U16 R6 , [ R2 .64+ 0x 20 ] ; @P 0 LDG . E . U16 R7 , [ R2 .64+ 0x 30 ] ; ; ... 48 x LDG.E.U16 total ; ... 0 x LDG/STG.E.128 remain Vectorized loads are replaced with multiple scalar loads.

TABLE IV: Effect of a bit-flip on the generated LLVM IR, PTX, and SASS

Component	Platform A	Platform B
CPU	Intel i7-4790 (Haswell)	Intel i7-8700 (Coffee Lake)
DRAM	8 GB Kingston DDR3 @ 1600	8 GB Corsair DDR4 @ 2400
OS	Ubuntu 20.04.6 LTS
Kernel	5.15.0-131	6.2.0-060200
GPU	NVIDIA RTX 4090 (24 GB GDDR6X) AMD RX 7900 XTX (24 GB GDDR6)
LLM Models	Qwen2.5-7B-Instruct [ 28 ] , Llama-3.1-8B-Instruct [ 29 ] Mistral-7B-Instruct-v0.3 [ 30 ] , Deepseek-llm-7B-chat [ 31 ]

TABLE V: Platforms for evaluation

Vendor	Model	$n=10$	$n=100$	$n=1000$
NVIDIA	DeepSeek	2/5	3/5	✓
	Llama	2/5	✓	✓
	Mistral	0/5	3/5	✓
	Qwen	1/5	1/5	✓
AMD	DeepSeek	✓	✓	✓
	Llama	1/5	✓	✓

TABLE VI: Number of metrics (out of five: TTFT, TPOT, ITL, throughput req , throughput tok ) with

p<0.05

under Welch’s t-test (baseline vs. flipped, 30 runs each). ✓ denotes all five metrics significant

Fig. 7: Batched inference slowdown in terms of throughput(requests/second), throughput(tokens/second), TTFT, TPOT, and ITL with varying number of requests for (a) NVIDIA and (b) AMD GPUs.

Fig. 8: Distribution of total inference time for n=1000 number of requests across 30 runs for (a) NVIDIA and (b) AMD GPUs. Although an increase of 4-16% does seems modest, in production inference servers which serve millions of requests [ 33 , 34 ] slowdown becomes quite apparent.

Platform	Baseline output	Flipped output
Prompt: A hospitalized 70-year-old woman develops thrombocytopenia on day 7 post orthopedic surgery $\ldots$ What is the most likely cause?
AMD	⬇ A : Immune - mediated thrombocytopenia B : Thrombotic thrombocytopenic purpura C : Drug - induced thrombocytopenia ...	⬇ ,000000000000000000000000000000000 0000000000000000000000000000000000 n000000000000000000000000000000000
NVIDIA	⬇ A ) Heparin - induced thrombocytopenia B ) Immune thrombocytopenia purpura C ) Thrombotic thrombocytopenic purpura ...	⬇ from theise from the , theise , theise from the , the from the , the from the , the from the , the from the from theise

TABLE VII: LLM output corruption after a bit-flip on AMD and NVIDIA GPUs. The same prompt produces coherent output under the baseline kernel and incoherent output under the suboptimal kernel on both platforms, demonstrating cross-architecture output corruption from a single target-independent bit-flip location.

LLM Security

Large language model (LLM) security studies attacks, privacy risks, and defenses across model training, inference, and deployment.

See all 318 papers

Apr 19, 2026

Bit-Flip Vulnerability of Shared KV-Cache Blocks in LLM Serving Systems

Rowhammer on GPU DRAM has enabled adversarial bit flips in model weights; shared KV-cache blocks in LLM serving systems present an analogous but previously unexamined target. In vLLM's Prefix Caching, these blocks exist as a single physical copy without integrity protection. Using software fault injection under ideal bit targeting, we characterize worst-case severity and identify three properties: (1) Silent divergence - 13 of 16 BF16 bit positions produce coherent but altered outputs, indistinguishable from legitimate responses without a clean baseline. (2) Selective propagation - only requests sharing the targeted prefix are affected. (3) Persistent accumulation - no temporal decay occurs, so cumulative damage grows linearly with subsequent requests. Together, these constitute a threat profile distinct from weight corruption: silent divergence and selective propagation enable detection evasion; persistent accumulation then proceeds unchecked, yielding damage amplification bounded only by how long the block remains cached. A checksum-based countermeasure detects any single-bit corruption at scheduling time, bounding cumulative damage to one batch independent of the block's cache lifetime, with negligible overhead. These results argue for integrity protection of prefix blocks before end-to-end exploitation is demonstrated.

Yuji Yamamoto, Satoshi Matsuura

Institute of Science Tokyo, Tokyo, Japan

LLM Serving Adversarial Attacks on LLMs LLM Security

Ask PDF

May 20, 2026

Trusted Weights, Treacherous Optimizations? Optimization-Triggered Backdoor Attacks on LLMs

Inference optimization aims to minimize the latency and resource consumption of LLM inference while preserving output quality, making large-scale deployment practical and cost-effective. However, optimized execution can introduce small numerical inconsistencies from the original model. We reveal that this inconsistency not only causes the model's outputs to diverge, but more critically can introduce hidden backdoors. The backdoor remains dormant under standard unoptimized execution and is activated only when inference optimization is enabled, allowing it to evade existing backdoor detection pipelines. We first introduce the Input-Specific Optimization Backdoor (ISOB) to demonstrate that optimization-induced differences can cause wrong predictions. However, ISOB remains input-specific and cannot establish a universal optimization-triggered backdoor. To overcome this limitation, we design the Universal Optimization Backdoor (UOB). The backdoored model stays benign under unoptimized execution but activates when inference optimization is enabled. We conduct extensive experiments across seven mainstream open-source LLMs, four tasks, and three optimization backends. UOB reaches up to 100% attack success while largely preserving clean accuracy. To mitigate this vulnerability, we design three defense methods that reduce the backdoor ASR to at most 0.02 while preserving clean accuracy. These results reveal inference optimization as a new LLM security attack surface and motivate defenses against test-deployment disagreement.

Yifei Wang, Yida Yang, Tianlin Li +3

Shanghai Jiao Tong University · Tongji University · Beihang University +1

LLM Backdoor Attacks Backdoor Attacks LLM Security Backdoor Defense

Ask PDF

Feb 13, 2026

TensorCommitments: A Lightweight Verifiable Inference for Large Language Models

Most large language models (LLMs) run on external clouds: users send a prompt, pay for inference, and must trust that the remote GPU executes the LLM without any adversarial tampering. We critically ask how to achieve verifiable LLM inference, where a prover (the service) must convince a verifier (the client) that an inference was run correctly without rerunning the LLM. Existing cryptographic works are too slow at the LLM scale, while non-cryptographic ones require a strong verifier GPU. We propose TensorCommitments (TCs), a tensor-native proof-of-inference scheme. TC binds the LLM inference to a commitment, an irreversible tag that breaks under tampering, organized in our multivariate Terkle Trees. For LLaMA2, TC adds only 0.97% prover and 0.12% verifier time over inference while improving robustness to tailored attacks by up to 48% over the best prior work requiring a verifier GPU.

Oguzhan Baser, Elahe Sadeghi, Eric Wang +5

Electrical and Computer Engineering, The University of Texas at Austin · Theseus AI Labs · Electrical and Computer Engineering, McGill University +1

LLM Inference LLM Security Verifiable Outsourced Computation

Ask PDF

Abstract

LLM Security

Explore similar work

Bit-Flip Vulnerability of Shared KV-Cache Blocks in LLM Serving Systems

Trusted Weights, Treacherous Optimizations? Optimization-Triggered Backdoor Attacks on LLMs

TensorCommitments: A Lightweight Verifiable Inference for Large Language Models