From 277cb686c6e1303f9f4e3f66108a75e171222351 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Mon, 28 Sep 2020 11:32:25 -0700 Subject: [PATCH 1/2] [SYCL][ESIMD] Fixed compiler crash in LowerESIMDVecArg pass This fixes potential compiler crash in LowerESIMDVecArg pass, which I encountered while writing a small test. Just to be clear, this doesn't happen in a real test, but potentially could happen. The problem arises when Global is used in simple instruction, not directly in ConstantExpr, e.g.: ``` @GlobalGRF_data = dso_local global %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" undef, align 16384 define void @no_crash(<2512 x i32> %simd_val) { %cast = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* ... } ``` It crashed in `ESIMDLowerVecArgPass::createNewConstantExpr`. --- llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp | 47 +++------------------ llvm/test/SYCLLowerIR/esimd_global_crash.ll | 23 ++++++++++ 2 files changed, 28 insertions(+), 42 deletions(-) create mode 100644 llvm/test/SYCLLowerIR/esimd_global_crash.ll diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp index cb07af7ae4a8a..e34e8d677d420 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp @@ -89,9 +89,6 @@ class ESIMDLowerVecArgPass { Function *rewriteFunc(Function &F); Type *getSimdArgPtrTyOrNull(Value *arg); void fixGlobals(Module &M); - void replaceConstExprWithGlobals(Module &M); - ConstantExpr *createNewConstantExpr(GlobalVariable *newGlobalVar, - Type *oldGlobalType, Value *old); void removeOldGlobals(); }; @@ -229,41 +226,6 @@ Function *ESIMDLowerVecArgPass::rewriteFunc(Function &F) { return NF; } -// Replace ConstantExpr if it contains old global variable. -ConstantExpr * -ESIMDLowerVecArgPass::createNewConstantExpr(GlobalVariable *NewGlobalVar, - Type *OldGlobalType, Value *Old) { - ConstantExpr *NewConstantExpr = nullptr; - - if (isa(Old)) { - NewConstantExpr = cast( - ConstantExpr::getBitCast(NewGlobalVar, OldGlobalType)); - return NewConstantExpr; - } - - auto InnerMost = createNewConstantExpr( - NewGlobalVar, OldGlobalType, cast(Old)->getOperand(0)); - - NewConstantExpr = cast( - cast(Old)->getWithOperandReplaced(0, InnerMost)); - - return NewConstantExpr; -} - -// Globals are part of ConstantExpr. This loop iterates over -// all such instances and replaces them with a new ConstantExpr -// consisting of new global vector* variable. -void ESIMDLowerVecArgPass::replaceConstExprWithGlobals(Module &M) { - for (auto &GlobalVars : OldNewGlobal) { - auto &G = *GlobalVars.first; - for (auto UseOfG : G.users()) { - auto NewGlobal = GlobalVars.second; - auto NewConstExpr = createNewConstantExpr(NewGlobal, G.getType(), UseOfG); - UseOfG->replaceAllUsesWith(NewConstExpr); - } - } -} - // This function creates new global variables of type vector* type // when old one is of simd* type. void ESIMDLowerVecArgPass::fixGlobals(Module &M) { @@ -288,16 +250,17 @@ void ESIMDLowerVecArgPass::fixGlobals(Module &M) { } } - replaceConstExprWithGlobals(M); - removeOldGlobals(); } // Remove old global variables from the program. void ESIMDLowerVecArgPass::removeOldGlobals() { for (auto &G : OldNewGlobal) { - G.first->removeDeadConstantUsers(); - G.first->eraseFromParent(); + auto OldGlob = G.first; + auto NewGlobal = G.second; + OldGlob->replaceAllUsesWith( + ConstantExpr::getBitCast(NewGlobal, OldGlob->getType())); + OldGlob->eraseFromParent(); } } diff --git a/llvm/test/SYCLLowerIR/esimd_global_crash.ll b/llvm/test/SYCLLowerIR/esimd_global_crash.ll new file mode 100644 index 0000000000000..0d97aad6add43 --- /dev/null +++ b/llvm/test/SYCLLowerIR/esimd_global_crash.ll @@ -0,0 +1,23 @@ +; NOTE: Assertions have been autogenerated by utils/update_test_checks.py +; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s + +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" +target triple = "spir64-unknown-unknown-sycldevice" + +%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> } + +; CHECK: @GlobalGRF_data = dso_local global <2512 x i32> undef, align 16384 +@GlobalGRF_data = dso_local global %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" undef, align 16384 + +define void @no_crash(<2512 x i32> %simd_val) { +; CHECK-LABEL: @no_crash( +; CHECK-NEXT: [[CAST:%.*]] = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* bitcast (<2512 x i32>* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"*) to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* +; CHECK-NEXT: [[GEP:%.*]] = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* [[CAST]], i64 0, i32 0 +; CHECK-NEXT: store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* [[GEP]], align 16384 +; CHECK-NEXT: ret void +; + %cast = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* + %gep = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* %cast, i64 0, i32 0 + store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* %gep, align 16384 + ret void +} From 143b2b8d0c5abec995577e4172a694890c9f000a Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Thu, 1 Oct 2020 10:21:49 -0700 Subject: [PATCH 2/2] Addressed code review comments --- llvm/test/SYCLLowerIR/esimd_global_crash.ll | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/llvm/test/SYCLLowerIR/esimd_global_crash.ll b/llvm/test/SYCLLowerIR/esimd_global_crash.ll index 0d97aad6add43..c25f1e4945282 100644 --- a/llvm/test/SYCLLowerIR/esimd_global_crash.ll +++ b/llvm/test/SYCLLowerIR/esimd_global_crash.ll @@ -1,23 +1,26 @@ ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py ; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s +; This test checks that there is no compiler crash when a Global +; is used in simple instruction, not directly in ConstantExpr. + 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" target triple = "spir64-unknown-unknown-sycldevice" -%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> } +%"class.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> } -; CHECK: @GlobalGRF_data = dso_local global <2512 x i32> undef, align 16384 -@GlobalGRF_data = dso_local global %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" undef, align 16384 +; CHECK: @Global = dso_local global <2512 x i32> undef, align 16384 +@Global = dso_local global %"class.cl::sycl::INTEL::gpu::simd" undef, align 16384 define void @no_crash(<2512 x i32> %simd_val) { ; CHECK-LABEL: @no_crash( -; CHECK-NEXT: [[CAST:%.*]] = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* bitcast (<2512 x i32>* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"*) to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* -; CHECK-NEXT: [[GEP:%.*]] = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* [[CAST]], i64 0, i32 0 +; CHECK-NEXT: [[CAST:%.*]] = addrspacecast %"class.cl::sycl::INTEL::gpu::simd"* bitcast (<2512 x i32>* @Global to %"class.cl::sycl::INTEL::gpu::simd"*) to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* +; CHECK-NEXT: [[GEP:%.*]] = getelementptr %"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* [[CAST]], i64 0, i32 0 ; CHECK-NEXT: store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* [[GEP]], align 16384 ; CHECK-NEXT: ret void ; - %cast = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* - %gep = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* %cast, i64 0, i32 0 + %cast = addrspacecast %"class.cl::sycl::INTEL::gpu::simd"* @Global to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* + %gep = getelementptr %"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* %cast, i64 0, i32 0 store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* %gep, align 16384 ret void }