llvm/test/CodeGen/SPIRV/image/sampler.ll

   1 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
   2 ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
   3
   4 ; CHECK-DAG: %[[#i32:]] = OpTypeInt 32 0
   5 ; CHECK-DAG: %[[#v4i32:]] = OpTypeVector %[[#i32]] 4
   6 ; CHECK-DAG: %[[#ptrv4i32:]] = OpTypePointer CrossWorkgroup %[[#v4i32]]
   7 ; CHECK-DAG: %[[#float:]] = OpTypeFloat 32
   8 ; CHECK-DAG: %[[#typesampled:]] = OpTypeSampledImage
   9 ; CHECK-DAG: %[[#const0:]] = OpConstant %[[#float]] 0
  10 ; CHECK: OpFunction
  11 ; CHECK: OpFunctionParameter
  12 ; CHECK: %[[#arg1:]] = OpFunctionParameter
  13 ; CHECK: %[[#arg2:]] = OpFunctionParameter
  14 ; CHECK: %[[#addr:]] = OpInBoundsPtrAccessChain
  15 ; CHECK: %[[#img:]] = OpSampledImage %[[#typesampled:]] %[[#arg1]] %[[#arg2]]
  16 ; CHECK: %[[#sample:]] = OpImageSampleExplicitLod %[[#v4i32]] %[[#img]] %[[#const0]] Lod %[[#const0]]
  17 ; CHECK: %[[#casted:]] = OpBitcast %[[#ptrv4i32]] %[[#addr]]
  18 ; CHECK: OpStore %[[#casted]] %[[#sample]] Aligned 16
  19
  20 %"class.sycl::_V1::vec" = type { <4 x i32> }
  21
  22 define weak_odr dso_local spir_kernel void @foo(ptr addrspace(1) align 16 %_arg_acc, target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %_arg_img, target("spirv.Sampler") %_arg_sampler) {
  23 entry:
  24   %data = getelementptr inbounds %"class.sycl::_V1::vec", ptr addrspace(1) %_arg_acc, i64 0
  25   %img = tail call spir_func target("spirv.SampledImage", void, 0, 0, 0, 0, 0, 0, 0) @_Z20__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %_arg_img, target("spirv.Sampler") %_arg_sampler)
  26   %sample = tail call spir_func <4 x i32> @_Z30__spirv_ImageSampleExplicitLod(target("spirv.SampledImage", void, 0, 0, 0, 0, 0, 0, 0) %img, float 0.000000e+00, i32 2, float 0.000000e+00)
  27   store <4 x i32> %sample, ptr addrspace(1) %data, align 16
  28   ret void
  29 }
  30
  31 declare dso_local spir_func target("spirv.SampledImage", void, 0, 0, 0, 0, 0, 0, 0) @_Z20__spirv_SampledImage(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"))
  32 declare dso_local spir_func <4 x i32> @_Z30__spirv_ImageSampleExplicitLod(target("spirv.SampledImage", void, 0, 0, 0, 0, 0, 0, 0), float, i32, float)