37 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -disable-llvm-passes -fcuda-is-device -fsanitize=null \3// RUN: -o - %s | FileCheck --enable-var-scope %s4 5// Check there are no assertions when trying to sanitize when globals have non-06// address spaces.7 8#define __device__ __attribute__((device))9 10//.11// CHECK: @.src = private unnamed_addr addrspace(4) constant [{{[0-9]+}} x i8] c12// CHECK: @0 = private unnamed_addr addrspace(1) constant { i16, i16, [7 x i8] } { i16 0, i16 7, [7 x i8] c"'char'\00" }13// 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 }14//.15// CHECK-LABEL: @_Z3fooPc(16// CHECK-NEXT: entry:17// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)18// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)19// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr20// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr21// CHECK-NEXT: store ptr [[P:%.*]], ptr [[P_ADDR_ASCAST]], align 822// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 823// CHECK-NEXT: [[TMP1:%.*]] = icmp ne ptr [[TMP0]], null, !nosanitize !424// CHECK-NEXT: br i1 [[TMP1]], label [[CONT:%.*]], label [[HANDLER_TYPE_MISMATCH:%.*]], !prof [[PROF4:![0-9]+]], !nosanitize !425// CHECK: handler.type_mismatch:26// CHECK-NEXT: [[TMP2:%.*]] = ptrtoint ptr [[TMP0]] to i64, !nosanitize !427// CHECK-NEXT: call void @__ubsan_handle_type_mismatch_v1_abort(ptr addrspace(1) @[[GLOB1:[0-9]+]], i64 [[TMP2]]) #[[ATTR2:[0-9]+]], !nosanitize !428// CHECK-NEXT: unreachable, !nosanitize !429// CHECK: cont:30// CHECK-NEXT: store i8 0, ptr [[TMP0]], align 131// CHECK-NEXT: ret i32 332//33__device__ int foo(char *p) {34 *p = 0;35 return 3;36}37