|
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 there is no compiler crash when a Global |
| 5 | +; is used in simple instruction, not directly in ConstantExpr. |
| 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 @no_crash(<2512 x i32> %simd_val) {
|
13 | 16 | ; CHECK-LABEL: @no_crash(
|
14 |
| -; CHECK-NEXT: [[CAST:%.*]] = 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)* |
15 |
| -; CHECK-NEXT: [[GEP:%.*]] = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* [[CAST]], i64 0, i32 0 |
| 17 | +; CHECK-NEXT: [[CAST:%.*]] = 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)* |
| 18 | +; CHECK-NEXT: [[GEP:%.*]] = getelementptr %"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* [[CAST]], i64 0, i32 0 |
16 | 19 | ; CHECK-NEXT: store <2512 x i32> [[SIMD_VAL:%.*]], <2512 x i32> addrspace(4)* [[GEP]], align 16384
|
17 | 20 | ; CHECK-NEXT: ret void
|
18 | 21 | ;
|
19 |
| - %cast = addrspacecast %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd"* @GlobalGRF_data to %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* |
20 |
| - %gep = getelementptr %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd", %"class._ZTSN2cl4sycl5INTEL3gpu4simdIiLi2512EEE.cl::sycl::INTEL::gpu::simd" addrspace(4)* %cast, i64 0, i32 0 |
| 22 | + %cast = addrspacecast %"class.cl::sycl::INTEL::gpu::simd"* @Global to %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* |
| 23 | + %gep = getelementptr %"class.cl::sycl::INTEL::gpu::simd", %"class.cl::sycl::INTEL::gpu::simd" addrspace(4)* %cast, i64 0, i32 0 |
21 | 24 | store <2512 x i32> %simd_val, <2512 x i32> addrspace(4)* %gep, align 16384
|
22 | 25 | ret void
|
23 | 26 | }
|
0 commit comments