|
1 | 1 | ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
|
2 | 2 | ; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s
|
3 | 3 |
|
| 4 | +; This test checks that undef initializer of a global variable is preserved |
| 5 | +; during ESIMDLowerVecArg transformation |
| 6 | + |
4 | 7 | 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"
|
5 | 8 | target triple = "spir64-unknown-unknown-sycldevice"
|
6 | 9 |
|
7 |
| -%"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> } |
| 10 | +%"class.cl::sycl::INTEL::gpu::simd" = type { <2512 x i32> } |
8 | 11 |
|
9 |
| -; CHECK: @GlobalGRF_data = dso_local global <2512 x i32> undef, align 16384 |
10 |
| -@GlobalGRF_data = dso_local global %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" undef, align 16384 |
| 12 | +; CHECK: @Global = dso_local global <2512 x i32> undef, align 16384 |
| 13 | +@Global = dso_local global %"class.cl::sycl::INTEL::gpu::simd" undef, align 16384 |
11 | 14 |
|
12 | 15 | define void @f(<2512 x i32> %simd_val) {
|
13 | 16 | ; CHECK-LABEL: @f(
|
14 |
| -; 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 |
| 17 | +; 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 |
15 | 18 | ; CHECK-NEXT: ret void
|
16 | 19 | ;
|
17 |
| - 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 |
| 20 | + 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 |
18 | 21 | ret void
|
19 | 22 | }
|
0 commit comments