102 lines · plain
1; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV2 3; CHECK-SPIRV-DAG: OpDecorate %[[#SC3:]] SpecId 34; CHECK-SPIRV-DAG: OpDecorate %[[#SC4:]] SpecId 45; CHECK-SPIRV-DAG: OpDecorate %[[#SC6:]] SpecId 66; CHECK-SPIRV-DAG: OpDecorate %[[#SC7:]] SpecId 77; CHECK-SPIRV-DAG: OpDecorate %[[#SC10:]] SpecId 108; CHECK-SPIRV-DAG: OpDecorate %[[#SC11:]] SpecId 119 10; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 3211; CHECK-SPIRV-DAG: %[[#Float:]] = OpTypeFloat 3212; CHECK-SPIRV-DAG: %[[#StructA:]] = OpTypeStruct %[[#Int]] %[[#Float]]13; CHECK-SPIRV-DAG: %[[#Array:]] = OpTypeArray %[[#StructA]] %[[#]]14; CHECK-SPIRV-DAG: %[[#Vector:]] = OpTypeVector %[[#Int]] 215; CHECK-SPIRV-DAG: %[[#Struct:]] = OpTypeStruct %[[#Vector]]16; CHECK-SPIRV-DAG: %[[#POD_TYPE:]] = OpTypeStruct %[[#Array]] %[[#Struct]]17 18%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" }19%struct._ZTS1A.A = type { i32, float }20%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> }21%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }22%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }23%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }24 25$_ZTS4Test = comdat any26 27define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD 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._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr comdat {28entry:29 %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 830 %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 031 %1 = load i64, i64* %0, align 832 %add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %133 %2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8*34 call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2)35 %3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)*36 37 %4 = call i32 @_Z20__spirv_SpecConstantii(i32 3, i32 1)38; CHECK-SPIRV-DAG: %[[#SC3]] = OpSpecConstant %[[#Int]] 139 40 %5 = call float @_Z20__spirv_SpecConstantif(i32 4, float 0.000000e+00)41; CHECK-SPIRV-DAG: %[[#SC4]] = OpSpecConstant %[[#Float]] 042 43 %6 = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %4, float %5)44; CHECK-SPIRV-DAG: %[[#SC_StructA0:]] = OpSpecConstantComposite %[[#StructA]] %[[#SC3]] %[[#SC4]]45 46 %7 = call i32 @_Z20__spirv_SpecConstantii(i32 6, i32 35)47; CHECK-SPIRV-DAG: %[[#SC6]] = OpSpecConstant %[[#Int]] 3548 49 %8 = call float @_Z20__spirv_SpecConstantif(i32 7, float 0.000000e+00)50; CHECK-SPIRV-DAG: %[[#SC7]] = OpSpecConstant %[[#Float]] 051 52 %9 = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %7, float %8)53; CHECK-SPIRV-DAG: %[[#SC_StructA1:]] = OpSpecConstantComposite %[[#StructA]] %[[#SC6]] %[[#SC7]]54 55 %10 = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %6, %struct._ZTS1A.A %9)56; CHECK-SPIRV-DAG: %[[#SC_Array:]] = OpSpecConstantComposite %[[#Array]] %[[#SC_StructA0]] %[[#SC_StructA1]]57 58 %11 = call i32 @_Z20__spirv_SpecConstantii(i32 10, i32 45)59; CHECK-SPIRV-DAG: %[[#SC10]] = OpSpecConstant %[[#Int]] 4560 61 %12 = call i32 @_Z20__spirv_SpecConstantii(i32 11, i32 55)62; CHECK-SPIRV-DAG: %[[#SC11]] = OpSpecConstant %[[#Int]] 5563 64 %13 = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 %11, i32 %12)65; CHECK-SPIRV-DAG: %[[#SC_Vector:]] = OpSpecConstantComposite %[[#Vector]] %[[#SC10]] %[[#SC11]]66 67 %14 = call %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32> %13)68; CHECK-SPIRV-DAG: %[[#SC_Struct:]] = OpSpecConstantComposite %[[#Struct]] %[[#SC_Vector]]69 70 %15 = call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A] %10, %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" %14)71; CHECK-SPIRV-DAG: %[[#SC_POD:]] = OpSpecConstantComposite %[[#POD_TYPE]] %[[#SC_Array]] %[[#SC_Struct]]72 73 store %struct._ZTS3POD.POD %15, %struct._ZTS3POD.POD addrspace(4)* %3, align 874; CHECK-SPIRV-DAG: OpStore %[[#]] %[[#SC_POD]]75 76 %16 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)*77 %17 = addrspacecast i8 addrspace(1)* %16 to i8 addrspace(4)*78 call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %17, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false)79 call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2)80 ret void81}82 83declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture)84 85declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture)86 87declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg)88 89declare i32 @_Z20__spirv_SpecConstantii(i32, i32)90 91declare float @_Z20__spirv_SpecConstantif(i32, float)92 93declare %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32, float)94 95declare [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A, %struct._ZTS1A.A)96 97declare <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32, i32)98 99declare %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32>)100 101declare %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec")102