1; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s 2 3; TODO: This test currently fails with LLVM_ENABLE_EXPENSIVE_CHECKS enabled 4; XFAIL: expensive_checks 5 6; CHECK: OpDecorate %[[#SpecConst:]] SpecId 0 7; CHECK: %[[#SpecConst]] = OpSpecConstant %[[#]] 70 8; CHECK: %[[#]] = OpPhi %[[#]] %[[#]] %[[#]] %[[#SpecConst]] %[[#]] 9 10%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } 11%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } 12 13$_ZTS6kernel = comdat any 14 15define weak_odr dso_local spir_kernel void @_ZTS6kernel(i8 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_3) local_unnamed_addr comdat { 16entry: 17 %0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %_arg_3, i64 0, i32 0, i32 0, i64 0 18 %1 = addrspacecast i64* %0 to i64 addrspace(4)* 19 %2 = load i64, i64 addrspace(4)* %1, align 8 20 br label %for.cond.i.i 21 22for.cond.i.i: ; preds = %for.body.i.i, %entry 23 %value.0.i.i = phi i8 [ -1, %entry ], [ %3, %for.body.i.i ] 24 %cmp.i.i = phi i1 [ true, %entry ], [ false, %for.body.i.i ] 25 br i1 %cmp.i.i, label %for.body.i.i, label %_ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlNS0_14kernel_handlerEE_clES4_.exit 26 27for.body.i.i: ; preds = %for.cond.i.i 28 %3 = call i8 @_Z20__spirv_SpecConstantia(i32 0, i8 70) 29 br label %for.cond.i.i 30 31_ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlNS0_14kernel_handlerEE_clES4_.exit: ; preds = %for.cond.i.i 32 %add.ptr.i = getelementptr inbounds i8, i8 addrspace(1)* %_arg_, i64 %2 33 %arrayidx.ascast.i.i = addrspacecast i8 addrspace(1)* %add.ptr.i to i8 addrspace(4)* 34 store i8 %value.0.i.i, i8 addrspace(4)* %arrayidx.ascast.i.i, align 1 35 ret void 36} 37 38declare i8 @_Z20__spirv_SpecConstantia(i32, i8) 39