From af13cc6e196ab453ac213ea50b538d9fc596fe02 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Mon, 28 Sep 2020 11:18:13 -0700 Subject: [PATCH 1/3] [SYCL][ESIMD] Precommit test for ESIMDLowerVecArg pass --- llvm/test/SYCLLowerIR/esimd_global_undef.ll | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) create mode 100644 llvm/test/SYCLLowerIR/esimd_global_undef.ll diff --git a/llvm/test/SYCLLowerIR/esimd_global_undef.ll b/llvm/test/SYCLLowerIR/esimd_global_undef.ll new file mode 100644 index 0000000000000..f462dbaeb409c --- /dev/null +++ b/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 +} From b63292d151a06992291bc3f2350db40ad8a8b684 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Mon, 28 Sep 2020 11:21:00 -0700 Subject: [PATCH 2/3] [SYCL][ESIMD] Preserving undef intializer for globals in ESIMDLowerVecArg pass --- llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp | 9 ++++++--- llvm/test/SYCLLowerIR/esimd_global_undef.ll | 2 +- 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp index 3488d9a9532d9..cb07af7ae4a8a 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp +++ b/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(G.getInitializer()) + ? static_cast(UndefValue::get(NewTy)) + : static_cast(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); diff --git a/llvm/test/SYCLLowerIR/esimd_global_undef.ll b/llvm/test/SYCLLowerIR/esimd_global_undef.ll index f462dbaeb409c..b162e7be968a7 100644 --- a/llvm/test/SYCLLowerIR/esimd_global_undef.ll +++ b/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) { From ee814a946c6c5be27ef22f917271dd79f4535f2e Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Tue, 29 Sep 2020 09:51:58 -0700 Subject: [PATCH 3/3] Addressed comments --- llvm/test/SYCLLowerIR/esimd_global_undef.ll | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/llvm/test/SYCLLowerIR/esimd_global_undef.ll b/llvm/test/SYCLLowerIR/esimd_global_undef.ll index b162e7be968a7..8cc6a46cae947 100644 --- a/llvm/test/SYCLLowerIR/esimd_global_undef.ll +++ b/llvm/test/SYCLLowerIR/esimd_global_undef.ll @@ -1,19 +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._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 @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: 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._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 + 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 }