diff --git a/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll b/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll new file mode 100644 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/CheckCapKernelWithoutKernel.ll @@ -0,0 +1,5 @@ +; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s + +@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 --- /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 + +@a = addrspace(1) global i8 0, align 1 +@b = addrspace(1) global i16 0, align 2 +@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 --- /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 --- 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 @@ 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 --- a/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll +++ b/llvm/test/CodeGen/SPIRV/constant/local-integers-constants.ll @@ -16,13 +16,13 @@ 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 @@ ; 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 --- /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 --- /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 --- /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 --- a/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/no_wrap.ll @@ -32,12 +32,12 @@ 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 --- /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 --- /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 --- /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" + +@ae = available_externally addrspace(1) global i32 79, align 4 +; SPIRV-DAG: OpName %[[#ae:]] "ae" + +@i1 = addrspace(1) global i32 1, align 4 +; SPIRV-DAG: OpName %[[#i1:]] "i1" + +@i2 = internal addrspace(1) global i32 2, align 4 +; SPIRV-DAG: OpName %[[#i2:]] "i2" + +@i3 = addrspace(1) global i32 3, align 4 +; SPIRV-DAG: OpName %[[#i3:]] "i3" + +@i4 = common addrspace(1) global i32 0, align 4 +; SPIRV-DAG: OpName %[[#i4:]] "i4" + +@i5 = internal addrspace(1) global i32 0, align 4 +; SPIRV-DAG: OpName %[[#i5:]] "i5" + +@color_table = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4 +; SPIRV-DAG: OpName %[[#color_table:]] "color_table" + +@noise_table = external addrspace(2) constant [256 x i32] +; SPIRV-DAG: OpName %[[#noise_table:]] "noise_table" + +@w = addrspace(1) constant i32 0, align 4 +; SPIRV-DAG: OpName %[[#w:]] "w" + +@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" + +@e = external addrspace(1) global i32 +; SPIRV-DAG: OpName %[[#e:]] "e" + +@f.t = internal addrspace(1) global i32 5, align 4 +; SPIRV-DAG: OpName %[[#f_t:]] "f.t" + +@f.stint = internal addrspace(1) global i32 0, align 4 +; SPIRV-DAG: OpName %[[#f_stint:]] "f.stint" + +@f.inside = internal addrspace(1) global i32 0, align 4 +; SPIRV-DAG: OpName %[[#f_inside:]] "f.inside" + +@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 --- /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 --- 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 --- 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 --- /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 --- /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 --- /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 --- /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 --- /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 +}