[SYCL][ESIMD] Preserve undef initializer for globals in ESIMDLowerVecArg pass (#2555)

Author: DenisBakhvalov

llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp

@@ -272,10 +272,13 @@ void ESIMDLowerVecArgPass::fixGlobals(Module &M) {
     if (NewTy && !G.user_empty()) {
       // Peel off ptr type that getSimdArgPtrTyOrNull applies
       NewTy = NewTy->getPointerElementType();
-      auto ZeroInit = ConstantAggregateZero::get(NewTy);
+      auto InitVal =
+          G.hasInitializer() && isa<UndefValue>(G.getInitializer())
+              ? static_cast<ConstantData *>(UndefValue::get(NewTy))
+              : static_cast<ConstantData *>(ConstantAggregateZero::get(NewTy));
       auto NewGlobalVar =
-          new GlobalVariable(NewTy, G.isConstant(), G.getLinkage(), ZeroInit,
-                             "", G.getThreadLocalMode(), G.getAddressSpace());
+          new GlobalVariable(NewTy, G.isConstant(), G.getLinkage(), InitVal, "",
+                             G.getThreadLocalMode(), G.getAddressSpace());
       NewGlobalVar->setExternallyInitialized(G.isExternallyInitialized());
       NewGlobalVar->copyAttributesFrom(&G);
       NewGlobalVar->takeName(&G);

llvm/test/SYCLLowerIR/esimd_global_undef.ll

@@ -0,0 +1,22 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s
+
+; This test checks that undef initializer of a global variable is preserved
+; during ESIMDLowerVecArg transformation
+
+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 @f(<2512 x i32> %simd_val) {
+; CHECK-LABEL: @f(
+; CHECK-NEXT:    store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* getelementptr (%"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* 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)*), i64 0, i32 0), align 16384
+; CHECK-NEXT:    ret void
+;
+  store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* getelementptr (%"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* addrspacecast (%"class.cl::sycl::INTEL::gpu::simd"* @Global to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)*), i64 0, i32 0), align 16384
+  ret void
+}