1; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s 2 3; CHECK-DAG: %[[#TypeImage:]] = OpTypeImage 4; CHECK-DAG: %[[#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 46define 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) { 47entry: 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 82declare spir_func i32 @_Z13get_global_idj(i32 noundef) 83 84declare 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 86declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)) 87