Exploring User Security and Privacy Attitudes and Concerns Toward the Use of General-Purpose LLM Chatbots for Mental Health

GPUHammer: Rowhammer Attacks on GPU Memories are Practical

Speakers: Chris S. Lin (University of Toronto); Joyce Qu (University of Toronto); Gururaj Saileshwar (University of Toronto)

Conference: USENIX Security

Paper Page: https://www.usenix.org/conference/usenixsecurity25/presentation/lin-shaopeng

Overview

The USENIX Security paper "GPUHammer: Rowhammer Attacks on GPU Memories are Practical" unveils a groundbreaking discovery: the first successful Rowhammer attack against discrete NVIDIA GPUs utilizing GDDR6 DRAM. Authored by Chris S. Lin, Joyce Qu, and Gururaj Saileshwar from the University of Toronto, this research addresses a critical gap in hardware security, extending the well-known Rowhammer vulnerability from traditional CPU-based DDR/LPDDR memories to the high-performance memory architectures of modern GPUs. With GPUs forming the backbone of emerging machine learning (ML) and high-performance computing (HPC) applications, understanding and mitigating such vulnerabilities is paramount for ensuring the integrity and security of these critical systems.

The paper meticulously details the unique challenges encountered when attempting Rowhammer attacks on GPUs, including proprietary memory mappings, higher memory latencies, faster refresh rates, and unknown in-DRAM mitigations. To overcome these hurdles, the researchers developed GPUHammer, a sophisticated attack framework that introduces novel techniques for reverse-engineering GDDR DRAM row mappings, optimizing memory access patterns for amplified hammering intensity, and precisely synchronizing attacks to bypass proprietary defenses. The efficacy of GPUHammer is demonstrated on an NVIDIA A6000 GPU with GDDR6 memory, where it successfully induces bit-flips and, more significantly, leverages these bit-flips to tamper with machine learning models, causing drastic accuracy degradation.

This work not only exposes a significant hardware vulnerability in widely deployed GPU architectures but also highlights the pressing need for robust system-level mitigations and algorithmic resilience in ML models. The findings have profound implications for cloud computing environments, where GPUs are often shared among multiple tenants, and for any application relying on the integrity of GPU memory. The responsible disclosure to NVIDIA and major cloud service providers underscores the severity and real-world applicability of this research, paving the way for future defensive strategies against this new class of hardware attacks.

Background

Rowhammer is a hardware vulnerability inherent in modern DRAM (Dynamic Random-Access Memory), where repeatedly accessing (or "hammering") a specific memory row, known as an aggressor row, can cause electrical disturbance that leads to unintended bit-flips in physically adjacent victim rows. This phenomenon, first publicly demonstrated in 2014, arises from the increasing density and shrinking cell sizes in DRAM, leading to charge leakage between closely packed memory cells. On CPUs, Rowhammer has been extensively studied on DDR and LPDDR memories, demonstrating its potential for severe security compromises, including data tampering and even kernel-level privilege escalation.

Despite the widespread recognition of Rowhammer on CPUs, its impact on GPUs, particularly those using GDDR (Graphics Double Data Rate) and HBM (High Bandwidth Memory), remained largely unexplored. This oversight is significant, as GPUs have become indispensable for high-performance computing, artificial intelligence, and graphics processing. The distinct architectural characteristics of GPUs, designed for throughput rather than latency, present several unique challenges for adapting existing Rowhammer attacks:

1. Proprietary Memory Mappings (C1): The mapping of virtual addresses to physical DRAM rows and banks on NVIDIA GPUs is proprietary and undocumented, making it difficult to identify addresses that conflict within the same DRAM bank—a prerequisite for Rowhammer.
2. High Memory Latency and Faster Refresh Rates (C2): GPUs typically exhibit higher memory latencies (up to 4x that of CPUs) and shorter DRAM refresh periods (e.g., 22ms for GDDR6 compared to 32-64ms for DDR5). These factors limit the achievable hammering intensity and the available time window to induce bit-flips before memory cells are refreshed.
3. Proprietary In-DRAM Mitigations (C3): GDDR memories likely incorporate proprietary in-DRAM defenses, such as Target Row Refresh (TRR), similar to those found in DDR4/DDR5. These mitigations track frequently accessed rows and proactively refresh neighbors to prevent bit-flips, requiring sophisticated attack patterns to bypass.

The motivation for studying GPU Rowhammer is further amplified by the increasing reliance on GPUs for machine learning inference. Prior research has shown that ML models are highly susceptible to bit-flips, which can lead to significant accuracy degradation, backdoor injection, or other malicious tampering. Given that ML inference is predominantly performed on GPUs, understanding their vulnerability to Rowhammer is critical for securing the entire ML ecosystem.

The attack model assumes an attacker capable of executing CUDA kernels with user-level privileges on the GPU. This is particularly relevant in multi-tenant cloud environments where GPUs are often time-sliced or spatially shared, allowing for co-location of an attacker's process with a victim's sensitive data within the same DRAM bank. The paper explicitly states that ECC (Error-Correcting Code) is assumed to be disabled, a common configuration in cloud environments where performance is prioritized, and ECC can introduce slowdowns and reduce memory capacity.

Key Findings

The GPUHammer research presents several pivotal findings that collectively demonstrate the practicality and impact of Rowhammer attacks on discrete NVIDIA GPUs:

• First Practical Rowhammer on Discrete GPUs: GPUHammer achieves the first successful Rowhammer attack on a discrete GPU, specifically targeting an NVIDIA A6000 with GDDR6 memory. This marks a significant expansion of the Rowhammer threat landscape beyond CPU-based systems.
• Bit-Flip Induction: The attack successfully injected up to 8 bit-flips across 4 different DRAM banks on the A6000, with at least one bit-flip observed in each hammered bank. All observed bit-flips were single-bit flips.
• Novel Reverse-Engineering Techniques: The researchers developed a method to directly reverse-engineer the mapping of virtual addresses to DRAM banks and rows without requiring access to physical addresses, which are typically hidden on NVIDIA GPUs. This involved using latency measurements from row-buffer conflicts to build a lookup table, overcoming a major challenge.
• High Activation Rate Hammering: GPUHammer introduces multi-warp hammering, a GPU-specific memory access optimization that leverages the parallel execution model of GPUs. This technique achieved activation rates of up to 620,000 activations per refresh window (tREFW), nearly 7 times higher than naive single-thread hammering and close to the theoretical maximum for GDDR6. This high intensity is crucial for crossing the Rowhammer Threshold (TRH) within the shorter refresh periods of GDDR6.
• Bypassing In-DRAM Mitigations: The attack employs a novel per-warp synchronization design to align hammering patterns with DRAM refresh commands (REFs). This synchronization is critical for consistently bypassing proprietary in-DRAM mitigations, such as TRR, by forcing the mitigation sampler into predictable states where specific aggressor rows evade tracking.
• Characterization of GDDR6 Properties: The research reverse-engineered key GDDR6 memory parameters, including a refresh interval (tREFI) of 1407ns (resulting in a refresh period of approximately 23ms) and a minimum TRH of 12.3K activations for the A6000.
• TRR Sampler Size Discovery: Through systematic hammering campaigns, the paper inferred the presence and capacity of a TRR-like mitigation in A6000's GDDR6 memory. Bit-flips were only consistently observed with 17 or more aggressor rows, suggesting the TRR sampler can track a maximum of 16 rows per bank.
• ML Model Accuracy Degradation Exploit: The most impactful finding is the demonstration of an end-to-end exploit targeting Deep Neural Network (DNN) weights. By inducing a single 0→1 bit-flip in the Most Significant Bit (MSB) of the exponent (E4) of FP16-represented weights, GPUHammer achieved a Relative Accuracy Drop (RAD) of up to 99% on prominent ImageNet models like AlexNet, VGG16, ResNet50, DenseNet161, and InceptionV3. This attack was shown to be highly efficient, with significant degradation (20% RAD) after just 2 attempts and over 99% RAD after 8 attempts.
• Non-Contiguous Row Layout: Analysis of critical aggressor rows suggested that physical DRAM rows might not be linearly laid out, with some bit-flips occurring more reliably from aggressors at Ri±2 rather than Ri±1, indicating internal remapping.
• Absence of Flips on Other GPUs: While successful on the A6000, no bit-flips were observed on an NVIDIA A100 (HBM2e) or an RTX 3080 (GDDR6). This is attributed to potential differences in TRH, architecture, or the presence of on-die ECC in HBM2e.

These findings collectively establish Rowhammer as a practical and severe threat to modern GPU systems, particularly those involved in sensitive ML workloads.

Technical Deep Dive

The success of GPUHammer hinges on overcoming several architectural and behavioral challenges inherent to NVIDIA GPUs and GDDR6 memory. The core technical contributions lie in establishing the necessary primitives for Rowhammer on GPUs, reverse-engineering memory mappings, achieving high activation rates, and precisely synchronizing with DRAM refresh commands.

Building Blocks of GPUHammer

A Rowhammer attack requires two fundamental primitives: cache eviction and row-buffer eviction.

1. Cache Eviction: To ensure every memory access hits the DRAM and triggers a row activation, data must be evicted from the GPU's on-chip L1 and L2 caches. For Ampere generation GPUs (sm_80 and later), NVIDIA introduced the discard instruction, which clears an address from the L2 cache. However, to bypass the L1 cache, the ld.global.volatile modifier for load instructions is crucial. The combination of discard.global.L2 and ld.u64.global.volatile ensures that memory accesses consistently incur the full DRAM latency (around 300ns on an A6000), confirming they bypass caches.

1. Row Buffer Eviction: Once data is fetched into the DRAM's row buffer, subsequent accesses to the same row will hit the row buffer, not trigger a new row activation. To force new activations, the current row must be evicted from the row buffer. This is achieved by accessing a different row within the same DRAM bank, which causes a row-buffer conflict and forces the current row to close. This mechanism is central to generating the rapid row activations needed for Rowhammer.

These primitives are integrated into a CUDA kernel that performs an n-sided Rowhammer attack, iteratively accessing N distinct aggressor addresses.

Reversing GPU DRAM Row Addressing

A major challenge (C1) is the proprietary and undocumented nature of NVIDIA's virtual-to-physical address mapping and physical-to-DRAM bank/row mapping. Unlike CPUs where physical addresses are accessible, GPUs obscure this information. GPUHammer addresses this by directly learning a lookup table that maps virtual addresses to DRAM bank and row IDs, bypassing the need to reverse-engineer the exact closed-form function.

The process involves:

1. Large Array Allocation: Allocating a large array (e.g., 47GB on a 48GB A6000) using cudaMalloc to span most of the GPU DRAM.
2. Conflict Testing: Fixing a reference address and iteratively testing other addresses (at 64B granularity initially, later optimized to 256B) for row-buffer conflicts. A higher memory access latency (e.g., 360-380ns vs. 320-370ns for non-conflicting) indicates a conflict within the same bank.
3. NUMA Effect Mitigation: The researchers identified Non-Uniform Memory Access (NUMA) effects in the A6000, causing latency variations. To filter out noise, they characterized single-access latencies for all addresses and only considered addresses with similar single-access latency (within ±10ns) to the reference address for conflict testing. This ensures that only addresses truly belonging to the same memory region/chip are compared.
4. Conflict-Set and Row-Set Generation: Addresses exhibiting conflicts are grouped into a Conflict-Set. From this, a Row-Set is derived by filtering out duplicates, ensuring each address in the Row-Set maps to a unique row within the bank.
5. Observations: This process revealed that on A6000 GDDR6, contiguous 256-byte chunks map to a single row/bank, and the row size is 2KB with 64K rows per bank. Crucially, for large allocations, the virtual-to-physical mapping remains consistent across reboots, allowing the generated Row-Sets to be reused.

Hammering with High Activation Rates

Achieving a sufficiently high activation rate (C2) is critical given the high memory latency (around 300ns) and short refresh period (23ms) of GDDR6. A naive single-thread hammering kernel achieves only ~106,000 activations in 32ms, which is insufficient for n-sided attacks where each aggressor receives far fewer activations (e.g., ~5K per row for a 20-sided attack, below the TRH of 10K-20K for modern DRAM).

GPUHammer overcomes this by exploiting GPU parallelism:

1. Multi-Thread Hammering: Assigning each aggressor to a dedicated thread within a single warp. While improving rates, lock-step execution still introduces idle time.
2. Multi-Warp Hammering: This is the key optimization. Each aggressor is handled by a thread in a different warp. When one warp stalls due to memory access, the Streaming Multiprocessor (SM) can schedule another warp, effectively masking memory latency and minimizing idle time at the memory controller.

This technique achieved 620,000 activations per tREFW (32ms) with just 8 aggressors, close to the theoretical maximum of 700,000 activations. This rate is nearly 7 times higher than single-thread hammering and ensures that individual aggressor rows receive close to 20K activations per tREFW even with many aggressors, exceeding the observed TRH.

Synchronization with Refresh Commands

Bypassing in-DRAM mitigations like TRR (C3) requires not just high activation rates but also synchronization of hammering patterns with DRAM refresh commands (REFs). TRR proactively refreshes victim rows, but its sampler can be fooled if the attack consistently presents aggressors in a way that some escape tracking. Prior CPU attacks synchronized by inserting NOPs, but __syncthreads on GPUs introduces unpredictable warp ordering.

GPUHammer's Per-Warp Synchronization Design addresses this:

1. It avoids __syncthreads to maintain warp order.
2. Instead, add instructions are inserted within each warp after a fixed number of hammering iterations (e.g., 1, 2, or 3 iterations depending on the n-sided pattern).
3. When these delays overlap across all active warps, an implicit "bubble" is created at the memory controller, allowing a REF command to be inserted in alignment with the hammering pattern.

This method ensures synchronization for various n-sided patterns (e.g., 8, 12, 16, 24 aggressors) by using k warps and m threads per warp, where k is kept at 8 or lower to maintain sufficient delay overlap. This synchronization allowed the researchers to determine that the tREFI for A6000 GDDR6 is 1407ns, leading to a full refresh period of 23ms. With this, the synchronized n-sided hammering can deliver up to 24 aggressor activations per tREFI, totaling 393K activations per refresh period.

Demo / Proof of Concept

GPUHammer culminates in a compelling proof-of-concept demonstrating the real-world impact of Rowhammer bit-flips on an NVIDIA A6000 GPU: tampering with Deep Neural Network (DNN) weights during inference. This exploit targets ML models to cause significant accuracy degradation, similar to previously simulated attacks like Terminal Brain Damage (TBD).

The attack involves:

1. Threat Model: Assumes a multi-tenant GPU environment where the attacker and victim are co-located and share the GPU in a time-multiplexed manner (e.g., 250ms time-slices enabled by NVIDIA GPU Operator/RunAI). ECC is disabled.
2. Profiling for Bit-Flips: The attacker first uses GPUHammer primitives to profile their own allocated memory, identifying the precise locations and directions (0→1 or 1→0) of vulnerable bit-flips.
3. Memory Massaging: To ensure victim data (DNN weights) maps to these vulnerable locations, the attacker exploits the behavior of the RAPIDS Memory Manager (RMM), a widely used, speed-optimized allocator in ML applications. The attacker allocates large memory regions, then frees a specific chunk containing a known flippy bit. Because RMM immediately reuses freed memory, when the victim (e.g., a PyTorch ML inference process) requests memory, that chunk is likely reallocated to the victim. RAPIDS' 256-byte allocation granularity provides fine-grained control over placement.
4. Target and Metrics: The attack targets FP16 weights of pre-trained ImageNet models (AlexNet, VGG16, ResNet50, DenseNet161, InceptionV3) running in PyTorch. The impact is measured by Relative Accuracy Drop (RAD), calculated as (Acc_pristine - Acc_corrupted) / Acc_pristine.
5. Exploit Execution and Results: The researchers conducted 50 attempts for each model, randomly shifting victim memory to map different weights to the vulnerable bit. The most impactful bit-flips were D1 and D3, which were 0→1 flips that mapped to the Most Significant Bit (MSB) of the exponent (E4) in the FP16 representation. A flip in this bit exponentially increases the weight's value, causing severe corruption.

• For bit-flips D1 and D3, the models' top-1 accuracy dropped from an average of 71.26% to less than 0.1%, resulting in an average RAD of 0.99 (99%).
• Other bit-flips (B1, B2) mapping to the mantissa had negligible impact.
• The attack was highly efficient: an average RAD of 20% was observed after just 2 exploit attempts, and over 99% RAD was achieved after only 8 attempts.
• While FP32 models were not successfully exploited for MSB flips in this study, the paper notes that with further profiling, such exploits are likely possible.

This demonstration provides concrete evidence that Rowhammer on GPUs is not just a theoretical concern but a practical vector for compromising critical ML workloads, with devastating consequences for model integrity and reliability.

Defensive Implications

The GPUHammer research highlights a critical security vulnerability and thus necessitates a multi-faceted approach to mitigation. Several strategies, both existing and proposed, can be employed to defend against such attacks:

1. Enabling ECC (Error-Correcting Code): For GDDR-based GPUs like the A6000, enabling memory-controller-based ECC is the most direct mitigation. ECC can detect and correct single bit-flips, effectively neutralizing the observed GPUHammer attacks. However, ECC is often disabled by default in these systems due to its associated costs: a 6.5% memory overhead and a 3% to 12% performance slowdown in memory bandwidth and ML applications. This trade-off between security and performance is a significant consideration for cloud providers and high-performance computing users.
2. Randomizing Virtual to Physical Mappings: GPUHammer relies on the consistency of virtual-to-physical memory mappings for large allocations to reuse pre-profiled row-sets. If the GPU driver were to randomize these mappings on each program run, attackers would be forced to re-profile the DRAM row mappings every time, significantly increasing the cost and complexity of the attack. While effective, this approach could potentially introduce performance overheads due to increased TLB (Translation Lookaside Buffer) fragmentation.
3. Making Memory Massaging Unpredictable: The current exploit leverages the RAPIDS Memory Manager's aggressive memory reuse policy to precisely map victim data to vulnerable locations. Adopting memory allocators that quarantine deallocated memory for a period or randomize reallocation patterns would make such precise memory massaging much harder. However, these techniques typically come with increased memory usage and potentially reduced performance, making them less desirable for performance-critical GPU workloads.
4. Adopting Modern Rowhammer Mitigations: Beyond the basic TRR-like mitigation identified in GDDR6, more advanced hardware-based Rowhammer defenses have been proposed for newer DRAM standards. These include Refresh Management (RFM) for HBM3 and DDR5, and Per Row Activation Counter (PRAC) specified in DDR5. Integrating such sophisticated mitigations into emerging GDDRx memories (e.g., GDDR7) is crucial to proactively address the evolving Rowhammer threat.
5. NVIDIA MIG and Confidential Computing: NVIDIA's Multi-instance GPU (MIG) and Confidential Computing (CC) offerings for cloud-based GPUs inherently provide some protection against the multi-tenant data tampering aspect of GPUHammer. MIG spatially partitions GPU memory, assigning separate memory channels to different users, while CC assigns entire GPUs to confidential VMs. Both configurations prevent the co-location of multi-tenant data within the same DRAM bank, which is a prerequisite for GPUHammer's exploits. However, the paper notes that Rowhammer exploits in such environments would require further exploration, implying that while direct data tampering between tenants might be thwarted, other attack vectors might still exist.
6. Algorithmic Resilience for ML Models: Beyond hardware and system-level mitigations, the vulnerability of ML models to bit-flips also calls for algorithmic solutions. Techniques like using early-exit mechanisms, honeypot neurons, or checksums for model weights can help detect or even tolerate minor corruptions, reducing the impact of successful bit-flips.

The decision on which mitigations to deploy will involve careful consideration of the security posture, performance requirements, and cost constraints of specific GPU deployments.

Key Takeaways

• Discrete NVIDIA GPUs are Vulnerable: GPUHammer demonstrates the first practical Rowhammer attack on discrete NVIDIA GPUs, specifically the A6000 with GDDR6 memory, capable of inducing bit-flips across multiple DRAM banks.
• Novel Techniques Overcome GPU Challenges: The attack overcomes unique GPU architectural challenges—proprietary address mappings, high memory latency, and faster refresh rates—through novel techniques like direct virtual-to-bank/row mapping, highly parallelized multi-warp hammering (achieving 620K activations per tREFW), and precise per-warp synchronization with DRAM refresh commands.
• ML Model Integrity at Risk: A single Rowhammer-induced bit-flip in the Most Significant Bit (MSB) of the exponent of FP16-represented DNN weights can cause catastrophic accuracy degradation (up to 99% Relative Accuracy Drop) across various state-of-the-art ML models.
• In-DRAM Mitigations are Bypassed: The GDDR6 memory in the A6000 contains a TRR-like mitigation that tracks approximately 16 rows per bank. GPUHammer successfully bypasses this defense using synchronized n-sided hammering patterns with 17 or more aggressors.
• Mitigations Exist with Trade-offs: While defenses like enabling ECC, randomizing virtual-to-physical mappings, and employing unpredictable memory allocators can mitigate the threat, they often come with performance overheads (e.g., up to 12% slowdown for ECC) or increased resource consumption.
• Urgent Need for GPU Security Enhancements: The findings underscore the critical need for enhanced hardware-level security in GPUs and for ML systems to develop greater resilience against physical memory attacks, especially given the widespread deployment of GPUs in multi-tenant cloud environments.

About the Speaker(s)

The research paper "GPUHammer: Rowhammer Attacks on GPU Memories are Practical" was authored by Chris S. Lin, Joyce Qu, and Gururaj Saileshwar, all affiliated with the University of Toronto. Their work contributes significantly to the field of hardware security, specifically addressing the under-explored area of Rowhammer vulnerabilities in graphics processing units. While the format specifies "speaker(s)," this is a peer-reviewed paper, and the authors are the primary contributors to the research and its presentation. Their collective expertise has brought to light a critical security concern in modern GPU architectures that has far-reaching implications for high-performance computing and artificial intelligence.