[SPIRV] Add tests to improve test coverage

Differential Revision: https://reviews.llvm.org/D132903

llvm/test/CodeGen/SPIRV/SampledImageRetType.ll

@@ -0,0 +1,29 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+%opencl.image1d_ro_t = type opaque
+; CHECK: %[[#image1d_t:]] = OpTypeImage
+%opencl.sampler_t = type opaque
+; CHECK: %[[#sampler_t:]] = OpTypeSampler
+; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage
+
+declare dso_local spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) local_unnamed_addr
+
+declare dso_local spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(i8 addrspace(4)*, i32, i32, float) local_unnamed_addr
+
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(2) constant <3 x i64>, align 32
+
+define weak_odr dso_local spir_kernel void @_ZTS17image_kernel_readILi1EE(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) {
+; CHECK: OpFunction
+; CHECK: %[[#image:]] = OpFunctionParameter %[[#image1d_t]]
+; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]]
+  %3 = load <3 x i64>, <3 x i64> addrspace(2)* @__spirv_BuiltInGlobalInvocationId, align 32
+  %4 = extractelement <3 x i64> %3, i64 0
+  %5 = trunc i64 %4 to i32
+  %6 = tail call spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)* %0, %opencl.sampler_t addrspace(2)* %1)
+  %7 = tail call spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(i8 addrspace(4)* %6, i32 %5, i32 2, float 0.000000e+00)
+
+; CHECK: %[[#sampled_image:]] = OpSampledImage %[[#sampled_image_t]] %[[#image]] %[[#sampler]]
+; CHECK: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#sampled_image]] %[[#]] {{.*}} %[[#]]
+
+  ret void
+}

llvm/test/CodeGen/SPIRV/event_no_group_cap.ll

@@ -0,0 +1,21 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; __kernel void test_fn( const __global char *src)
+; {
+;     wait_group_events(0, NULL);
+; }
+
+; CHECK-NOT: OpCapability Groups
+; CHECK: OpGroupWaitEvents
+
+%opencl.event_t = type opaque
+
+define dso_local spir_kernel void @test_fn(i8 addrspace(1)* noundef %src) {
+entry:
+  %src.addr = alloca i8 addrspace(1)*, align 8
+  store i8 addrspace(1)* %src, i8 addrspace(1)** %src.addr, align 8
+  call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 noundef 0, %opencl.event_t* addrspace(4)* noundef null)
+  ret void
+}
+
+declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 noundef, %opencl.event_t* addrspace(4)* noundef)

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

@@ -0,0 +1,89 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK:     %[[#TypeImage:]] = OpTypeImage
+; CHECK:     %[[#TypeSampler:]] = OpTypeSampler
+; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]]
+; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]]
+
+; CHECK:     %[[#srcimg:]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK:     %[[#sampler:]] = OpFunctionParameter %[[#TypeSampler]]
+
+; CHECK:     %[[#srcimg_addr:]] = OpVariable %[[#TypeImagePtr]]
+; CHECK:     %[[#sampler_addr:]] = OpVariable %[[#TypeSamplerPtr]]
+
+; CHECK:     OpStore %[[#srcimg_addr]] %[[#srcimg]]
+; CHECK:     OpStore %[[#sampler_addr]] %[[#sampler]]
+
+; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
+; CHECK:     %[[#sampler_val:]] = OpLoad %[[#]] %[[#sampler_addr]]
+
+; CHECK:     %[[#]] = OpSampledImage %[[#]] %[[#srcimg_val]] %[[#sampler_val]]
+; CHECK-NEXT: OpImageSampleExplicitLod
+
+; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
+; CHECK:     %[[#]] = OpImageQuerySizeLod %[[#]] %[[#srcimg_val]]
+
+;; Excerpt from opencl-c-base.h
+;; typedef float float4 __attribute__((ext_vector_type(4)));
+;; typedef int int2 __attribute__((ext_vector_type(2)));
+;; typedef __SIZE_TYPE__ size_t;
+;;
+;; Excerpt from opencl-c.h to speed up compilation.
+;; #define __ovld __attribute__((overloadable))
+;; #define __purefn __attribute__((pure))
+;; #define __cnfn __attribute__((const))
+;; size_t __ovld __cnfn get_global_id(unsigned int dimindx);
+;; int __ovld __cnfn get_image_width(read_only image2d_t image);
+;; float4 __purefn __ovld read_imagef(read_only image2d_t image, sampler_t sampler, int2 coord);
+;;
+;;
+;; __kernel void test_fn(image2d_t srcimg, sampler_t sampler, global float4 *results) {
+;;   int tid_x = get_global_id(0);
+;;   int tid_y = get_global_id(1);
+;;   results[tid_x + tid_y * get_image_width(srcimg)] = read_imagef(srcimg, sampler, (int2){tid_x, tid_y});
+;; }
+
+%opencl.image2d_ro_t = type opaque
+%opencl.sampler_t = type opaque
+
+define dso_local spir_kernel void @test_fn(%opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.sampler_t addrspace(2)* %sampler, <4 x float> addrspace(1)* noundef %results) {
+entry:
+  %srcimg.addr = alloca %opencl.image2d_ro_t addrspace(1)*, align 4
+  %sampler.addr = alloca %opencl.sampler_t addrspace(2)*, align 4
+  %results.addr = alloca <4 x float> addrspace(1)*, align 4
+  %tid_x = alloca i32, align 4
+  %tid_y = alloca i32, align 4
+  %.compoundliteral = alloca <2 x i32>, align 8
+  store %opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  store %opencl.sampler_t addrspace(2)* %sampler, %opencl.sampler_t addrspace(2)** %sampler.addr, align 4
+  store <4 x float> addrspace(1)* %results, <4 x float> addrspace(1)** %results.addr, align 4
+  %call = call spir_func i32 @_Z13get_global_idj(i32 noundef 0)
+  store i32 %call, i32* %tid_x, align 4
+  %call1 = call spir_func i32 @_Z13get_global_idj(i32 noundef 1)
+  store i32 %call1, i32* %tid_y, align 4
+  %0 = load %opencl.image2d_ro_t addrspace(1)*, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  %1 = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** %sampler.addr, align 4
+  %2 = load i32, i32* %tid_x, align 4
+  %vecinit = insertelement <2 x i32> undef, i32 %2, i32 0
+  %3 = load i32, i32* %tid_y, align 4
+  %vecinit2 = insertelement <2 x i32> %vecinit, i32 %3, i32 1
+  store <2 x i32> %vecinit2, <2 x i32>* %.compoundliteral, align 8
+  %4 = load <2 x i32>, <2 x i32>* %.compoundliteral, align 8
+  %call3 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(%opencl.image2d_ro_t addrspace(1)* %0, %opencl.sampler_t addrspace(2)* %1, <2 x i32> noundef %4)
+  %5 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %results.addr, align 4
+  %6 = load i32, i32* %tid_x, align 4
+  %7 = load i32, i32* %tid_y, align 4
+  %8 = load %opencl.image2d_ro_t addrspace(1)*, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  %call4 = call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %8)
+  %mul = mul nsw i32 %7, %call4
+  %add = add nsw i32 %6, %mul
+  %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %5, i32 %add
+  store <4 x float> %call3, <4 x float> addrspace(1)* %arrayidx, align 16
+  ret void
+}
+
+declare spir_func i32 @_Z13get_global_idj(i32 noundef)
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <2 x i32> noundef)
+
+declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)

llvm/test/CodeGen/SPIRV/image.ll

@@ -0,0 +1,31 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:     %[[#VOID_TY:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV-NOT: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV:     OpImageSampleExplicitLod
+; CHECK-SPIRV:     OpImageWrite
+
+%opencl.image2d_t = type opaque
+
+define spir_kernel void @image_copy(%opencl.image2d_t addrspace(1)* readnone %image1, %opencl.image2d_t addrspace(1)* %image2) !kernel_arg_access_qual !1 {
+entry:
+  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %conv = trunc i64 %call to i32
+  %call1 = tail call spir_func i64 @_Z13get_global_idj(i32 1)
+  %conv2 = trunc i64 %call1 to i32
+  %vecinit = insertelement <2 x i32> undef, i32 %conv, i32 0
+  %vecinit3 = insertelement <2 x i32> %vecinit, i32 %conv2, i32 1
+  %call4 = tail call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)* %image1, i32 20, <2 x i32> %vecinit3)
+  tail call spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(%opencl.image2d_t addrspace(1)* %image2, <2 x i32> %vecinit3, <4 x float> %call4)
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
+
+declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)*, i32, <2 x i32>)
+
+declare spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(%opencl.image2d_t addrspace(1)*, <2 x i32>, <4 x float>)
+
+!1 = !{!"read_only", !"write_only"}

llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll

@@ -0,0 +1,28 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:     %[[#TypeImage:]] = OpTypeImage
+; CHECK-SPIRV-NOT: OpTypeImage
+; CHECK-SPIRV:     %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#ParamID:]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionCall %[[#]] %[[#]] %[[#ParamID]]
+
+%opencl.image2d_ro_t = type opaque
+
+define spir_func void @f0(%opencl.image2d_ro_t addrspace(1)* %v2, <2 x float> %v3) {
+entry:
+  ret void
+}
+
+define spir_func void @f1(%opencl.image2d_ro_t addrspace(1)* %v2, <2 x float> %v3) {
+entry:
+  ret void
+}
+
+define spir_kernel void @test(%opencl.image2d_ro_t addrspace(1)* %v1) {
+entry:
+  call spir_func void @f0(%opencl.image2d_ro_t addrspace(1)* %v1, <2 x float> <float 1.000000e+00, float 5.000000e+00>)
+  ret void
+}

llvm/test/CodeGen/SPIRV/image_dim.ll

@@ -0,0 +1,12 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability Sampled1D
+; CHECK-SPIRV-DAG: OpCapability SampledBuffer
+
+%opencl.image1d_t = type opaque
+%opencl.image1d_buffer_t = type opaque
+
+define spir_kernel void @image_d(%opencl.image1d_t addrspace(1)* %image1d_td6, %opencl.image1d_buffer_t addrspace(1)* %image1d_buffer_td8) {
+entry:
+  ret void
+}

llvm/test/CodeGen/SPIRV/image_store.ll

@@ -0,0 +1,23 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; Image types may be represented in two ways while translating to SPIR-V:
+;; - OpenCL form, for example, '%opencl.image2d_ro_t',
+;; - SPIR-V form, for example, '%spirv.Image._void_1_0_0_0_0_0_0',
+;; but it is still one type which should be translated to one SPIR-V type.
+;;
+;; The test checks that the code below is successfully translated and only one
+;; SPIR-V type for images is generated.
+
+; CHECK:     OpTypeImage
+; CHECK-NOT: OpTypeImage
+
+%opencl.image2d_ro_t = type opaque
+%spirv.Image._void_1_0_0_0_0_0_0 = type opaque
+
+define spir_kernel void @read_image(%opencl.image2d_ro_t addrspace(1)* %srcimg) {
+entry:
+  %srcimg.addr = alloca %opencl.image2d_ro_t addrspace(1)*, align 8
+  %spirvimg.addr = alloca %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, align 8
+  store %opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 8
+  ret void
+}

llvm/test/CodeGen/SPIRV/instructions/float-fast-flags.ll

@@ -0,0 +1,95 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; DISABLED-CHECK-DAG: OpName [[FNEG:%.+]] "scalar_fneg"
+; CHECK-DAG: OpName [[FADD:%.+]] "test_fadd"
+; CHECK-DAG: OpName [[FSUB:%.+]] "test_fsub"
+; CHECK-DAG: OpName [[FMUL:%.+]] "test_fmul"
+; CHECK-DAG: OpName [[FDIV:%.+]] "test_fdiv"
+; CHECK-DAG: OpName [[FREM:%.+]] "test_frem"
+; CHECK-DAG: OpName [[FMA:%.+]] "test_fma"
+
+; CHECK-DAG: [[F32Ty:%.+]] = OpTypeFloat 32
+; CHECK-DAG: [[FNTy:%.+]] = OpTypeFunction [[F32Ty]] [[F32Ty]] [[F32Ty]]
+
+
+; CHECK:      [[FADD]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFAdd [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fadd(float %a, float %b) {
+    %c = fadd nnan ninf float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FSUB]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFSub [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fsub(float %a, float %b) {
+    %c = fsub fast float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FMUL]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFMul [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks]
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fmul(float %a, float %b) {
+    %c = fmul contract float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FDIV]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFDiv [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fdiv(float %a, float %b) {
+    %c = fdiv arcp nsz float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FREM]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFRem [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_frem(float %a, float %b) {
+    %c = frem nsz float %a, %b
+    ret float %c
+}
+
+
+declare float @llvm.fma.f32(float, float, float)
+
+; CHECK:      [[FMA]] = OpFunction
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[C:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[R:%.+]] = OpExtInst [[F32Ty]] {{%.+}} fma [[A]] [[B]] [[C]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fma(float %a, float %b, float %c) {
+    %r = call float @llvm.fma.f32(float %a, float %b, float %c)
+    ret float %r
+}