From 1cd3ad78b54c638f2783a9c9f22ed6dd700ba302 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Tue, 10 Dec 2024 09:11:04 -0800 Subject: [PATCH] [SYCL][DeviceASAN] Fix AcceeChain to a matrix for bfloat16 If element type is bfloat16 its storage will be accessed via GEP. So need to skip this GEP during AccessChain handling. Signed-off-by: Sidorov, Dmitry --- llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp | 3 ++- .../AddressSanitizer/SPIRV/ignore_target_ext_type.ll | 6 ++++-- 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 955bcc2d3ff27..9da7458710ca0 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -1623,7 +1623,8 @@ static TargetExtType *getTargetExtType(Type *Ty) { // store float %1, ptr %call, align 4 // clang-format on static bool isJointMatrixAccess(Value *V) { - if (auto *CI = dyn_cast(V)) { + auto *ActualV = V->stripInBoundsOffsets(); + if (auto *CI = dyn_cast(ActualV)) { for (Value *Op : CI->args()) { if (auto *AI = dyn_cast(Op->stripInBoundsOffsets())) if (auto *TargetTy = getTargetExtType(AI->getAllocatedType())) diff --git a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/ignore_target_ext_type.ll b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/ignore_target_ext_type.ll index 82b2db0d74cc7..d8954dd0f49a9 100644 --- a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/ignore_target_ext_type.ll +++ b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/ignore_target_ext_type.ll @@ -26,11 +26,13 @@ entry: %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 - %1 = load float, ptr %call.i35, align 4 + %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) - store float %1, ptr %call.i42, align 4 + %3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0 + store i16 %2, ptr %3, align 4 ret void }