Update lit

Author: haonanya1

llvm/include/llvm/Transforms/Instrumentation/SanitizerCommonUtils.h

@@ -29,10 +29,6 @@ constexpr unsigned kSpirOffloadGenericAS = 4;
 
 TargetExtType *getTargetExtType(Type *Ty);
 bool isJointMatrixAccess(Value *V);
-bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I,
-                             bool InstrumentLocalPtr,
-                             bool InstrumentPrivatePtr);
-
 } // namespace SanitizerCommonUtils
 } // namespace llvm
 

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1619,6 +1619,53 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) {
   return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false");
 }
 
+static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) {
+  // Skip SPIR-V built-in varibles
+  auto *OrigValue = Addr->stripInBoundsOffsets();
+  if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
+    return true;
+
+  GlobalVariable *GV = dyn_cast<GlobalVariable>(OrigValue);
+  if (GV && isUnsupportedDeviceGlobal(GV))
+    return true;
+
+  // Ignore load/store for target ext type since we can't know exactly what size
+  // it is.
+  if (auto *SI = dyn_cast<StoreInst>(Inst))
+    if (SanitizerCommonUtils::getTargetExtType(
+            SI->getValueOperand()->getType()) ||
+        SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand()))
+      return true;
+
+  if (auto *LI = dyn_cast<LoadInst>(Inst))
+    if (SanitizerCommonUtils::getTargetExtType(Inst->getType()) ||
+        SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand()))
+      return true;
+
+  Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
+  switch (PtrTy->getPointerAddressSpace()) {
+  case SanitizerCommonUtils::kSpirOffloadPrivateAS: {
+    if (!ClSpirOffloadPrivates)
+      return true;
+    // Skip kernel arguments
+    return Inst->getFunction()->getCallingConv() == CallingConv::SPIR_KERNEL &&
+           isa<Argument>(Addr);
+  }
+  case SanitizerCommonUtils::kSpirOffloadGlobalAS: {
+    return !ClSpirOffloadGlobals;
+  }
+  case SanitizerCommonUtils::kSpirOffloadLocalAS: {
+    if (!ClSpirOffloadLocals)
+      return true;
+    return Addr->getName().starts_with("__Asan");
+  }
+  case SanitizerCommonUtils::kSpirOffloadGenericAS: {
+    return !ClSpirOffloadGenerics;
+  }
+  }
+  return true;
+}
+
 void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore,
                                              Value *Addr,
                                              SmallVectorImpl<Value *> &Args) {
@@ -1847,8 +1894,7 @@ bool AddressSanitizer::isInterestingAlloca(const AllocaInst &AI) {
 bool AddressSanitizer::ignoreAccess(Instruction *Inst, Value *Ptr) {
   // SPIR has its own rules to filter the instrument accesses
   if (TargetTriple.isSPIROrSPIRV()) {
-    if (SanitizerCommonUtils::isUnsupportedSPIRAccess(
-            Ptr, Inst, ClSpirOffloadLocals, ClSpirOffloadPrivates))
+    if (isUnsupportedSPIRAccess(Ptr, Inst))
       return true;
   } else {
     // Instrument accesses from different address spaces only for AMDGPU.

llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp

@@ -1729,6 +1729,44 @@ static unsigned TypeSizeToSizeIndex(TypeSize TS) {
   return Log2_32_Ceil((TypeSizeFixed + 7) / 8);
 }
 
+static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) {
+  if (isa<Instruction>(Addr) &&
+      cast<Instruction>(Addr)->getMetadata(LLVMContext::MD_nosanitize)) {
+    return true;
+  }
+
+  // Skip SPIR-V built-in varibles
+  auto *OrigValue = Addr->stripInBoundsOffsets();
+  assert(OrigValue != nullptr);
+  if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
+    return true;
+
+  // Ignore load/store for target ext type since we can't know exactly what size
+  // it is.
+  if (auto *SI = dyn_cast<StoreInst>(I))
+    if (SanitizerCommonUtils::getTargetExtType(
+            SI->getValueOperand()->getType()) ||
+        SanitizerCommonUtils::isJointMatrixAccess(SI->getPointerOperand()))
+      return true;
+
+  if (auto *LI = dyn_cast<LoadInst>(I))
+    if (SanitizerCommonUtils::getTargetExtType(I->getType()) ||
+        SanitizerCommonUtils::isJointMatrixAccess(LI->getPointerOperand()))
+      return true;
+
+  Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
+  switch (PtrTy->getPointerAddressSpace()) {
+  case kSpirOffloadPrivateAS:
+    return !ClSpirOffloadPrivates;
+  case kSpirOffloadLocalAS:
+    return !ClSpirOffloadLocals;
+  case kSpirOffloadGenericAS:
+    return false;
+  }
+
+  return false;
+}
+
 static void setNoSanitizedMetadataSPIR(Instruction &I) {
   const Value *Addr = nullptr;
   if (const auto *LI = dyn_cast<LoadInst>(&I))
@@ -1780,8 +1818,7 @@ static void setNoSanitizedMetadataSPIR(Instruction &I) {
     }
   }
 
-  if (Addr && SanitizerCommonUtils::isUnsupportedSPIRAccess(
-                  Addr, &I, ClSpirOffloadLocals, ClSpirOffloadPrivates))
+  if (Addr && isUnsupportedSPIRAccess(Addr, &I))
     I.setNoSanitizeMetadata();
 }
 

llvm/lib/Transforms/Instrumentation/SanitizerCommonUtils.cpp

@@ -59,45 +59,5 @@ bool isJointMatrixAccess(Value *V) {
   }
   return false;
 }
-
-bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I,
-                             bool InstrumentLocalPtr,
-                             bool InstrumentPrivatePtr) {
-  if (isa<Instruction>(Addr) &&
-      cast<Instruction>(Addr)->getMetadata(LLVMContext::MD_nosanitize)) {
-    return true;
-  }
-
-  // Skip SPIR-V built-in varibles
-  auto *OrigValue = Addr->stripInBoundsOffsets();
-  assert(OrigValue != nullptr);
-  if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
-    return true;
-
-  // Ignore load/store for target ext type since we can't know exactly what size
-  // it is.
-  if (auto *SI = dyn_cast<StoreInst>(I))
-    if (getTargetExtType(SI->getValueOperand()->getType()) ||
-        isJointMatrixAccess(SI->getPointerOperand()))
-      return true;
-
-  if (auto *LI = dyn_cast<LoadInst>(I))
-    if (getTargetExtType(I->getType()) ||
-        isJointMatrixAccess(LI->getPointerOperand()))
-      return true;
-
-  Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
-  switch (PtrTy->getPointerAddressSpace()) {
-  case kSpirOffloadPrivateAS:
-    return !InstrumentPrivatePtr;
-  case kSpirOffloadLocalAS:
-    return !InstrumentLocalPtr;
-  case kSpirOffloadGenericAS:
-    return false;
-  }
-
-  return false;
-}
-
 } // namespace SanitizerCommonUtils
 } // namespace llvm

llvm/test/Instrumentation/MemorySanitizer/SPIRV/ignore_target_ext_type.ll

@@ -3,42 +3,23 @@
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
 target triple = "spir64-unknown-unknown"
 
-%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1) }
+%"class.sycl::_V1::ext::oneapi::bfloat16" = type { i16 }
+%"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix" = type { target("spirv.CooperativeMatrixKHR", i16, 3, 16, 32, 0) }
 
-; Function Attrs: sanitize_address
-define spir_kernel void @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE() #0 {
-entry:
-; CHECK-LABEL-DAG: @_ZTS4multIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm16ELm32EE
-; CHECK-NOT: MyAlloc
-  %a = alloca [2 x %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix"], i32 0, align 8
-  br label %for.cond10.i
-
-for.cond10.i:                                     ; preds = %for.cond10.i, %entry
-  %0 = load target("spirv.JointMatrixINTEL", i16, 16, 32, 0, 3, 0, 1), ptr null, align 8
-  store target("spirv.JointMatrixINTEL", float, 16, 16, 3, 3, 2) zeroinitializer, ptr null, align 8
-; CHECK-NOT: call void @asan_load
-; CHECK-NOT: call void @asan_store
-  br label %for.cond10.i
-}
+; CHECK-LABEL: @test
+; CHECK-NOT: call i64 @__msan_get_shadow
+declare dso_local spir_func noundef ptr addrspace(4) @_Z19__spirv_AccessChainIN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm32ELN5__spv9MatrixUseE0ELNS5_5Scope4FlagE3EEPT_PPNS5_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr addrspace(4) noundef, i64 noundef)
 
-; Function Attrs: sanitize_address
-define spir_kernel void @AccessChain() #0 {
+define weak_odr dso_local spir_kernel void @test() {
 entry:
-; CHECK-LABEL-DAG: @AccessChain
-  %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8
-  %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0
-  %call.i35 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
-  %1 = getelementptr inbounds { i16 }, ptr %call.i35, i64 0, i32 0
-; CHECK-NOT: call void @__asan_load
-; CHECK-NOT: call void @__asan_store
-  %2 = load i16, ptr %1, align 4
-  %call.i42 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
-  %3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0
-  store i16 %2, ptr %3, align 4
+  %sub_a.i = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8
+  %element.i = alloca %"class.sycl::_V1::ext::oneapi::bfloat16", align 2
+  %0 = getelementptr inbounds { i16 }, ptr %element.i, i64 0, i32 0
+  %spvm.i = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %sub_a.i, i64 0, i32 0
+  %addrcast = addrspacecast ptr %spvm.i to ptr addrspace(4)
+  %call.i67 = call spir_func noundef ptr addrspace(4) @_Z19__spirv_AccessChainIN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm32ELN5__spv9MatrixUseE0ELNS5_5Scope4FlagE3EEPT_PPNS5_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr addrspace(4) noundef %addrcast, i64 1)
+  %gep = getelementptr inbounds nuw { i16 }, ptr addrspace(4) %call.i67, i64 0, i32 0
+  %val = load i16, ptr %0, align 2
+  store i16 %val, ptr addrspace(4) %gep, align 2
   ret void
 }
-
-declare spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr, i64)
-
-attributes #0 = { sanitize_address }
-