[SYCL][ESIMD] Fixed compiler crash in LowerESIMDVecArg pass (#2556)

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`.

Author: DenisBakhvalov

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<GlobalVariable>(Old)) {
-    NewConstantExpr = cast<ConstantExpr>(
-        ConstantExpr::getBitCast(NewGlobalVar, OldGlobalType));
-    return NewConstantExpr;
-  }
-
-  auto InnerMost = createNewConstantExpr(
-      NewGlobalVar, OldGlobalType, cast<ConstantExpr>(Old)->getOperand(0));
-
-  NewConstantExpr = cast<ConstantExpr>(
-      cast<ConstantExpr>(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();
   }
 }
 

llvm/test/SYCLLowerIR/esimd_global_crash.ll

@@ -0,0 +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.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> }
+
+; 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.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.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
+}