llvm/test/CodeGen/SPIRV/image-unoptimized.ll

   1 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
   2
   3 ; CHECK:     %[[#TypeImage:]] = OpTypeImage
   4 ; CHECK:     %[[#TypeSampler:]] = OpTypeSampler
   5 ; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]]
   6 ; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]]
   7
   8 ; CHECK:     %[[#srcimg:]] = OpFunctionParameter %[[#TypeImage]]
   9 ; CHECK:     %[[#sampler:]] = OpFunctionParameter %[[#TypeSampler]]
  10
  11 ; CHECK:     %[[#srcimg_addr:]] = OpVariable %[[#TypeImagePtr]]
  12 ; CHECK:     %[[#sampler_addr:]] = OpVariable %[[#TypeSamplerPtr]]
  13
  14 ; CHECK:     OpStore %[[#srcimg_addr]] %[[#srcimg]]
  15 ; CHECK:     OpStore %[[#sampler_addr]] %[[#sampler]]
  16
  17 ; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
  18 ; CHECK:     %[[#sampler_val:]] = OpLoad %[[#]] %[[#sampler_addr]]
  19
  20 ; CHECK:     %[[#]] = OpSampledImage %[[#]] %[[#srcimg_val]] %[[#sampler_val]]
  21 ; CHECK-NEXT: OpImageSampleExplicitLod
  22
  23 ; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
  24 ; CHECK:     %[[#]] = OpImageQuerySizeLod %[[#]] %[[#srcimg_val]]
  25
  26 ;; Excerpt from opencl-c-base.h
  27 ;; typedef float float4 __attribute__((ext_vector_type(4)));
  28 ;; typedef int int2 __attribute__((ext_vector_type(2)));
  29 ;; typedef __SIZE_TYPE__ size_t;
  30 ;;
  31 ;; Excerpt from opencl-c.h to speed up compilation.
  32 ;; #define __ovld __attribute__((overloadable))
  33 ;; #define __purefn __attribute__((pure))
  34 ;; #define __cnfn __attribute__((const))
  35 ;; size_t __ovld __cnfn get_global_id(unsigned int dimindx);
  36 ;; int __ovld __cnfn get_image_width(read_only image2d_t image);
  37 ;; float4 __purefn __ovld read_imagef(read_only image2d_t image, sampler_t sampler, int2 coord);
  38 ;;
  39 ;;
  40 ;; __kernel void test_fn(image2d_t srcimg, sampler_t sampler, global float4 *results) {
  41 ;;   int tid_x = get_global_id(0);
  42 ;;   int tid_y = get_global_id(1);
  43 ;;   results[tid_x + tid_y * get_image_width(srcimg)] = read_imagef(srcimg, sampler, (int2){tid_x, tid_y});
  44 ;; }
  45
  46 define dso_local spir_kernel void @test_fn(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcimg, target("spirv.Sampler") %sampler, <4 x float> addrspace(1)* noundef %results) {
  47 entry:
  48   %srcimg.addr = alloca target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), align 4
  49   %sampler.addr = alloca target("spirv.Sampler"), align 4
  50   %results.addr = alloca <4 x float> addrspace(1)*, align 4
  51   %tid_x = alloca i32, align 4
  52   %tid_y = alloca i32, align 4
  53   %.compoundliteral = alloca <2 x i32>, align 8
  54   store target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcimg, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
  55   store target("spirv.Sampler") %sampler, target("spirv.Sampler")* %sampler.addr, align 4
  56   store <4 x float> addrspace(1)* %results, <4 x float> addrspace(1)** %results.addr, align 4
  57   %call = call spir_func i32 @_Z13get_global_idj(i32 noundef 0)
  58   store i32 %call, i32* %tid_x, align 4
  59   %call1 = call spir_func i32 @_Z13get_global_idj(i32 noundef 1)
  60   store i32 %call1, i32* %tid_y, align 4
  61   %0 = load target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
  62   %1 = load target("spirv.Sampler"), target("spirv.Sampler")* %sampler.addr, align 4
  63   %2 = load i32, i32* %tid_x, align 4
  64   %vecinit = insertelement <2 x i32> undef, i32 %2, i32 0
  65   %3 = load i32, i32* %tid_y, align 4
  66   %vecinit2 = insertelement <2 x i32> %vecinit, i32 %3, i32 1
  67   store <2 x i32> %vecinit2, <2 x i32>* %.compoundliteral, align 8
  68   %4 = load <2 x i32>, <2 x i32>* %.compoundliteral, align 8
  69   %call3 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1, <2 x i32> noundef %4)
  70   %5 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %results.addr, align 4
  71   %6 = load i32, i32* %tid_x, align 4
  72   %7 = load i32, i32* %tid_y, align 4
  73   %8 = load target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
  74   %call4 = call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %8)
  75   %mul = mul nsw i32 %7, %call4
  76   %add = add nsw i32 %6, %mul
  77   %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %5, i32 %add
  78   store <4 x float> %call3, <4 x float> addrspace(1)* %arrayidx, align 16
  79   ret void
  80 }
  81
  82 declare spir_func i32 @_Z13get_global_idj(i32 noundef)
  83
  84 declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <2 x i32> noundef)
  85
  86 declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))