{"id":"ef8ba128-2c14-4b0c-9cbd-38294c555548","arxiv_id":"2507.08166","paper_version":1,"verdict":"CONDITIONAL","confidence":"MODERATE","novelty_score":8.0,"correctness_risk":"medium","formal_verification":"none","parameter_count":5,"one_line_summary":"First successful Rowhammer attack on discrete GPU GDDR6 memory (NVIDIA A6000), with 8 bit-flips and a single-bit-flip ML accuracy degradation from 80% to below 1%.","lead":"GPUHammer demonstrates the first Rowhammer attack on a discrete NVIDIA GPU, flipping bits in GDDR6 memory on an A6000. The paper then shows that a single such bit-flip can reduce an image classifier's accuracy from about 80% to below 1%.","discovery_kind":"new_application","skeptic_critique":{"model":"deepseek-v4-flash","headline":"Memory massaging in §8.2 relies on a cross-process handoff that RMM's per-process pool cannot provide; the ML tampering exploit lacks a demonstrated delivery mechanism.","rationale":"The paper's strongest claim has two parts: (1) first successful Rowhammer bit-flips in GDDR6 on a discrete GPU, and (2) an end-to-end ML accuracy degradation exploit using those bit-flips. I find part (1) largely supported by the reported 8 reproducible bit-flips, the TRH and TRR sampler characterizations, and the careful cache/discard and latency methodology. Part (2), however, depends on the memory massaging handoff in §8.2, and that step appears to fail under the stated multi-tenant, separate-process threat model. RMM is a per-process pool allocator; when the attacker frees a chunk, it stays in the attacker's pool and is not available to a different process. The paper's own comparison with cudaMalloc makes this worse, because cudaMalloc—the mechanism that could return memory to the driver for another process to obtain—is explicitly deferred. This is not a disagreement with community consensus; it is an internal inconsistency between the threat model and the described allocation mechanism. The reader's weakest assumption about virtual-to-physical mapping stability is also plausible and related, but the RMM issue is more immediately fatal to the ML exploit claim: even if mapping is perfectly stable, a separate victim process will not receive attacker-freed pool memory. A focused two-process experiment can settle whether the handoff occurs; if it does not, the ML tampering contribution should be reported as simulated or conjectural rather than demonstrated. The bit-flip contribution and the mapping/refresh reverse-engineering remain valuable and would justify a conditional accept pending artifact verification and this massaging check.","tokens_in":27465,"tokens_out":14169,"duration_ms":165656,"concrete_test":"Implement the §8.2 massaging procedure as two separate CUDA processes: process A (attacker) allocates a large region via RMM, frees the chunk containing a characterized bit-flip (e.g., D1), and keeps its RMM pool alive; process B (victim) then requests memory via RMM/PyTorch. Use the §4 row-conflict timing method to check whether B's allocation maps to the freed physical chunk and to attacker-controlled adjacent aggressor rows. Repeat using the released artifact. If B does not land on the freed chunk, the memory massaging step is not achievable under the stated multi-tenant threat model, and Table 4 cannot be attributed to the described attack. A complementary check is to trace RMM pool behavior in the artifact to confirm whether freed blocks are returned to the CUDA driver.","verdict_should_be":"CONDITIONAL","load_bearing_attack":"The end-to-end exploit's memory massaging mechanism is internally inconsistent with RMM's design. Section 8.2 says the attacker 'frees the chunk that contains a known flippy bit' and that the victim process (e.g., PyTorch) later receives that chunk, with RAPIDS Memory Manager (RMM) 'immediately reus[ing] any freed memory.' But RMM's pool is per-process/per-context: memory freed into the attacker's RMM pool is retained in that pool and is not returned to the CUDA driver, so a separate victim process cannot be allocated that chunk. The paper explicitly contrasts RMM with cudaMalloc and leaves cudaMalloc-based attacks for future work; however, cudaMalloc is the path that actually returns memory to the driver and could enable cross-process handoff. Without a shared cross-process allocator, the attacker cannot place victim weights onto a characterized flippy bit while retaining access to the aggressor rows needed to hammer it. Consequently, Table 4's accuracy degradation results—central to the abstract's claim that the attacker can 'tamper with ML models'—have no demonstrated Rowhammer delivery path. The bit-flip finding itself is not invalidated, but the ML exploit contribution rests on an unsupported allocation handoff.","agreement_with_reader":"partial"},"referee_report":{"model":"deepseek-v4-flash","summary":"GPUHammer reports the first Rowhammer bit-flips in discrete GPU GDDR6 DRAM. On an NVIDIA RTX A6000, the authors reverse-engineer the virtual-address to DRAM bank/row mapping using latency-based conflict testing, develop multi-warp hammering that approaches ~620K activations per refresh window, synchronize many-sided hammering patterns to tREFI, and observe 8 single-bit flips across 4 DRAM banks. They additionally characterize Rowhammer threshold, TRR sampler size, bit-flip direction, and data-pattern dependence, and claim an end-to-end ML tampering exploit that degrades ImageNet model accuracy via a single bit-flip in the MSB of an FP16 weight exponent.","tokens_in":27692,"tokens_out":6970,"duration_ms":75647,"significance":"If confirmed, the bit-flip result is a significant first: it demonstrates read-disturbance faults in GDDR6 on a discrete NVIDIA GPU, extending Rowhammer beyond CPU DDR/LPDDR memories, and it contributes a feasible methodology for reverse-engineering proprietary GPU DRAM mappings via unprivileged CUDA kernels. The paper also ships an artifact and follows responsible disclosure. However, the evidence base is thin: 8 bit-flips on a single A6000 unit, with key characterizations based on best-of-50 results and no error bars. More importantly, the claimed ML-tampering exploit lacks a working memory-massaging delivery mechanism, as the RMM pool is per-process. The core bit-flip finding is plausible and publishable, but the exploit contribution and several hardware-property conclusions need substantial repair.","major_comments":[{"comment":"The memory-massaging step described in Section 8.2 is not supported by RMM's design. RMM's pooling memory resource is per-process: memory freed into the attacker's pool is retained in that pool and is not returned to the CUDA driver, so a separate PyTorch process cannot subsequently receive that chunk. The statement that 'this chunk is allocated to the victim process (e.g., PyTorch)' therefore lacks a mechanism. The experiments in Section 8.3 need either a demonstrated cross-process handoff (e.g., via cudaMalloc-based reuse or CUDA IPC) or must be re-scoped to a same-process or simulated setting. As written, Table 4 and Takeaway 5 overstate the practical ML-tampering exploit, which is central to the abstract's claim that an attacker can 'tamper with ML models'.","section":"Section 8.2"},{"comment":"The conclusion that the TRR-like sampler in A6000 GDDR6 tracks at most 16 rows per bank rests on the fraction of hammers triggering bit-flips from 50 attempts per n-sided pattern on a single victim bit (A1), with no confidence intervals or repetitions on another bank or another chip. Given that only 8 bit-flips were observed in total, the abrupt threshold at n=17 should be reported with error bars and repeated on at least one more bank and ideally another unit before being stated as a hardware property.","section":"Section 7.4, Figure 12"},{"comment":"The attack evaluation is based on a single A6000 unit and only 8 bit-flips (one per row), with no run-to-run variation, no multi-chip results, and no error bars on bit-flip counts, TRH, or activation rates. Since Rowhammer susceptibility is known to vary with DRAM process and temperature, the title/abstract claim that GPU Rowhammer attacks 'are practical' is broader than the evidence supports. Please either test additional units and report distributions, or scope the claims to the specific tested sample.","section":"Section 7.1, Table 1"},{"comment":"The attack's repeatability depends on the empirical observation that, for large allocations, the mapping of offsets within the array to DRAM rows remains the same across system/program restarts. This behavior is not guaranteed by any hardware or software contract and was observed only on one system. Driver updates, memory pressure, fragmentation, or different GPU instances could invalidate precomputed row-sets. The paper acknowledges this as a possible mitigation but does not evaluate robustness; this dependency should be explicitly stated as a load-bearing limitation, and ideally tested under at least one perturbation (e.g., driver version or allocation size).","section":"Section 4.3, Observation 4 and Section 10"}],"minor_comments":[{"comment":"The delay-generation loop 'for (size_t i = delay; i--;) dummy_sum += dummy;' accumulates into a local variable that is never used, which is dead code under standard compiler optimization and may be eliminated, undermining the intended per-warp synchronization delay. Use a volatile accumulator or a side-effecting operation to preserve the delay.","section":"Appendix A, Listing 4"},{"comment":"The reported RAD values are the highest across 50 attempts; please also report the median and distribution (e.g., min/max or percentiles) so the reader can assess typical rather than best-case attack impact.","section":"Section 8.3, Table 4"},{"comment":"The introduction states a refresh period of 22 ms (or 32 ms), while Observation 5 concludes 23 ms from tREFI = 1407 ns and 16K tREFIs. Please reconcile these numbers and cite the measurement details consistently.","section":"Section 6, Observation 5"},{"comment":"The claim that 8-warp configurations achieve 'strong synchronization' is based on a single visual flat-line at ~1407 ns; adding per-configuration variance or repeated measurements would strengthen the synchronisation analysis.","section":"Figure 10"}],"recommendation":"major_revision","confidential_remarks":"The core bit-flip demonstration is likely publishable and is the paper's main contribution, but the ML exploit currently contains a correctness error: the RMM-based cross-process handoff cannot work because RMM pools are per-process. This is not a presentation issue but a load-bearing flaw in the exploit story. I would encourage the authors to either provide a real cross-process delivery mechanism or explicitly reframe the ML results as simulations of bit-flip effects rather than demonstrated attacks."},"author_rebuttal":null,"desk_editor":{"model":"deepseek-v4-flash","letter":"Read GPUHammer. The core finding—8 bit-flips on a single A6000 GDDR6, first on discrete GPU—is real and worthwhile. The reverse-engineering of virtual-to-DRAM row mapping without physical addresses is clever, and the multi-warp + REF-synchronized hammering is a genuine engineering step past CPU-based attacks. If it holds up on more hardware, this opens a new attack surface for cloud ML.\n\nThe soft spot is the ML exploit. The stress-test concern is right: Section 8.2 describes freeing a chunk into RMM and having a victim process allocate it, but RMM's pool is per-process, so that memory never returns to the driver. The cross-process handoff doesn't work as written. That doesn't invalidate the bit-flip finding, but it does mean Table 4 and the 'up to 80% accuracy drop' claim currently have no demonstrated Rowhammer delivery path. The authors need to either fix the massaging (same-process attacker? cudaMalloc? some driver-level mechanism?) or present the ML damage as a simulated consequence of the bit-flips rather than an end-to-end exploit.\n\nAlso worth tightening: single GPU unit, no error bars, best-of-50 in the ML experiments, and some internal inconsistencies in activation rates (500K vs 620K vs 393K ACTs per refresh) and refresh period (22ms vs 23ms). The TRR sampler size of 16 is an inference, not a measurement, and the '8 bit-flips' is a small sample. None of these are fatal to the main claim, but they are exactly the kind of thing that needs a second unit and an artifact before publication.\n\nI'd send this to peer review, not desk reject. The bit-flip result is the kind of empirical finding the community needs; it's just that the paper is overclaiming the exploit. A serious referee can separate the two. Recommend major revision.","headline":"First real shot at Rowhammer on discrete GDDR6, but the ML exploit's delivery mechanism doesn't hold together.","tokens_in":28268,"tokens_out":3578,"would_cite":true,"duration_ms":36967,"reading_group":"yes","serious_thinker":"yes","would_accept_peer_review":true},"rs_alignment":null,"lean_confirmation":null,"pith_extraction":{"msc":[],"pacs":[],"model":"deepseek-v4-flash","headline":"This paper demonstrates the first Rowhammer attack on a discrete NVIDIA GPU, flipping bits in GDDR6 DRAM and using one flipped model-weight bit to collapse ML accuracy to near zero.","keywords":["Rowhammer","GPU security","GDDR6","DRAM bit-flips","NVIDIA GPU","deep learning model tampering","memory massaging","refresh synchronization"],"falsifier":"Recompute the row-sets on the same A6000 after a reboot, a driver update, or memory pressure, then repeat the hammering campaign without re-profiling; if the previously flippable addresses no longer flip, the attack's reusable row-set assumption fails. A second decisive test is to run the same campaign on a second A6000 unit and see whether any bit-flips occur at all.","tokens_in":27223,"feed_emoji":"⚡","tokens_out":7916,"duration_ms":76423,"temperature":0.7,"pith_summary":"GPUHammer claims that Rowhammer is practical on discrete NVIDIA GPUs, not just CPUs. It reverses the proprietary mapping from GPU virtual addresses to GDDR6 DRAM banks and rows, then hammers rows with a parallel CUDA kernel synchronized to refresh commands. On an NVIDIA A6000 with ECC off, this produced the first reported Rowhammer bit-flips in GDDR6 memory: up to 8 single-bit flips across 4 DRAM banks. The paper then shows that one such flip in the MSB of an FP16 weight's exponent can cut ImageNet top-1 accuracy from about 80% to below 0.5% across five common models. If correct, this makes multi-tenant GPU inference a realistic Rowhammer target.","feed_headline":"First Rowhammer bit-flips hit GDDR6 in a discrete GPU","feed_subtitle":"An unprivileged CUDA kernel flips bits across 4 DRAM banks and drops ImageNet accuracy from 80% to under 0.5%.","key_machinery":"Three mechanisms carry the attack. First, since GPU physical addresses are hidden, the authors learn a lookup table at 256-byte granularity from virtual offsets to DRAM banks and rows by timing row-buffer conflicts, filtering out non-uniform memory-access latency by comparing only addresses with similar single-access latency. Second, the PTX discard instruction plus ld.global.volatile forces every access to miss all caches, and assigning one aggressor row per warp lets multiple warps overlap their memory round trips, raising activations from about 90K to roughly 620K per refresh window, near the theoretical ceiling. Third, per-warp add-instruction delays create an aligned idle bubble at the memory controller so that hammering rounds synchronize with refresh commands; keeping at most 8 warps preserves that synchronization and lets 17-24 aggressor rows overflow the 16-entry TRR sampler.","core_discovery":"The paper's central claim is that an unprivileged CUDA kernel can induce read-disturbance bit-flips in GDDR6 DRAM on a discrete NVIDIA GPU, by recovering the virtual-address-to-DRAM-row mapping, maximizing activation rates through multi-warp hammering, and synchronizing n-sided aggressor patterns to the 1407ns refresh interval to defeat a TRR-like mitigation that tracks at most 16 rows per bank. On the A6000, this yields 8 bit-flips across all 4 hammered banks with a minimum Rowhammer threshold of 12.3K activations, and a single 0-to-1 flip at bit position 6 of a byte maps to the MSB of the FP16 exponent in the weight layout used by common ML runtimes, degrading ResNet50 accuracy from 80.26% to 0.08% (RAD of 0.99). The paper also reports no bit-flips on an A100 with HBM2e or an RTX 3080 with GDDR6, attributing the difference to chip variation, higher thresholds, or on-die ECC.","pith_inferences":["If the same discard/volatile primitives and 256-byte mapping granularity hold on other GDDR6 parts, the reverse-engineering and hammering methods likely transfer to those GPUs, even though bit-flips were only observed on the A6000 in this study.","The observed data-dependence of bit-flip frequency on victim and aggressor bytes suggests a possible Rowhammer-based side channel that could read GPU memory contents; the paper notes this connection but does not implement it.","On-die ECC in HBM2e, HBM3, and GDDR7 may hide single-bit Rowhammer errors, but ECC can mis-correct when multiple bits in one codeword flip, so future GPU Rowhammer may become harder rather than impossible.","Cloud GPU operators cannot conclude safety from the A100 or RTX 3080 results, since Rowhammer thresholds and TRR sampler sizes vary by chip, batch, and temperature; each GPU SKU would need its own empirical test."],"forward_implications":["Multi-tenant GPU platforms with GDDR6 and ECC disabled should treat co-located CUDA kernels as capable of corrupting other tenants' GPU memory.","FP16 ML inference is an actionable target: a single exponent-MSB flip can collapse top-1 accuracy from 80% to 0.02% or below, as shown on AlexNet, VGG16, ResNet50, DenseNet161, and InceptionV3.","Enabling memory-controller ECC on the A6000 prevents the observed single-bit flips at a measured cost of 3-10% inference slowdown and up to 12% bandwidth loss.","Randomizing the GPU driver's virtual-to-physical mapping, quarantining freed memory in the allocators commonly used by ML workloads, or adopting RFM/PRAC-style mitigations in GDDR would each raise the attack's cost."],"supporting_citations":[{"why":"Supplies the latency-based method for reverse-engineering DRAM bank and row addressing that GPUHammer adapts to GPU virtual addresses.","marker":"[59]"},{"why":"Establishes n-sided hammering patterns that overflow TRR sampling, the basis for the 8- to 24-sided aggressor patterns used here.","marker":"[19]"},{"why":"Provides synchronized many-sided hammering with refresh-aligned patterns, which GPUHammer adapts to multi-warp kernels via inserted delays.","marker":"[13]"},{"why":"Gives the experimental methodology and threshold baselines for characterizing Rowhammer thresholds in modern DRAM.","marker":"[37]"},{"why":"Demonstrates that a single bit-flip in the MSB of a DNN weight exponent causes severe accuracy loss, the exploit target this paper realizes on GPU memory.","marker":"[26]"},{"why":"Documents the roughly 300ns GPU memory latency that motivates the parallelized hammering kernels.","marker":"[44]"},{"why":"Defines GDDR6 timing parameters such as tRC and refresh behavior used to compute theoretical activation rates and refresh intervals.","marker":"[34]"}],"fun_headline_variants":["GPUHammer exposes first Rowhammer bit-flips in GDDR6","Rowhammer on GPUs practical: 8 bit-flips, 80% accuracy drop","Unprivileged CUDA kernel flips GDDR6 bits on A6000","GPU rowhammer attack: 8 bit-flips, ML model accuracy collapses","First practical Rowhammer on discrete GPU: GDDR6 bit-flips"],"cache_read_input_tokens":3200,"weakest_assumption_plain":"The whole attack leans on an empirical observation from a single machine: for large allocations, the address mapping that connects what a GPU program sees to the actual memory rows stays fixed across runs, so row-sets profiled once can be reused later.","fun_headline_variants_meta":{"raw":{"variants":["GPUHammer exposes first Rowhammer bit-flips in GDDR6","Rowhammer on GPUs practical: 8 bit-flips, 80% accuracy drop","Unprivileged CUDA kernel flips GDDR6 bits on A6000","GPU rowhammer attack: 8 bit-flips, ML model accuracy collapses","First practical Rowhammer on discrete GPU: GDDR6 bit-flips"]},"model":"deepseek-v4-flash","effort":"low","cost_usd":0.000902,"raw_usage":{"total_tokens":3913,"prompt_tokens":1005,"completion_tokens":2908,"prompt_tokens_details":{"cached_tokens":384},"prompt_cache_hit_tokens":384,"prompt_cache_miss_tokens":621,"completion_tokens_details":{"reasoning_tokens":2804}},"tokens_in":621,"tokens_out":2908,"duration_ms":22363,"temperature":1.0,"reasoning_tokens":2804,"cache_read_input_tokens":384,"cache_creation_input_tokens":0},"cache_creation_input_tokens":0},"created_at":"2026-08-06T18:25:32.632102+00:00","model_set":{"reader":"deepseek-v4-flash"},"falsifier":"Recompute the row-sets on the same A6000 after a reboot, a driver update, or memory pressure, then repeat the hammering campaign without re-profiling; if the previously flippable addresses no longer flip, the attack's reusable row-set assumption fails. A second decisive test is to run the same campaign on a second A6000 unit and see whether any bit-flips occur at all.","supporting_citations":[{"cited_title":"DRAMA: Exploiting DRAM addressing for Cross-CPU attacks","cited_arxiv_id":null,"evidence_quote":"Supplies the latency-based method for reverse-engineering DRAM bank and row addressing that GPUHammer adapts to GPU virtual addresses."},{"cited_title":"Trrespass: Exploiting the many sides of target row refresh","cited_arxiv_id":null,"evidence_quote":"Establishes n-sided hammering patterns that overflow TRR sampling, the basis for the 8- to 24-sided aggressor patterns used here."},{"cited_title":"SMASH: Synchronized many-sided rowhammer attacks from JavaScript","cited_arxiv_id":null,"evidence_quote":"Provides synchronized many-sided hammering with refresh-aligned patterns, which GPUHammer adapts to multi-warp kernels via inserted delays."},{"cited_title":"Kim, Minesh Patel, A","cited_arxiv_id":null,"evidence_quote":"Gives the experimental methodology and threshold baselines for characterizing Rowhammer thresholds in modern DRAM."},{"cited_title":"Terminal brain damage: Exposing the graceless degradation in deep neural net- works under hardware fault attacks","cited_arxiv_id":null,"evidence_quote":"Demonstrates that a single bit-flip in the MSB of a DNN weight exponent causes severe accuracy loss, the exploit target this paper realizes on GPU memory."},{"cited_title":"Measuring gpu memory latency","cited_arxiv_id":null,"evidence_quote":"Documents the roughly 300ns GPU memory latency that motivates the parallelized hammering kernels."},{"cited_title":"Graphics Double Data Rate (GDDR6) SGRAM Standard (JESD250D)","cited_arxiv_id":null,"evidence_quote":"Defines GDDR6 timing parameters such as tRC and refresh behavior used to compute theoretical activation rates and refresh intervals."}],"review_version":1}