[SYCL][ESIMD] Preserving undef intializer for globals in ESIMDLowerVecArg pass

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

@@ -6,7 +6,7 @@ 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> zeroinitializer, align 16384
+; 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 @f(<2512 x i32> %simd_val) {