llvm/test/CodeGen/SPIRV/transcoding/SpecConstantComposite.ll

   1 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
   2
   3 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC3:]] SpecId 3
   4 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC4:]] SpecId 4
   5 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC6:]] SpecId 6
   6 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC7:]] SpecId 7
   7 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC10:]] SpecId 10
   8 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC11:]] SpecId 11
   9
  10 ; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 32
  11 ; CHECK-SPIRV-DAG: %[[#Float:]] = OpTypeFloat 32
  12 ; CHECK-SPIRV-DAG: %[[#StructA:]] = OpTypeStruct %[[#Int]] %[[#Float]]
  13 ; CHECK-SPIRV-DAG: %[[#Array:]] = OpTypeArray %[[#StructA]] %[[#]]
  14 ; CHECK-SPIRV-DAG: %[[#Vector:]] = OpTypeVector %[[#Int]] 2
  15 ; 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 any
  26
  27 define 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 {
  28 entry:
  29   %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8
  30   %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
  31   %1 = load i64, i64* %0, align 8
  32   %add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1
  33   %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]] 1
  39
  40   %5 = call float @_Z20__spirv_SpecConstantif(i32 4, float 0.000000e+00)
  41 ; CHECK-SPIRV-DAG: %[[#SC4]] = OpSpecConstant %[[#Float]] 0
  42
  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]] 35
  48
  49   %8 = call float @_Z20__spirv_SpecConstantif(i32 7, float 0.000000e+00)
  50 ; CHECK-SPIRV-DAG: %[[#SC7]] = OpSpecConstant %[[#Float]] 0
  51
  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]] 45
  60
  61   %12 = call i32 @_Z20__spirv_SpecConstantii(i32 11, i32 55)
  62 ; CHECK-SPIRV-DAG: %[[#SC11]] = OpSpecConstant %[[#Int]] 55
  63
  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 8
  74 ; 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 void
  81 }
  82
  83 declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture)
  84
  85 declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture)
  86
  87 declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg)
  88
  89 declare i32 @_Z20__spirv_SpecConstantii(i32, i32)
  90
  91 declare float @_Z20__spirv_SpecConstantif(i32, float)
  92
  93 declare %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32, float)
  94
  95 declare [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A, %struct._ZTS1A.A)
  96
  97 declare <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32, i32)
  98
  99 declare %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32>)
 100
 101 declare %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec")