[PATCH] D138439: clang: Fix cast failure when using -fsanitize=undefined for HIP

[Matt Arsenault via Phabricator via cfe-commits]

arsenm created this revision.
arsenm added reviewers: yaxunl, samsonov, tra, bkramer.
Herald added a subscriber: arichardson.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.

This was assuming a direct reference to the global variable. The
constant string is placed in addrspace 4, and has a constexpr
addrspacecast to the generic address space.


https://reviews.llvm.org/D138439

Files:
  clang/lib/CodeGen/CGExpr.cpp
  clang/test/CodeGenHIP/sanitize-undefined-null.hip


Index: clang/test/CodeGenHIP/sanitize-undefined-null.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGenHIP/sanitize-undefined-null.hip
@@ -0,0 +1,36 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -disable-llvm-passes -fcuda-is-device -fsanitize=null \
+// RUN:   -o - %s | FileCheck --enable-var-scope %s
+
+// Check there are no assertions when trying to sanitize when globals have non-0
+// address spaces.
+
+#define __device__ __attribute__((device))
+
+//.
+// CHECK: @.src = private unnamed_addr addrspace(4) constant [{{[0-9]+}} x i8] c
+// CHECK: @0 = private unnamed_addr addrspace(1) constant { i16, i16, [7 x i8] } { i16 0, i16 7, [7 x i8] c"'char'\00" }
+// CHECK: @1 = private unnamed_addr addrspace(1) global { { ptr, i32, i32 }, ptr addrspace(1), i8, i8 } { { ptr, i32, i32 } { ptr addrspacecast (ptr addrspace(4) @.src to ptr), i32 {{[0-9]+}}, i32 3 }, ptr addrspace(1) @0, i8 1, i8 1 }
+//.
+// CHECK-LABEL: @_Z3fooPc(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT:    [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// CHECK-NEXT:    store ptr [[P:%.*]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = icmp ne ptr [[TMP0]], null, !nosanitize !3
+// CHECK-NEXT:    br i1 [[TMP1]], label [[CONT:%.*]], label [[HANDLER_TYPE_MISMATCH:%.*]], !prof [[PROF4:![0-9]+]], !nosanitize !3
+// CHECK:       handler.type_mismatch:
+// CHECK-NEXT:    [[TMP2:%.*]] = ptrtoint ptr [[TMP0]] to i64, !nosanitize !3
+// CHECK-NEXT:    call void @__ubsan_handle_type_mismatch_v1_abort(ptr addrspace(1) @[[GLOB1:[0-9]+]], i64 [[TMP2]]) #[[ATTR2:[0-9]+]], !nosanitize !3
+// CHECK-NEXT:    unreachable, !nosanitize !3
+// CHECK:       cont:
+// CHECK-NEXT:    store i8 0, ptr [[TMP0]], align 1
+// CHECK-NEXT:    ret i32 3
+//
+__device__ int foo(char *p) {
+  *p = 0;
+  return 3;
+}
Index: clang/lib/CodeGen/CGExpr.cpp
===================================================================
--- clang/lib/CodeGen/CGExpr.cpp
+++ clang/lib/CodeGen/CGExpr.cpp
@@ -3157,7 +3157,8 @@
     auto FilenameGV =
         CGM.GetAddrOfConstantCString(std::string(FilenameString), ".src");
     CGM.getSanitizerMetadata()->disableSanitizerForGlobal(
-                          cast<llvm::GlobalVariable>(FilenameGV.getPointer()));
+        cast<llvm::GlobalVariable>(
+            FilenameGV.getPointer()->stripPointerCasts()));
     Filename = FilenameGV.getPointer();
     Line = PLoc.getLine();
     Column = PLoc.getColumn();
@@ -3325,13 +3326,15 @@
     // Emit handler arguments and create handler function type.
     if (!StaticArgs.empty()) {
       llvm::Constant *Info = llvm::ConstantStruct::getAnon(StaticArgs);
-      auto *InfoPtr =
-          new llvm::GlobalVariable(CGM.getModule(), Info->getType(), false,
-                                   llvm::GlobalVariable::PrivateLinkage, Info);
+      auto *InfoPtr = new llvm::GlobalVariable(
+          CGM.getModule(), Info->getType(), false,
+          llvm::GlobalVariable::PrivateLinkage, Info, "", nullptr,
+          llvm::GlobalVariable::NotThreadLocal,
+          CGM.getDataLayout().getDefaultGlobalsAddressSpace());
       InfoPtr->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
       CGM.getSanitizerMetadata()->disableSanitizerForGlobal(InfoPtr);
-      Args.push_back(Builder.CreateBitCast(InfoPtr, Int8PtrTy));
-      ArgTypes.push_back(Int8PtrTy);
+      Args.push_back(EmitCastToVoidPtr(InfoPtr));
+      ArgTypes.push_back(Args.back()->getType());
     }
 
     for (size_t i = 0, n = DynamicArgs.size(); i != n; ++i) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D138439.476910.patch
Type: text/x-patch
Size: 3951 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20221121/572cce21/attachment-0001.bin>