[llvm] bb6a437 - [SPIRV] Add tests to improve test coverage
Andrey Tretyakov via llvm-commits
llvm-commits at lists.llvm.org
Sun Aug 28 12:21:33 PDT 2022
Author: Andrey Tretyakov
Date: 2022-08-28T22:18:11+03:00
New Revision: bb6a437306c4f20d069a84e4cc48e632cacb389f
URL: https://github.com/llvm/llvm-project/commit/bb6a437306c4f20d069a84e4cc48e632cacb389f
DIFF: https://github.com/llvm/llvm-project/commit/bb6a437306c4f20d069a84e4cc48e632cacb389f.diff
LOG: [SPIRV] Add tests to improve test coverage
Differential Revision: https://reviews.llvm.org/D132562
Added:
llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll
llvm/test/CodeGen/SPIRV/capability-integers.ll
llvm/test/CodeGen/SPIRV/capability-kernel.ll
llvm/test/CodeGen/SPIRV/empty-opencl32.ll
llvm/test/CodeGen/SPIRV/empty-opencl64.ll
llvm/test/CodeGen/SPIRV/empty.ll
llvm/test/CodeGen/SPIRV/half_extension.ll
llvm/test/CodeGen/SPIRV/half_no_extension.ll
llvm/test/CodeGen/SPIRV/linkage-types.ll
llvm/test/CodeGen/SPIRV/transcoding/NoSignedUnsignedWrap.ll
llvm/test/CodeGen/SPIRV/transcoding/fdiv.ll
llvm/test/CodeGen/SPIRV/transcoding/fmul.ll
llvm/test/CodeGen/SPIRV/transcoding/frem.ll
llvm/test/CodeGen/SPIRV/transcoding/fsub.ll
llvm/test/CodeGen/SPIRV/transcoding/vec8.ll
Modified:
llvm/test/CodeGen/SPIRV/constant/local-float-point-constants.ll
llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll
llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll
llvm/test/CodeGen/SPIRV/transcoding/ReqdSubgroupSize.ll
llvm/test/CodeGen/SPIRV/transcoding/fadd.ll
Removed:
################################################################################
diff --git a/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll b/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll
new file mode 100644
index 0000000000000..c61b3d90eb25a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll
@@ -0,0 +1,5 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+ at a = addrspace(2) constant i32 1, align 4
+
+; CHECK-DAG: OpCapability Kernel
diff --git a/llvm/test/CodeGen/SPIRV/capability-integers.ll b/llvm/test/CodeGen/SPIRV/capability-integers.ll
new file mode 100644
index 0000000000000..259199a548cf9
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/capability-integers.ll
@@ -0,0 +1,17 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpCapability Int8
+; CHECK-DAG: OpCapability Int16
+; CHECK-DAG: OpCapability Int64
+
+; CHECK-DAG: %[[#]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#]] = OpTypeInt 16 0
+; CHECK-DAG: %[[#]] = OpTypeInt 64 0
+
+ at a = addrspace(1) global i8 0, align 1
+ at b = addrspace(1) global i16 0, align 2
+ at c = addrspace(1) global i64 0, align 8
+
+define spir_kernel void @test_atomic_fn() {
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/capability-kernel.ll b/llvm/test/CodeGen/SPIRV/capability-kernel.ll
new file mode 100644
index 0000000000000..03ea58c985adb
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/capability-kernel.ll
@@ -0,0 +1,32 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpCapability Addresses
+
+; CHECK-DAG: OpCapability Linkage
+define spir_func void @func_export(i32 addrspace(1)* nocapture %a) {
+entry:
+; CHECK-DAG: OpCapability Int64
+ %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+ %cmp = icmp eq i64 %call, 0
+ br i1 %cmp, label %if.then, label %if.end
+
+if.then: ; preds = %entry
+ store i32 1, i32 addrspace(1)* %a, align 4
+ br label %if.end
+
+if.end: ; preds = %if.then, %entry
+ ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
+
+; CHECK-DAG: OpCapability Kernel
+; CHECK-NOT: OpCapability Shader
+; CHECK-NOT: OpCapability Float64
+define spir_kernel void @func_kernel(i32 addrspace(1)* %a) {
+entry:
+ tail call spir_func void @func_import(i32 addrspace(1)* %a)
+ ret void
+}
+
+declare spir_func void @func_import(i32 addrspace(1)*)
diff --git a/llvm/test/CodeGen/SPIRV/constant/local-float-point-constants.ll b/llvm/test/CodeGen/SPIRV/constant/local-float-point-constants.ll
index 3a11ceb0d97cb..1bb795f14ab00 100644
--- a/llvm/test/CodeGen/SPIRV/constant/local-float-point-constants.ll
+++ b/llvm/test/CodeGen/SPIRV/constant/local-float-point-constants.ll
@@ -12,23 +12,23 @@ define double @getConstantFP64() {
ret double 0x4f2de42b8c68f3f1
}
-; Capabilities
+;; Capabilities
; CHECK-DAG: OpCapability Float16
; CHECK-DAG: OpCapability Float64
; CHECK-NOT: DAG-FENCE
-; Names:
+;; Names:
; CHECK-DAG: OpName %[[#GET_FP16:]] "getConstantFP16"
; CHECK-DAG: OpName %[[#GET_FP32:]] "getConstantFP32"
; CHECK-DAG: OpName %[[#GET_FP64:]] "getConstantFP64"
; CHECK-NOT: DAG-FENCE
-; Types and Constants:
-; NOTE: These tests don't actually check the values of the constants because
-; their representation isn't defined for textual output.
-; TODO: Test constant representation using binary output.
+;; Types and Constants:
+;; NOTE: These tests don't actually check the values of the constants because
+;; their representation isn't defined for textual output.
+;; TODO: Test constant representation using binary output.
; CHECK-DAG: %[[#FP16:]] = OpTypeFloat 16
; CHECK-DAG: %[[#FP32:]] = OpTypeFloat 32
; CHECK-DAG: %[[#FP64:]] = OpTypeFloat 64
diff --git a/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll b/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll
index cc14327bb3174..0530d1ffa4c62 100644
--- a/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll
+++ b/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll
@@ -16,13 +16,13 @@ define i64 @getLargeConstantI64() {
ret i64 34359738368
}
-; Capabilities:
+;; Capabilities:
; CHECK-DAG: OpCapability Int16
; CHECK-DAG: OpCapability Int64
; CHECK-NOT: DAG-FENCE
-; Names:
+;; Names:
; CHECK-DAG: OpName %[[#GET_I16:]] "getConstantI16"
; CHECK-DAG: OpName %[[#GET_I32:]] "getConstantI32"
; CHECK-DAG: OpName %[[#GET_I64:]] "getConstantI64"
@@ -30,7 +30,7 @@ define i64 @getLargeConstantI64() {
; CHECK-NOT: DAG-FENCE
-; Types and Constants:
+;; Types and Constants:
; CHECK-DAG: %[[#I16:]] = OpTypeInt 16 0
; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0
; CHECK-DAG: %[[#I64:]] = OpTypeInt 64 0
diff --git a/llvm/test/CodeGen/SPIRV/empty-opencl32.ll b/llvm/test/CodeGen/SPIRV/empty-opencl32.ll
new file mode 100644
index 0000000000000..a373781e290bd
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/empty-opencl32.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+;; FIXME: ensure Magic Number, version number, generator's magic number, "bound" and "schema" are at least present
+
+;; Ensure the required Capabilities are listed.
+; CHECK-DAG: OpCapability Kernel
+; CHECK-DAG: OpCapability Addresses
+
+;; Ensure one, and only one, OpMemoryModel is defined.
+; CHECK: OpMemoryModel Physical32 OpenCL
+; CHECK-NOT: OpMemoryModel
diff --git a/llvm/test/CodeGen/SPIRV/empty-opencl64.ll b/llvm/test/CodeGen/SPIRV/empty-opencl64.ll
new file mode 100644
index 0000000000000..d10196525368d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/empty-opencl64.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; FIXME: ensure Magic Number, version number, generator's magic number, "bound" and "schema" are at least present
+
+;; Ensure the required Capabilities are listed.
+; CHECK-DAG: OpCapability Kernel
+; CHECK-DAG: OpCapability Addresses
+
+;; Ensure one, and only one, OpMemoryModel is defined.
+; CHECK: OpMemoryModel Physical64 OpenCL
+; CHECK-NOT: OpMemoryModel
diff --git a/llvm/test/CodeGen/SPIRV/empty.ll b/llvm/test/CodeGen/SPIRV/empty.ll
new file mode 100644
index 0000000000000..fdcf316b02566
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/empty.ll
@@ -0,0 +1,10 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: OpCapability Addresses
+; CHECK: "foo"
+define spir_kernel void @foo(i32 addrspace(1)* %a) {
+entry:
+ %a.addr = alloca i32 addrspace(1)*, align 4
+ store i32 addrspace(1)* %a, i32 addrspace(1)** %a.addr, align 4
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll b/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll
index e59d2e34e7bca..a988b0a3bfff7 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll
@@ -32,12 +32,12 @@ define i32 @no_wrap_test(i32 %a, i32 %b) {
ret i32 %e
}
-; CHECK: %[[#NO_WRAP_TEST]] = OpFunction %[[#I32]] None %[[#FN]]
+; CHECK: %[[#NO_WRAP_TEST]] = OpFunction %[[#I32]] None %[[#FN]]
; CHECK-NEXT: %[[#A]] = OpFunctionParameter %[[#I32]]
; CHECK-NEXT: %[[#B]] = OpFunctionParameter %[[#I32]]
-; CHECK: OpLabel
-; CHECK: %[[#C]] = OpIMul %[[#I32]] %[[#A]] %[[#B]]
-; CHECK: %[[#D]] = OpIMul %[[#I32]] %[[#A]] %[[#B]]
-; CHECK: %[[#E]] = OpIAdd %[[#I32]] %[[#C]] %[[#D]]
-; CHECK: OpReturnValue %[[#E]]
+; CHECK: OpLabel
+; CHECK: %[[#C]] = OpIMul %[[#I32]] %[[#A]] %[[#B]]
+; CHECK: %[[#D]] = OpIMul %[[#I32]] %[[#A]] %[[#B]]
+; CHECK: %[[#E]] = OpIAdd %[[#I32]] %[[#C]] %[[#D]]
+; CHECK: OpReturnValue %[[#E]]
; CHECK-NEXT: OpFunctionEnd
diff --git a/llvm/test/CodeGen/SPIRV/half_extension.ll b/llvm/test/CodeGen/SPIRV/half_extension.ll
new file mode 100644
index 0000000000000..b30e5514c95be
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/half_extension.ll
@@ -0,0 +1,31 @@
+;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable
+;; half test()
+;; {
+;; half x = 0.1f;
+;; x += 2.0f;
+;; half y = x + x;
+;; return y;
+;; }
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability Float16Buffer
+; CHECK-SPIRV-DAG: OpCapability Float16
+
+define spir_func half @test() {
+entry:
+ %x = alloca half, align 2
+ %y = alloca half, align 2
+ store half 0xH2E66, half* %x, align 2
+ %0 = load half, half* %x, align 2
+ %conv = fpext half %0 to float
+ %add = fadd float %conv, 2.000000e+00
+ %conv1 = fptrunc float %add to half
+ store half %conv1, half* %x, align 2
+ %1 = load half, half* %x, align 2
+ %2 = load half, half* %x, align 2
+ %add2 = fadd half %1, %2
+ store half %add2, half* %y, align 2
+ %3 = load half, half* %y, align 2
+ ret half %3
+}
diff --git a/llvm/test/CodeGen/SPIRV/half_no_extension.ll b/llvm/test/CodeGen/SPIRV/half_no_extension.ll
new file mode 100644
index 0000000000000..a5b0ec9c92d23
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/half_no_extension.ll
@@ -0,0 +1,30 @@
+;; __kernel void test( __global float4 *p, __global half *f )
+;; {
+;; __private float4 data;
+;; data = p[0];
+;; vstorea_half4_rtp( data, 0, f );
+;; }
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpCapability Float16Buffer
+; CHECK-SPIRV-NOT: OpCapability Float16
+
+define spir_kernel void @test(<4 x float> addrspace(1)* %p, half addrspace(1)* %f) {
+entry:
+ %p.addr = alloca <4 x float> addrspace(1)*, align 8
+ %f.addr = alloca half addrspace(1)*, align 8
+ %data = alloca <4 x float>, align 16
+ store <4 x float> addrspace(1)* %p, <4 x float> addrspace(1)** %p.addr, align 8
+ store half addrspace(1)* %f, half addrspace(1)** %f.addr, align 8
+ %0 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %p.addr, align 8
+ %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %0, i64 0
+ %1 = load <4 x float>, <4 x float> addrspace(1)* %arrayidx, align 16
+ store <4 x float> %1, <4 x float>* %data, align 16
+ %2 = load <4 x float>, <4 x float>* %data, align 16
+ %3 = load half addrspace(1)*, half addrspace(1)** %f.addr, align 8
+ call spir_func void @_Z17vstorea_half4_rtpDv4_fmPU3AS1Dh(<4 x float> %2, i64 0, half addrspace(1)* %3)
+ ret void
+}
+
+declare spir_func void @_Z17vstorea_half4_rtpDv4_fmPU3AS1Dh(<4 x float>, i64, half addrspace(1)*)
diff --git a/llvm/test/CodeGen/SPIRV/linkage-types.ll b/llvm/test/CodeGen/SPIRV/linkage-types.ll
new file mode 100644
index 0000000000000..e941768204ba8
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/linkage-types.ll
@@ -0,0 +1,112 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIRV
+
+; SPIRV: OpCapability Linkage
+; SPIRV: OpEntryPoint Kernel %[[#kern:]] "kern"
+
+ at ae = available_externally addrspace(1) global i32 79, align 4
+; SPIRV-DAG: OpName %[[#ae:]] "ae"
+
+ at i1 = addrspace(1) global i32 1, align 4
+; SPIRV-DAG: OpName %[[#i1:]] "i1"
+
+ at i2 = internal addrspace(1) global i32 2, align 4
+; SPIRV-DAG: OpName %[[#i2:]] "i2"
+
+ at i3 = addrspace(1) global i32 3, align 4
+; SPIRV-DAG: OpName %[[#i3:]] "i3"
+
+ at i4 = common addrspace(1) global i32 0, align 4
+; SPIRV-DAG: OpName %[[#i4:]] "i4"
+
+ at i5 = internal addrspace(1) global i32 0, align 4
+; SPIRV-DAG: OpName %[[#i5:]] "i5"
+
+ at color_table = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4
+; SPIRV-DAG: OpName %[[#color_table:]] "color_table"
+
+ at noise_table = external addrspace(2) constant [256 x i32]
+; SPIRV-DAG: OpName %[[#noise_table:]] "noise_table"
+
+ at w = addrspace(1) constant i32 0, align 4
+; SPIRV-DAG: OpName %[[#w:]] "w"
+
+ at f.color_table = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4
+; SPIRV-DAG: OpName %[[#f_color_table:]] "f.color_table"
+
+ at e = external addrspace(1) global i32
+; SPIRV-DAG: OpName %[[#e:]] "e"
+
+ at f.t = internal addrspace(1) global i32 5, align 4
+; SPIRV-DAG: OpName %[[#f_t:]] "f.t"
+
+ at f.stint = internal addrspace(1) global i32 0, align 4
+; SPIRV-DAG: OpName %[[#f_stint:]] "f.stint"
+
+ at f.inside = internal addrspace(1) global i32 0, align 4
+; SPIRV-DAG: OpName %[[#f_inside:]] "f.inside"
+
+ at f.b = internal addrspace(2) constant float 1.000000e+00, align 4
+; SPIRV-DAG: OpName %[[#f_b:]] "f.b"
+
+; SPIRV-DAG: OpName %[[#foo:]] "foo"
+; SPIRV-DAG: OpName %[[#f:]] "f"
+; SPIRV-DAG: OpName %[[#g:]] "g"
+; SPIRV-DAG: OpName %[[#inline_fun:]] "inline_fun"
+
+; SPIRV-DAG: OpDecorate %[[#ae]] LinkageAttributes "ae" Import
+; SPIRV-DAG: OpDecorate %[[#e]] LinkageAttributes "e" Import
+; SPIRV-DAG: OpDecorate %[[#f]] LinkageAttributes "f" Export
+; SPIRV-DAG: OpDecorate %[[#w]] LinkageAttributes "w" Export
+; SPIRV-DAG: OpDecorate %[[#i1]] LinkageAttributes "i1" Export
+; SPIRV-DAG: OpDecorate %[[#i3]] LinkageAttributes "i3" Export
+; SPIRV-DAG: OpDecorate %[[#i4]] LinkageAttributes "i4" Export
+; SPIRV-DAG: OpDecorate %[[#foo]] LinkageAttributes "foo" Import
+; SPIRV-DAG: OpDecorate %[[#inline_fun]] LinkageAttributes "inline_fun" Export
+; SPIRV-DAG: OpDecorate %[[#color_table]] LinkageAttributes "color_table" Export
+; SPIRV-DAG: OpDecorate %[[#noise_table]] LinkageAttributes "noise_table" Import
+
+; SPIRV: %[[#foo]] = OpFunction %[[#]]
+declare spir_func void @foo()
+
+; SPIRV: %[[#f]] = OpFunction %[[#]]
+define spir_func void @f() {
+entry:
+ %q = alloca i32, align 4
+ %r = alloca i32, align 4
+ %0 = load i32, i32 addrspace(1)* @i2, align 4
+ store i32 %0, i32* %q, align 4
+ %1 = load i32, i32 addrspace(1)* @i3, align 4
+ store i32 %1, i32 addrspace(1)* @i5, align 4
+ %2 = load i32, i32 addrspace(1)* @e, align 4
+ store i32 %2, i32* %r, align 4
+ %3 = load i32, i32 addrspace(2)* getelementptr inbounds ([256 x i32], [256 x i32] addrspace(2)* @noise_table, i32 0, i32 0), align 4
+ store i32 %3, i32* %r, align 4
+ %4 = load i32, i32 addrspace(2)* getelementptr inbounds ([2 x i32], [2 x i32] addrspace(2)* @f.color_table, i32 0, i32 0), align 4
+ store i32 %4, i32* %r, align 4
+ %call = call spir_func i32 @g()
+ call spir_func void @inline_fun()
+ ret void
+}
+
+; SPIRV: %[[#g]] = OpFunction %[[#]]
+define internal spir_func i32 @g() {
+entry:
+ call spir_func void @foo()
+ ret i32 25
+}
+
+; SPIRV: %[[#inline_fun]] = OpFunction %[[#]]
+;; "linkonce_odr" is lost in translation !
+define linkonce_odr spir_func void @inline_fun() {
+entry:
+ %t = alloca i32 addrspace(1)*, align 4
+ store i32 addrspace(1)* @i1, i32 addrspace(1)** %t, align 4
+ ret void
+}
+
+; SPIRV: %[[#kern]] = OpFunction %[[#]]
+define spir_kernel void @kern() {
+entry:
+ call spir_func void @f()
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/NoSignedUnsignedWrap.ll b/llvm/test/CodeGen/SPIRV/transcoding/NoSignedUnsignedWrap.ll
new file mode 100644
index 0000000000000..5a67dc2599421
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/NoSignedUnsignedWrap.ll
@@ -0,0 +1,29 @@
+;; Source
+;; int square(unsigned short a) {
+;; return a * a;
+;; }
+;; Command
+;; clang -cc1 -triple spir -emit-llvm -O2 -o NoSignedUnsignedWrap.ll test.cl
+;;
+;; Positive tests:
+;;
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-NEGATIVE
+;;
+;; Negative tests:
+;;
+;; Check that backend is able to skip nsw/nuw attributes if extension is
+;; disabled implicitly or explicitly and if max SPIR-V version is lower then 1.4
+
+; CHECK-SPIRV-DAG: OpDecorate %[[#]] NoSignedWrap
+; CHECK-SPIRV-DAG: OpDecorate %[[#]] NoUnsignedWrap
+;
+; CHECK-SPIRV-NEGATIVE-NOT: OpExtension "SPV_KHR_no_integer_wrap_decoration"
+; CHECK-SPIRV-NEGATIVE-NOT: OpDecorate %[[#]] NoSignedWrap
+; CHECK-SPIRV-NEGATIVE-NOT: OpDecorate %[[#]] NoUnsignedWrap
+
+define spir_func i32 @square(i16 zeroext %a) local_unnamed_addr {
+entry:
+ %conv = zext i16 %a to i32
+ %mul = mul nuw nsw i32 %conv, %conv
+ ret i32 %mul
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ReqdSubgroupSize.ll b/llvm/test/CodeGen/SPIRV/transcoding/ReqdSubgroupSize.ll
index 04486b79705b6..2bc0a95e56be2 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/ReqdSubgroupSize.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ReqdSubgroupSize.ll
@@ -1,7 +1,7 @@
-; Check translation of intel_reqd_sub_group_size metadata to SubgroupSize
-; execution mode and back. The IR is producded from the following OpenCL C code:
-; kernel __attribute__((intel_reqd_sub_group_size(8)))
-; void foo() {}
+;; Check translation of intel_reqd_sub_group_size metadata to SubgroupSize
+;; execution mode and back. The IR is producded from the following OpenCL C code:
+;; kernel __attribute__((intel_reqd_sub_group_size(8)))
+;; void foo() {}
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fadd.ll b/llvm/test/CodeGen/SPIRV/transcoding/fadd.ll
index 0cb9d99eaefb9..78d9a23266558 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/fadd.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fadd.ll
@@ -1,12 +1,12 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; CHECK-SPIRV: OpName %[[#r1:]] "r1"
-; CHECK-SPIRV: OpName %[[#r2:]] "r2"
-; CHECK-SPIRV: OpName %[[#r3:]] "r3"
-; CHECK-SPIRV: OpName %[[#r4:]] "r4"
-; CHECK-SPIRV: OpName %[[#r5:]] "r5"
-; CHECK-SPIRV: OpName %[[#r6:]] "r6"
-; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
; CHECK-SPIRV-NOT: OpDecorate %[[#r1]] FPFastMathMode
; CHECK-SPIRV-DAG: OpDecorate %[[#r2]] FPFastMathMode NotNaN
; CHECK-SPIRV-DAG: OpDecorate %[[#r3]] FPFastMathMode NotInf
@@ -14,14 +14,14 @@
; CHECK-SPIRV-DAG: OpDecorate %[[#r5]] FPFastMathMode AllowRecip
; CHECK-SPIRV-DAG: OpDecorate %[[#r6]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|Fast
; CHECK-SPIRV-DAG: OpDecorate %[[#r7]] FPFastMathMode NotNaN|NotInf
-; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
-; CHECK-SPIRV: %[[#r1]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r2]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r3]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r4]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r5]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r6]] = OpFAdd %[[#float]]
-; CHECK-SPIRV: %[[#r7]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFAdd %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFAdd %[[#float]]
define spir_kernel void @testFAdd(float %a, float %b) {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fdiv.ll b/llvm/test/CodeGen/SPIRV/transcoding/fdiv.ll
new file mode 100644
index 0000000000000..d0ed5640e7066
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fdiv.ll
@@ -0,0 +1,36 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV-NOT: OpDecorate %[[#r1]] FPFastMathMode
+; CHECK-SPIRV-DAG: OpDecorate %[[#r2]] FPFastMathMode NotNaN
+; CHECK-SPIRV-DAG: OpDecorate %[[#r3]] FPFastMathMode NotInf
+; CHECK-SPIRV-DAG: OpDecorate %[[#r4]] FPFastMathMode NSZ
+; CHECK-SPIRV-DAG: OpDecorate %[[#r5]] FPFastMathMode AllowRecip
+; CHECK-SPIRV-DAG: OpDecorate %[[#r6]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|Fast
+; CHECK-SPIRV-DAG: OpDecorate %[[#r7]] FPFastMathMode NotNaN|NotInf
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFDiv %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFDiv %[[#float]]
+
+define spir_kernel void @testFDiv(float %a, float %b) local_unnamed_addr {
+entry:
+ %r1 = fdiv float %a, %b
+ %r2 = fdiv nnan float %a, %b
+ %r3 = fdiv ninf float %a, %b
+ %r4 = fdiv nsz float %a, %b
+ %r5 = fdiv arcp float %a, %b
+ %r6 = fdiv fast float %a, %b
+ %r7 = fdiv nnan ninf float %a, %b
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fmul.ll b/llvm/test/CodeGen/SPIRV/transcoding/fmul.ll
new file mode 100644
index 0000000000000..886077a67b4e6
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fmul.ll
@@ -0,0 +1,36 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV-NOT: OpDecorate %[[#r1]] FPFastMathMode
+; CHECK-SPIRV-DAG: OpDecorate %[[#r2]] FPFastMathMode NotNaN
+; CHECK-SPIRV-DAG: OpDecorate %[[#r3]] FPFastMathMode NotInf
+; CHECK-SPIRV-DAG: OpDecorate %[[#r4]] FPFastMathMode NSZ
+; CHECK-SPIRV-DAG: OpDecorate %[[#r5]] FPFastMathMode AllowRecip
+; CHECK-SPIRV-DAG: OpDecorate %[[#r6]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|Fast
+; CHECK-SPIRV-DAG: OpDecorate %[[#r7]] FPFastMathMode NotNaN|NotInf
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFMul %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFMul %[[#float]]
+
+define spir_kernel void @testFMul(float %a, float %b) local_unnamed_addr {
+entry:
+ %r1 = fmul float %a, %b
+ %r2 = fmul nnan float %a, %b
+ %r3 = fmul ninf float %a, %b
+ %r4 = fmul nsz float %a, %b
+ %r5 = fmul arcp float %a, %b
+ %r6 = fmul fast float %a, %b
+ %r7 = fmul nnan ninf float %a, %b
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/frem.ll b/llvm/test/CodeGen/SPIRV/transcoding/frem.ll
new file mode 100644
index 0000000000000..ecb8f6f950cab
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/frem.ll
@@ -0,0 +1,36 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV-NOT: OpDecorate %[[#r1]] FPFastMathMode
+; CHECK-SPIRV-DAG: OpDecorate %[[#r2]] FPFastMathMode NotNaN
+; CHECK-SPIRV-DAG: OpDecorate %[[#r3]] FPFastMathMode NotInf
+; CHECK-SPIRV-DAG: OpDecorate %[[#r4]] FPFastMathMode NSZ
+; CHECK-SPIRV-DAG: OpDecorate %[[#r5]] FPFastMathMode AllowRecip
+; CHECK-SPIRV-DAG: OpDecorate %[[#r6]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|Fast
+; CHECK-SPIRV-DAG: OpDecorate %[[#r7]] FPFastMathMode NotNaN|NotInf
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFRem %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFRem %[[#float]]
+
+define spir_kernel void @testFRem(float %a, float %b) local_unnamed_addr {
+entry:
+ %r1 = frem float %a, %b
+ %r2 = frem nnan float %a, %b
+ %r3 = frem ninf float %a, %b
+ %r4 = frem nsz float %a, %b
+ %r5 = frem arcp float %a, %b
+ %r6 = frem fast float %a, %b
+ %r7 = frem nnan ninf float %a, %b
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fsub.ll b/llvm/test/CodeGen/SPIRV/transcoding/fsub.ll
new file mode 100644
index 0000000000000..99d0d0eb84f95
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fsub.ll
@@ -0,0 +1,36 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV-NOT: OpDecorate %[[#r1]] FPFastMathMode
+; CHECK-SPIRV-DAG: OpDecorate %[[#r2]] FPFastMathMode NotNaN
+; CHECK-SPIRV-DAG: OpDecorate %[[#r3]] FPFastMathMode NotInf
+; CHECK-SPIRV-DAG: OpDecorate %[[#r4]] FPFastMathMode NSZ
+; CHECK-SPIRV-DAG: OpDecorate %[[#r5]] FPFastMathMode AllowRecip
+; CHECK-SPIRV-DAG: OpDecorate %[[#r6]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|Fast
+; CHECK-SPIRV-DAG: OpDecorate %[[#r7]] FPFastMathMode NotNaN|NotInf
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFSub %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFSub %[[#float]]
+
+define spir_kernel void @testFSub(float %a, float %b) local_unnamed_addr {
+entry:
+ %r1 = fsub float %a, %b
+ %r2 = fsub nnan float %a, %b
+ %r3 = fsub ninf float %a, %b
+ %r4 = fsub nsz float %a, %b
+ %r5 = fsub arcp float %a, %b
+ %r6 = fsub fast float %a, %b
+ %r7 = fsub nnan ninf float %a, %b
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/vec8.ll b/llvm/test/CodeGen/SPIRV/transcoding/vec8.ll
new file mode 100644
index 0000000000000..e055e0a1231c4
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/vec8.ll
@@ -0,0 +1,15 @@
+;; This test verifies that the Vector16 capability is correctly added
+;; if an OpenCL kernel uses a vector of eight elements.
+;;
+;; Source:
+;; __kernel void test( int8 v ) {}
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpCapability Vector16
+
+define spir_kernel void @test(<8 x i32> %v) {
+ %1 = alloca <8 x i32>, align 32
+ store <8 x i32> %v, <8 x i32>* %1, align 32
+ ret void
+}
More information about the llvm-commits
mailing list