[clang] [llvm] [ASan] Do not instrument HIP/CUDA fatbin wrapper sections (PR #194928)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 29 18:11:07 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/194928
>From f6debd8002b9f6e16f4991201be89b9222ae46b1 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 29 Apr 2026 11:52:48 -0400
Subject: [PATCH] [ASan] Do not instrument HIP/CUDA fatbin wrapper sections
HIP and CUDA runtimes consume these sections as packed wrapper arrays, so ASan redzones would break the section ABI.
Fixes: ROCM-23813
---
clang/docs/HIPSupport.rst | 18 +++++++++
clang/lib/CodeGen/CGCUDANV.cpp | 1 +
clang/test/CodeGenCUDA/device-stub.cu | 4 ++
clang/test/Driver/linker-wrapper-image.c | 40 +++++++++++++++----
.../Frontend/Offloading/OffloadWrapper.cpp | 1 +
.../llvm-offload-wrapper/offload-wrapper.ll | 6 +--
6 files changed, 59 insertions(+), 11 deletions(-)
diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index e7f00202c34e7..82070a4042679 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -246,6 +246,24 @@ the C runtime before ``main`` and it:
during program termination and use the stored handle to unregister the fat
binary from the HIP runtime.
+The compiler-generated wrapper object is placed in the ``.hipFatBinSegment``
+section. Its layout is:
+
+.. code-block:: c++
+
+ struct {
+ uint32_t magic;
+ uint32_t version;
+ void *image;
+ void *reserved;
+ };
+
+The section is an internal ABI between Clang-generated host objects, the HIP
+runtime, and binary tools that need to find HIP fat binaries. Consumers may
+walk the raw section contents as a packed array of wrapper records, so compiler
+and instrumentation passes must preserve the wrapper object layout in the
+section.
+
The module destructor (for example ``__hip_module_dtor`` or a
``.hip.fatbin_unreg`` function) loads the stored handle, checks that it is
non-null, calls ``__hipUnregisterFatBinary`` to unregister the fat binary from
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index f08040d1d3d15..9f38aeed5b6d3 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -888,6 +888,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
/*constant*/ true);
FatbinWrapper->setSection(FatbinSectionName);
+ CGM.getSanitizerMetadata()->disableSanitizerForGlobal(FatbinWrapper);
// There is only one HIP fat binary per linked module, however there are
// multiple constructor functions. Make sure the fat binary is registered
diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu
index 8695433f6df10..38c7dc711ef1d 100644
--- a/clang/test/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CodeGenCUDA/device-stub.cu
@@ -45,6 +45,9 @@
// RUN: -fcuda-include-gpubinary %t -o - -x hip\
// RUN: | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
+// RUN: -fcuda-include-gpubinary %t -fsanitize=address -o - -x hip\
+// RUN: | FileCheck %s --check-prefix=HIP-ASAN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
// RUN: -fcuda-include-gpubinary %t -o - -DNOGLOBALS -x hip \
// RUN: | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
@@ -173,6 +176,7 @@ __device__ void device_use() {
// LNX-SAME: ptr null }
// CUDA-SAME: section ".nvFatBinSegment"
// HIP-SAME: section ".hipFatBinSegment"
+// HIP-ASAN: @__hip_fatbin_wrapper = internal constant {{.*}}, section ".hipFatBinSegment", no_sanitize_address
// * variable to save GPU binary handle after initialization
// CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global ptr null
// HIPNEF: @__[[PREFIX]]_gpubin_handle_{{[0-9a-f]+}} = internal global ptr null, align 8
diff --git a/clang/test/Driver/linker-wrapper-image.c b/clang/test/Driver/linker-wrapper-image.c
index dd346bce225e8..9d3c65fd36760 100644
--- a/clang/test/Driver/linker-wrapper-image.c
+++ b/clang/test/Driver/linker-wrapper-image.c
@@ -59,8 +59,14 @@
// CUDA-COFF: @__start_llvm_offload_entries = weak_odr hidden constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "llvm_offload_entries$OA"
// CUDA-COFF-NEXT: @__stop_llvm_offload_entries = weak_odr hidden constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "llvm_offload_entries$OZ"
-// CUDA: @.fatbin_image = internal constant [0 x i8] zeroinitializer, section ".nv_fatbin"
-// CUDA-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1180844977, i32 1, ptr @.fatbin_image, ptr null }, section ".nvFatBinSegment", align 8
+// CUDA: @.fatbin_image = internal constant
+// CUDA-SAME: section ".nv_fatbin"
+// CUDA-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper
+// CUDA-SAME: { i32 1180844977, i32 1, ptr @.fatbin_image, ptr null }
+// CUDA-SAME: section ".nvFatBinSegment"
+// CUDA-SAME: no_sanitize_address
+// CUDA-SAME: no_sanitize_hwaddress
+// CUDA-SAME: align 8
// CUDA-NEXT: @.cuda.binary_handle = internal global ptr null
// CUDA: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 101, ptr @.cuda.fatbin_reg, ptr null }]
@@ -176,16 +182,34 @@
// HIP-MACHO-NEXT: @"\01section$end$__LLVM$offload_entries" = external hidden constant [0 x %struct.__tgt_offload_entry]
// HIP-MACHO-NEXT: @"__dummy.__LLVM,offload_entries" = internal constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "__LLVM,offload_entries"
-// HIP-ELF: @.fatbin_image = internal constant [0 x i8] zeroinitializer, section ".hip_fatbin"
-// HIP-ELF-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section ".hipFatBinSegment", align 8
+// HIP-ELF: @.fatbin_image = internal constant
+// HIP-ELF-SAME: section ".hip_fatbin"
+// HIP-ELF-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper
+// HIP-ELF-SAME: { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }
+// HIP-ELF-SAME: section ".hipFatBinSegment"
+// HIP-ELF-SAME: no_sanitize_address
+// HIP-ELF-SAME: no_sanitize_hwaddress
+// HIP-ELF-SAME: align 8
// HIP-ELF-NEXT: @.hip.binary_handle = internal global ptr null
-// HIP-COFF: @.fatbin_image = internal constant [0 x i8] zeroinitializer, section ".hip_fatbin"
-// HIP-COFF-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section ".hipFatBinSegment", align 8
+// HIP-COFF: @.fatbin_image = internal constant
+// HIP-COFF-SAME: section ".hip_fatbin"
+// HIP-COFF-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper
+// HIP-COFF-SAME: { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }
+// HIP-COFF-SAME: section ".hipFatBinSegment"
+// HIP-COFF-SAME: no_sanitize_address
+// HIP-COFF-SAME: no_sanitize_hwaddress
+// HIP-COFF-SAME: align 8
// HIP-COFF-NEXT: @.hip.binary_handle = internal global ptr null
-// HIP-MACHO: @.fatbin_image = internal constant [0 x i8] zeroinitializer, section "__HIP,__hip_fatbin"
-// HIP-MACHO-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section "__HIP,__fatbin", align 8
+// HIP-MACHO: @.fatbin_image = internal constant
+// HIP-MACHO-SAME: section "__HIP,__hip_fatbin"
+// HIP-MACHO-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper
+// HIP-MACHO-SAME: { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }
+// HIP-MACHO-SAME: section "__HIP,__fatbin"
+// HIP-MACHO-SAME: no_sanitize_address
+// HIP-MACHO-SAME: no_sanitize_hwaddress
+// HIP-MACHO-SAME: align 8
// HIP-MACHO-NEXT: @.hip.binary_handle = internal global ptr null
// HIP: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 101, ptr @.hip.fatbin_reg, ptr null }]
diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
index b6b25c7bbad19..ded603a1e00e3 100644
--- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
@@ -318,6 +318,7 @@ GlobalVariable *createFatbinDesc(Module &M, ArrayRef<char> Image, bool IsHIP,
FatbinInitializer, ".fatbin_wrapper" + Suffix);
FatbinDesc->setSection(FatbinWrapperSection);
FatbinDesc->setAlignment(Align(8));
+ FatbinDesc->setNoSanitizeMetadata();
return FatbinDesc;
}
diff --git a/llvm/test/tools/llvm-offload-wrapper/offload-wrapper.ll b/llvm/test/tools/llvm-offload-wrapper/offload-wrapper.ll
index b9d2d86eda375..e3d507ba50547 100644
--- a/llvm/test/tools/llvm-offload-wrapper/offload-wrapper.ll
+++ b/llvm/test/tools/llvm-offload-wrapper/offload-wrapper.ll
@@ -31,7 +31,7 @@
; HIP-NEXT: @__dummy.llvm_offload_entries = internal constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "llvm_offload_entries", align 8
; HIP-NEXT: @llvm.compiler.used = appending global [1 x ptr] [ptr @__dummy.llvm_offload_entries], section "llvm.metadata"
; HIP-NEXT: @.fatbin_image = internal constant {{.*}}, section ".hip_fatbin"
-; HIP-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section ".hipFatBinSegment", align 8
+; HIP-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section ".hipFatBinSegment", no_sanitize_address, no_sanitize_hwaddress, align 8
; HIP-NEXT: @.hip.binary_handle = internal global ptr null
; HIP-NEXT: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 101, ptr @.hip.fatbin_reg, ptr null }]
@@ -59,7 +59,7 @@
; HIP-MACHO-NEXT: @"__dummy.__LLVM,offload_entries" = internal constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "__LLVM,offload_entries", align 8
; HIP-MACHO-NEXT: @llvm.compiler.used = appending global [1 x ptr] [ptr @"__dummy.__LLVM,offload_entries"], section "llvm.metadata"
; HIP-MACHO-NEXT: @.fatbin_image = internal constant {{.*}}, section "__HIP,__hip_fatbin"
-; HIP-MACHO-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section "__HIP,__fatbin", align 8
+; HIP-MACHO-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1212764230, i32 1, ptr @.fatbin_image, ptr null }, section "__HIP,__fatbin", no_sanitize_address, no_sanitize_hwaddress, align 8
; HIP-MACHO-NEXT: @.hip.binary_handle = internal global ptr null
; HIP-MACHO-NEXT: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 101, ptr @.hip.fatbin_reg, ptr null }]
@@ -87,7 +87,7 @@
; CUDA-NEXT: @__dummy.llvm_offload_entries = internal constant [0 x %struct.__tgt_offload_entry] zeroinitializer, section "llvm_offload_entries", align 8
; CUDA-NEXT: @llvm.compiler.used = appending global [1 x ptr] [ptr @__dummy.llvm_offload_entries], section "llvm.metadata"
; CUDA-NEXT: @.fatbin_image = internal constant {{.*}}, section ".nv_fatbin"
-; CUDA-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1180844977, i32 1, ptr @.fatbin_image, ptr null }, section ".nvFatBinSegment", align 8
+; CUDA-NEXT: @.fatbin_wrapper = internal constant %fatbin_wrapper { i32 1180844977, i32 1, ptr @.fatbin_image, ptr null }, section ".nvFatBinSegment", no_sanitize_address, no_sanitize_hwaddress, align 8
; CUDA-NEXT: @.cuda.binary_handle = internal global ptr null
; CUDA-NEXT: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 101, ptr @.cuda.fatbin_reg, ptr null }]
More information about the cfe-commits
mailing list