[SYCL][ESIMD] Precommit test for ESIMDLowerVecArg pass

Author: DenisBakhvalov

llvm/test/SYCLLowerIR/esimd_global_undef.ll

@@ -0,0 +1,19 @@
+; 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> zeroinitializer, 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) {
+; CHECK-LABEL: @f(
+; CHECK-NEXT:    store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* getelementptr (%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* 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)*), i64 0, i32 0), align 16384
+; CHECK-NEXT:    ret void
+;
+  store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* getelementptr (%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* addrspacecast (%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)*), i64 0, i32 0), align 16384
+  ret void
+}