diff options
Diffstat (limited to 'llvm/test/CodeGen/SPIRV')
32 files changed, 1071 insertions, 90 deletions
diff --git a/llvm/test/CodeGen/SPIRV/GlobalVarAddrspace.ll b/llvm/test/CodeGen/SPIRV/GlobalVarAddrspace.ll new file mode 100644 index 0000000..2bccfde --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/GlobalVarAddrspace.ll @@ -0,0 +1,23 @@ +; This test case checks that LLVM -> SPIR-V translation produces valid +; SPIR-V module, where a global variable, defined with non-default +; address space, have correct non-function storage class. +; +; No additional checks are needed in addition to simple translation +; to SPIR-V. In case of an error newly produced SPIR-V module validation +; would fail due to spirv-val that detects problematic SPIR-V code from +; translator and reports it as the following error: +; +; "Variables can not have a function[7] storage class outside of a function". +; +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: %[[#Ptr:]] = OpTypePointer CrossWorkgroup %[[#]] +; CHECK: %[[#]] = OpVariable %[[#Ptr]] CrossWorkgroup %[[#]] + +@G = addrspace(1) global i1 true + +define spir_func i1 @f(i1 %0) { + store i1 %0, ptr addrspace(1) @G, align 1 + ret i1 %0 +} diff --git a/llvm/test/CodeGen/SPIRV/SamplerArgNonKernel.ll b/llvm/test/CodeGen/SPIRV/SamplerArgNonKernel.ll new file mode 100644 index 0000000..5b3a5d8 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/SamplerArgNonKernel.ll @@ -0,0 +1,37 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +;CHECK: OpEntryPoint Kernel %[[#KernelId:]] +;CHECK: %[[#image2d_t:]] = OpTypeImage +;CHECK: %[[#sampler_t:]] = OpTypeSampler +;CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage + +define spir_func float @test(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %Img, target("spirv.Sampler") %Smp) { +;CHECK-NOT: %[[#KernelId]] = OpFunction %[[#]] +;CHECK: OpFunction +;CHECK: %[[#image:]] = OpFunctionParameter %[[#image2d_t]] +;CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]] +entry: + %call = call spir_func <4 x i32> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %Img, target("spirv.Sampler") %Smp, <2 x i32> zeroinitializer) +;CHECK: %[[#sampled_image:]] = OpSampledImage %[[#sampled_image_t]] %[[#image]] %[[#sampler]] +;CHECK: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#sampled_image]] %[[#]] Lod %[[#]] + + %0 = extractelement <4 x i32> %call, i32 0 + %conv = sitofp i32 %0 to float + ret float %conv +} + +declare spir_func <4 x i32> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), i32, <2 x i32>) + +define spir_kernel void @test2(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %Img, target("spirv.Sampler") %Smp, ptr addrspace(1) %result) { +;CHECK: %[[#KernelId]] = OpFunction %[[#]] +entry: + %call = call spir_func float @test(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %Img, target("spirv.Sampler") %Smp) + %0 = load float, ptr addrspace(1) %result, align 4 + %add = fadd float %0, %call + store float %add, ptr addrspace(1) %result, align 4 + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/SpecConstants/spec-constant-length-array.ll b/llvm/test/CodeGen/SPIRV/SpecConstants/spec-constant-length-array.ll new file mode 100644 index 0000000..fccddd7 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/SpecConstants/spec-constant-length-array.ll @@ -0,0 +1,56 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_variable_length_array %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpCapability VariableLengthArrayINTEL +; CHECK: OpExtension "SPV_INTEL_variable_length_array" + +; CHECK-DAG: OpDecorate %[[#]] SpecId 0 +; CHECK-DAG: OpDecorate %[[#]] SpecId 1 +; CHECK-DAG: OpDecorate %[[#]] SpecId 2 +; CHECK-DAG: OpDecorate %[[#A0:]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#A1:]] Alignment 2 +; CHECK-DAG: OpDecorate %[[#A2:]] Alignment 16 + +; CHECK: %[[#VOID_TY:]] = OpTypeVoid +; CHECK: %[[#FUNC_TY:]] = OpTypeFunction %[[#VOID_TY]] +; CHECK-DAG: %[[#I64:]] = OpTypeInt 64 0 +; CHECK-DAG: %[[#I32:]] = OpTypeInt 32 0 +; CHECK-DAG: %[[#I8:]] = OpTypeInt 8 0 +; CHECK-DAG: %[[#F64:]] = OpTypeFloat 64 +; CHECK-DAG: %[[#STRUCT_TY:]] = OpTypeStruct %[[#F64]] %[[#F64]] +; CHECK-DAG: %[[#PTR_STRUCT:]] = OpTypePointer Function %[[#STRUCT_TY]] +; CHECK-DAG: %[[#PTR_I8:]] = OpTypePointer Function %[[#I8]] +; CHECK-DAG: %[[#F32:]] = OpTypeFloat 32 +; CHECK-DAG: %[[#PTR_F32:]] = OpTypePointer Function %[[#F32]] + +; CHECK-DAG: %[[#SC0:]] = OpSpecConstant %[[#I64]] 1 +; CHECK-DAG: %[[#SC1:]] = OpSpecConstant %[[#I32]] 2 +; CHECK-DAG: %[[#SC2:]] = OpSpecConstant %[[#I8]] 4 + +; CHECK: %[[#]] = OpFunction %[[#VOID_TY]] None %[[#FUNC_TY]] +; CHECK: %[[#LABEL:]] = OpLabel + +; CHECK: %[[#A0]] = OpVariableLengthArrayINTEL %[[#PTR_F32]] %[[#SC0]] +; CHECK: %[[#A1]] = OpVariableLengthArrayINTEL %[[#PTR_I8]] %[[#SC1]] +; CHECK: %[[#A2]] = OpVariableLengthArrayINTEL %[[#PTR_STRUCT]] %[[#SC2]] + +%struct_type = type { double, double } + +define spir_kernel void @test() { + entry: + %length0 = call i64 @_Z20__spirv_SpecConstantix(i32 0, i64 1), !SYCL_SPEC_CONST_SYM_ID !0 + %length1 = call i32 @_Z20__spirv_SpecConstantii(i32 1, i32 2), !SYCL_SPEC_CONST_SYM_ID !1 + %length2 = call i8 @_Z20__spirv_SpecConstantic(i32 2, i8 4), !SYCL_SPEC_CONST_SYM_ID !2 + %scla0 = alloca float, i64 %length0, align 4 + %scla1 = alloca i8, i32 %length1, align 2 + %scla2 = alloca %struct_type, i8 %length2, align 16 + ret void +} + +declare i8 @_Z20__spirv_SpecConstantic(i32, i8) +declare i32 @_Z20__spirv_SpecConstantii(i32, i32) +declare i64 @_Z20__spirv_SpecConstantix(i32, i64) + +!0 = !{!"i64_spec_const", i32 0} +!1 = !{!"i32_spec_const", i32 1} +!2 = !{!"i8_spec_const", i32 2} diff --git a/llvm/test/CodeGen/SPIRV/align-duplicate.ll b/llvm/test/CodeGen/SPIRV/align-duplicate.ll new file mode 100644 index 0000000..8a8d8ae --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/align-duplicate.ll @@ -0,0 +1,16 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; Test that duplicate align information does not result in SPIR-V validation +; errors due to duplicate Alignment Decorations. + +;CHECK: OpDecorate %[[#Var:]] Alignment +;CHECK: %[[#Var]] = OpVariable %[[#]] + +define spir_func void @f() { + %res = alloca i16, align 2, !spirv.Decorations !1 + ret void +} + +!1 = !{!2} +!2 = !{i32 44, i32 2} diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll new file mode 100644 index 0000000..251b48f --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/array_type.ll @@ -0,0 +1,78 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %} + +; CHECK: OpCapability Kernel +; CHECK-NOT: OpCapability Shader +; CHECK-DAG: %[[#float16:]] = OpTypeFloat 16 +; CHECK-DAG: %[[#SyclHalfTy:]] = OpTypeStruct %[[#float16]] +; CHECK-DAG: %[[#i16:]] = OpTypeInt 16 +; CHECK-DAG: %[[#i32:]] = OpTypeInt 32 +; CHECK-DAG: %[[#i64:]] = OpTypeInt 64 +; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#i64]] +; CHECK-DAG: %[[#ConstOne:]] = OpConstant %[[#i64]] 1 +; CHECK-DAG: %[[#ConstFive:]] = OpConstant %[[#i16]] 5 +; CHECK-DAG: %[[#SyclHalfTyPtr:]] = OpTypePointer Function %[[#SyclHalfTy]] +; CHECK-DAG: %[[#i32Ptr:]] = OpTypePointer Function %[[#i32]] +; CHECK-DAG: %[[#StorePtrTy:]] = OpTypePointer Function %[[#i16]] + +%"class.sycl::_V1::detail::half_impl::half" = type { half } + +; Function Attrs: mustprogress norecurse nounwind +define spir_kernel void @foo(ptr %p){ +; CHECK: OpFunction +; CHECK: %[[#Ptr:]] = OpFunctionParameter +; CHECK: OpLabel +; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]] +; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]] +; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]] +; CHECK: OpReturn +entry: + %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 0, i64 0 + store i16 5, ptr %0 + ret void +} + +; Function Attrs: mustprogress norecurse nounwind +define spir_kernel void @foo2(ptr %p){ +; CHECK: OpFunction +; CHECK: %[[#Ptr:]] = OpFunctionParameter +; CHECK: OpLabel +; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]] +; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]] +; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]] +; CHECK: OpReturn +entry: + %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 1, i64 1 + store i16 5, ptr %0 + ret void +} + +; Function Attrs: mustprogress norecurse nounwind +define spir_kernel void @foo3(ptr %p){ +; CHECK: OpFunction +; CHECK: %[[#Ptr:]] = OpFunctionParameter +; CHECK: OpLabel +; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]] +; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]] +; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]] +; CHECK: OpReturn +entry: + %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 0, i64 0 + store i16 5, ptr %0 + ret void +} + +; Function Attrs: mustprogress norecurse nounwind +define spir_kernel void @foo4(ptr %p){ +; CHECK: OpFunction +; CHECK: %[[#Ptr:]] = OpFunctionParameter +; CHECK: OpLabel +; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]] +; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]] +; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]] +; CHECK: OpReturn +entry: + %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 1, i64 1 + store i16 5, ptr %0 + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/duplicate-types.ll b/llvm/test/CodeGen/SPIRV/duplicate-types.ll new file mode 100644 index 0000000..df1ae04 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/duplicate-types.ll @@ -0,0 +1,16 @@ +; Check that we don't end up with duplicated array types in TypeMap. +; No FileCheck needed, we only want to check the absence of errors. +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: %[[#]] = OpTypeArray %[[#]] %[[#]] +; CHECK-NOT: OpTypeArray + +%duplicate = type { [2 x ptr addrspace(4)] } + +define spir_kernel void @foo() { +entry: + alloca [2 x ptr addrspace(4)], align 8 + alloca %duplicate, align 8 + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/entry-point-interfaces.ll b/llvm/test/CodeGen/SPIRV/entry-point-interfaces.ll new file mode 100644 index 0000000..f1e0927 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/entry-point-interfaces.ll @@ -0,0 +1,31 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpEntryPoint Kernel %[[#Func:]] "test" %[[#Interface1:]] %[[#Interface2:]] %[[#Interface3:]] %[[#Interface4:]] +; CHECK-DAG: OpName %[[#Func]] "test" +; CHECK-DAG: OpName %[[#Interface1]] "var" +; CHECK-DAG: OpName %[[#Interface3]] "var2" +; CHECK-DAG: OpName %[[#Interface2]] "var.const" +; CHECK-DAG: OpName %[[#Interface4]] "var2.const" +; CHECK-DAG: %[[#TypeInt:]] = OpTypeInt 32 0 +; CHECK-DAG: %[[#Const1:]] = OpConstant %[[#TypeInt]] 1 +; CHECK-DAG: %[[#Const2:]] = OpConstant %[[#TypeInt]] 3 + +; CHECK: %[[#Interface1]] = OpVariable %[[#]] UniformConstant %[[#Const1]] +; CHECK: %[[#Interface3]] = OpVariable %[[#]] UniformConstant %[[#Const2]] +; CHECK: %[[#Interface2]] = OpVariable %[[#]] UniformConstant %[[#Const1]] +; CHECK: %[[#Interface4]] = OpVariable %[[#]] UniformConstant %[[#Const2]] + +@var = dso_local addrspace(2) constant i32 1, align 4 +@var2 = dso_local addrspace(2) constant i32 3, align 4 +@var.const = private unnamed_addr addrspace(2) constant i32 1, align 4 +@var2.const = private unnamed_addr addrspace(2) constant i32 3, align 4 + +define dso_local spir_kernel void @test() { +entry: + %0 = load i32, ptr addrspace(2) @var.const, align 4 + %1 = load i32, ptr addrspace(2) @var2.const, align 4 + %mul = mul nsw i32 %0, %1 + %mul1 = mul nsw i32 %mul, 2 + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll index 40e2aff..7adb039 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll @@ -1,18 +1,11 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION -; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %} ; CHECK-EXTENSION: OpCapability BitInstructions ; CHECK-EXTENSION-NEXT: OpExtension "SPV_KHR_bit_instructions" ; CHECK-EXTENSION-NOT: OpCabilitity Shader -; CHECK-NO-EXTENSION: OpCapability Shader -; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions -; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions" - - ; CHECK-EXTENSION: %[[#int:]] = OpTypeInt 32 ; CHECK-EXTENSION: OpBitReverse %[[#int]] -; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32 -; CHECK-NO-EXTENSION: OpBitReverse %[[#int]] define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr { entry: diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll index 65cccc8..3bd1bd6 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll @@ -1,12 +1,8 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION -; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION -; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val %} +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %} ; ; CHECK-EXTENSION: Capability BitInstructions ; CHECK-EXTENSION: Extension "SPV_KHR_bit_instructions" -; CHECK-NO-EXTENSION-NOT: Capability BitInstructions -; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions" -; CHECK-NO-EXTENSION: Capability Shader ; ; CHECK-EXTENSION: %[[#]] = OpFunction %[[#]] None %[[#]] ; CHECK-EXTENSION: %[[#reversebase:]] = OpFunctionParameter %[[#]] @@ -15,24 +11,11 @@ ; kernel void testBitReverse_SPIRVFriendly(long4 b, global long4 *res) { ; *res = bit_reverse(b); ; } -define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) nocapture align 32 %res) #3 { +define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) %res) { entry: %call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> %b) - store <4 x i64> %call, ptr addrspace(1) %res, align 32 + store <4 x i64> %call, ptr addrspace(1) %res ret void } -declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>) #4 - - -attributes #3 = { nounwind } -attributes #4 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) } - -!llvm.module.flags = !{!0} -!opencl.ocl.version = !{!1} -!opencl.spir.version = !{!1} -!llvm.ident = !{!2} - -!0 = !{i32 1, !"wchar_size", i32 4} -!1 = !{i32 2, i32 0} -!2 = !{!"clang version 20.0.0git (https://github.com/llvm/llvm-project.git cc61409d353a40f62d3a137f3c7436aa00df779d)"} +declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>) diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll new file mode 100644 index 0000000..61ef273 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll @@ -0,0 +1,16 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} +; +; CHECK-NO-EXTENSION-NOT: Capability BitInstructions +; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions" +; CHECK-NO-EXTENSION: Capability Shader + +define internal spir_func void @testBitReverse_SPIRVFriendly() #3 { +entry: + %call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> <i64 1, i64 2, i64 3, i64 4>) + ret void +} + +declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>) + +attributes #3 = { nounwind "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll new file mode 100644 index 0000000..452df0a --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll @@ -0,0 +1,23 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %} + + +; CHECK-NO-EXTENSION: OpCapability Shader +; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions +; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions" +; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32 +; CHECK-NO-EXTENSION: OpBitReverse %[[#int]] + +define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) local_unnamed_addr { +entry: + %call = tail call i32 @llvm.bitreverse.i32(i32 %b) + store i32 %call, ptr %res, align 4 + ret void +} + +define void @main() #1 { + ret void +} + +declare i32 @llvm.bitreverse.i32(i32) +attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/get_global_size.ll b/llvm/test/CodeGen/SPIRV/get_global_size.ll new file mode 100644 index 0000000..959371a7 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/get_global_size.ll @@ -0,0 +1,50 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: %[[#int32:]] = OpTypeInt 32 0 +; CHECK: %[[#int64:]] = OpTypeInt 64 0 +; CHECK: %[[#vec3:]] = OpTypeVector %[[#int64]] 3 +; CHECK: %[[#ptr_input_vec3:]] = OpTypePointer Input %[[#vec3]] +; CHECK: %[[#global_size_var:]] = OpVariable %[[#ptr_input_vec3]] Input + +; CHECK: %[[#load_gs1:]] = OpLoad %[[#vec3]] %[[#global_size_var]] Aligned 1 +; CHECK: %[[#extract3:]] = OpCompositeExtract %[[#int64]] %[[#load_gs1]] 0 + +; CHECK: %[[#bitcast1:]] = OpBitcast %[[#]] %[[#]] +; CHECK: %[[#load_out1:]] = OpLoad %[[#]] %[[#bitcast1]] Aligned 8 +; CHECK: %[[#gep1:]] = OpInBoundsPtrAccessChain %[[#]] %[[#load_out1]] %[[#]] +; CHECK: OpStore %[[#gep1]] %[[#extract3]] Aligned 8 + +; CHECK: %[[#load_param_x:]] = OpLoad %[[#int32]] %[[#]] +; CHECK: %[[#load_gs2:]] = OpLoad %[[#vec3]] %[[#global_size_var]] Aligned 1 +; CHECK: %[[#dyn_extract:]] = OpVectorExtractDynamic %[[#int64]] %[[#load_gs2]] %[[#load_param_x]] +; CHECK: %[[#cmp:]] = OpULessThan %[[#]] %[[#load_param_x]] %[[#]] +; CHECK: %[[#select2:]] = OpSelect %[[#int64]] %[[#cmp]] %[[#dyn_extract]] %[[#]] +; CHECK: %[[#bitcast2:]] = OpBitcast %[[#]] %[[#]] +; CHECK: %[[#load_out2:]] = OpLoad %[[#]] %[[#bitcast2]] Aligned 8 +; CHECK: %[[#gep2:]] = OpInBoundsPtrAccessChain %[[#]] %[[#load_out2]] %[[#]] +; CHECK: OpStore %[[#gep2]] %[[#select2]] Aligned 8 + +define dso_local spir_kernel void @ggs(ptr noundef align 8 %out, i32 noundef %x) { +entry: + %out.addr = alloca ptr, align 8 + %x.addr = alloca i32, align 4 + store ptr %out, ptr %out.addr, align 8 + store i32 %x, ptr %x.addr, align 4 + %call = call i64 @_Z15get_global_sizej(i32 noundef 0) + %0 = load ptr, ptr %out.addr, align 8 + %arrayidx = getelementptr inbounds i64, ptr %0, i64 0 + store i64 %call, ptr %arrayidx, align 8 + %call1 = call i64 @_Z15get_global_sizej(i32 noundef 3) + %1 = load ptr, ptr %out.addr, align 8 + %arrayidx2 = getelementptr inbounds i64, ptr %1, i64 1 + store i64 %call1, ptr %arrayidx2, align 8 + %2 = load i32, ptr %x.addr, align 4 + %call3 = call i64 @_Z15get_global_sizej(i32 noundef %2) + %3 = load ptr, ptr %out.addr, align 8 + %arrayidx4 = getelementptr inbounds i64, ptr %3, i64 2 + store i64 %call3, ptr %arrayidx4, align 8 + ret void +} + +declare i64 @_Z15get_global_sizej(i32 noundef) diff --git a/llvm/test/CodeGen/SPIRV/hlsl-resources/ImplicitBinding.ll b/llvm/test/CodeGen/SPIRV/hlsl-resources/ImplicitBinding.ll new file mode 100644 index 0000000..00e9185 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/hlsl-resources/ImplicitBinding.ll @@ -0,0 +1,75 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv1.6-vulkan1.3-library %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv1.6-vulkan1.3-library %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %} + +@.str = private unnamed_addr constant [2 x i8] c"b\00", align 1 +@.str.2 = private unnamed_addr constant [2 x i8] c"c\00", align 1 +@.str.4 = private unnamed_addr constant [2 x i8] c"d\00", align 1 +@.str.6 = private unnamed_addr constant [2 x i8] c"e\00", align 1 +@.str.8 = private unnamed_addr constant [2 x i8] c"f\00", align 1 +@.str.10 = private unnamed_addr constant [2 x i8] c"g\00", align 1 +@.str.12 = private unnamed_addr constant [2 x i8] c"h\00", align 1 +@.str.14 = private unnamed_addr constant [2 x i8] c"i\00", align 1 + +; CHECK-DAG: OpName [[b:%[0-9]+]] "b" +; CHECK-DAG: OpName [[c:%[0-9]+]] "c" +; CHECK-DAG: OpName [[d:%[0-9]+]] "d" +; CHECK-DAG: OpName [[e:%[0-9]+]] "e" +; CHECK-DAG: OpName [[f:%[0-9]+]] "f" +; CHECK-DAG: OpName [[g:%[0-9]+]] "g" +; CHECK-DAG: OpName [[h:%[0-9]+]] "h" +; CHECK-DAG: OpName [[i:%[0-9]+]] "i" +; CHECK-DAG: OpDecorate [[b]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[b]] Binding 1 +; CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c]] Binding 0 +; CHECK-DAG: OpDecorate [[d]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[d]] Binding 3 +; CHECK-DAG: OpDecorate [[e]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[e]] Binding 2 +; CHECK-DAG: OpDecorate [[f]] DescriptorSet 10 +; CHECK-DAG: OpDecorate [[f]] Binding 1 +; CHECK-DAG: OpDecorate [[g]] DescriptorSet 10 +; CHECK-DAG: OpDecorate [[g]] Binding 0 +; CHECK-DAG: OpDecorate [[h]] DescriptorSet 10 +; CHECK-DAG: OpDecorate [[h]] Binding 3 +; CHECK-DAG: OpDecorate [[i]] DescriptorSet 10 +; CHECK-DAG: OpDecorate [[i]] Binding 2 + + +define void @main() local_unnamed_addr #0 { +entry: + %0 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefromimplicitbinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 0, i32 0, i32 1, i32 0, i1 false, ptr nonnull @.str) + %1 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefrombinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 0, i32 0, i32 1, i32 0, i1 false, ptr nonnull @.str.2) + %2 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefromimplicitbinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 1, i32 0, i32 1, i32 0, i1 false, ptr nonnull @.str.4) + %3 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefrombinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 0, i32 2, i32 1, i32 0, i1 false, ptr nonnull @.str.6) + %4 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefrombinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 10, i32 1, i32 1, i32 0, i1 false, ptr nonnull @.str.8) + %5 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefromimplicitbinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 2, i32 10, i32 1, i32 0, i1 false, ptr nonnull @.str.10) + %6 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefromimplicitbinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 3, i32 10, i32 1, i32 0, i1 false, ptr nonnull @.str.12) + %7 = tail call target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) @llvm.spv.resource.handlefrombinding.tspirv.SignedImage_i32_5_2_0_0_2_0t(i32 10, i32 2, i32 1, i32 0, i1 false, ptr nonnull @.str.14) + %8 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %1, i32 0) + %9 = load i32, ptr addrspace(11) %8, align 4 + %10 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %2, i32 0) + %11 = load i32, ptr addrspace(11) %10, align 4 + %add.i = add nsw i32 %11, %9 + %12 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %3, i32 0) + %13 = load i32, ptr addrspace(11) %12, align 4 + %add4.i = add nsw i32 %add.i, %13 + %14 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %4, i32 0) + %15 = load i32, ptr addrspace(11) %14, align 4 + %add6.i = add nsw i32 %add4.i, %15 + %16 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %5, i32 0) + %17 = load i32, ptr addrspace(11) %16, align 4 + %add8.i = add nsw i32 %add6.i, %17 + %18 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %6, i32 0) + %19 = load i32, ptr addrspace(11) %18, align 4 + %add10.i = add nsw i32 %add8.i, %19 + %20 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %7, i32 0) + %21 = load i32, ptr addrspace(11) %20, align 4 + %add12.i = add nsw i32 %add10.i, %21 + %22 = tail call noundef align 4 dereferenceable(4) ptr addrspace(11) @llvm.spv.resource.getpointer.p11.tspirv.SignedImage_i32_5_2_0_0_2_0t(target("spirv.SignedImage", i32, 5, 2, 0, 0, 2, 0) %0, i32 0) + store i32 %add12.i, ptr addrspace(11) %22, align 4 + ret void +} + + +attributes #0 = { "hlsl.numthreads"="1,1,1" "hlsl.shader"="compute" }
\ No newline at end of file diff --git a/llvm/test/CodeGen/SPIRV/layout.ll b/llvm/test/CodeGen/SPIRV/layout.ll new file mode 100644 index 0000000..94fa432 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/layout.ll @@ -0,0 +1,84 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpCapability Kernel +; CHECK: OpCapability Addresses +; CHECK: OpCapability GenericPointer +; CHECK: OpCapability Int64 +; CHECK: OpCapability Int8 +; CHECK: OpCapability Linkage + +; CHECK: OpExtInstImport "OpenCL.std" +; CHECK: OpMemoryModel Physical64 OpenCL +; CHECK: OpEntryPoint Kernel %[[#]] "foo" %[[#]] +; CHECK: OpSource OpenCL_C 200000 + +; CHECK-DAG: OpName %[[#]] +; CHECK-DAG: OpDecorate %[[#]] + + +; CHECK: %[[#I8:]] = OpTypeInt 8 0 +; CHECK: %[[#PTR_CW_I8:]] = OpTypePointer CrossWorkgroup %[[#I8]] +; CHECK: %[[#I32:]] = OpTypeInt 32 0 +; CHECK: %[[#VEC4:]] = OpTypeVector %[[#I32]] 4 +; CHECK: %[[#VOID:]] = OpTypeVoid +; CHECK: %[[#FUNC_TYPE0:]] = OpTypeFunction %[[#VOID]] %[[#PTR_CW_I8]] %[[#VEC4]] +; CHECK: %[[#FUNC_TYPE1:]] = OpTypeFunction %[[#VOID]] %[[#PTR_CW_I8]] +; CHECK: %[[#VEC3:]] = OpTypeVector %[[#I32]] 3 +; CHECK: %[[#FUNC_TYPE2:]] = OpTypeFunction %[[#VOID]] %[[#PTR_CW_I8]] %[[#VEC3]] +; CHECK: %[[#PTR_GEN_I8:]] = OpTypePointer Generic %[[#I8]] +; CHECK: %[[#STRUCT_B:]] = OpTypeStruct %[[#I32]] %[[#PTR_GEN_I8]] +; CHECK: %[[#STRUCT_C:]] = OpTypeStruct %[[#I32]] %[[#STRUCT_B]] +; CHECK: %[[#STRUCT_A:]] = OpTypeStruct %[[#I32]] %[[#STRUCT_C]] +; CHECK: %[[#F32:]] = OpTypeFloat 32 +; CHECK: %[[#CONST_2:]] = OpConstant %[[#I32]] 2 +; CHECK: %[[#ARRAY_F:]] = OpTypeArray %[[#F32]] %[[#CONST_2]] +; CHECK: %[[#ARRAY_I:]] = OpTypeArray %[[#I32]] %[[#CONST_2]] +; CHECK: %[[#PTR_CW_STRUCT_A:]] = OpTypePointer CrossWorkgroup %[[#STRUCT_A]] +; CHECK: %[[#PTR_UC_VEC4:]] = OpTypePointer UniformConstant %[[#VEC4]] +; CHECK: %[[#PTR_UC_ARRAY_F:]] = OpTypePointer UniformConstant %[[#ARRAY_F]] +; CHECK: %[[#PTR_CW_PTR_CW_I8:]] = OpTypePointer CrossWorkgroup %[[#PTR_CW_I8]] +; CHECK: %[[#I64:]] = OpTypeInt 64 0 +; CHECK: %[[#PTR_CW_ARRAY_I:]] = OpTypePointer CrossWorkgroup %[[#ARRAY_I]] + +; CHECK: %[[#NULL_I32:]] = OpConstantNull %[[#I32]] +; CHECK: %[[#CONST_I64_4:]] = OpConstant %[[#I64]] 4 +; CHECK: %[[#CONST_I32_1:]] = OpConstant %[[#I32]] 1 +; CHECK: %[[#COMP_I32:]] = OpConstantComposite %[[#ARRAY_I]] %[[#CONST_I32_1]] %[[#CONST_2]] + +; CHECK: %[[#VAR_V:]] = OpVariable %[[#PTR_CW_ARRAY_I]] CrossWorkgroup %[[#COMP_I32]] +; CHECK: %[[#SPECCONSTOP:]] = OpSpecConstantOp %[[#PTR_CW_I8]] InBoundsPtrAccessChain %[[#VAR_V]] %[[#NULL_I32]] %[[#CONST_I64_4]] +; CHECK: %[[#VAR_S:]] = OpVariable %[[#PTR_CW_PTR_CW_I8]] CrossWorkgroup %[[#SPECCONSTOP]] +; CHECK: %[[#NULL_ARRAY_F:]] = OpConstantNull %[[#ARRAY_F]] +; CHECK: %[[#VAR_F:]] = OpVariable %[[#PTR_UC_ARRAY_F]] UniformConstant %[[#NULL_ARRAY_F]] +; CHECK: %[[#NULL_STRUCT_A:]] = OpConstantNull %[[#STRUCT_A]] +; CHECK: %[[#VAR_A:]] = OpVariable %[[#PTR_CW_STRUCT_A]] CrossWorkgroup %[[#NULL_STRUCT_A]] + +; CHECK: %[[#FN_BAR1:]] = OpFunction %[[#VOID]] None %[[#FUNC_TYPE1]] +; CHECK: %[[#P_BAR1:]] = OpFunctionParameter %[[#PTR_CW_I8]] +; CHECK: OpFunctionEnd + +@v = addrspace(1) global [2 x i32] [i32 1, i32 2], align 4 +@s = addrspace(1) global ptr addrspace(1) getelementptr inbounds ([2 x i32], ptr addrspace(1) @v, i32 0, i32 1), align 4 + +%struct.A = type { i32, %struct.C } +%struct.C = type { i32, %struct.B } +%struct.B = type { i32, ptr addrspace(4) } + +@f = addrspace(2) constant [2 x float] zeroinitializer, align 4 +@b = external addrspace(2) constant <4 x i32> +@a = common addrspace(1) global %struct.A zeroinitializer, align 4 + +define spir_kernel void @foo(ptr addrspace(1) %a, <4 x i32> %vec_in) { +entry: + call spir_func void @bar1(ptr addrspace(1) %a) + %extractVec = shufflevector <4 x i32> %vec_in, <4 x i32> %vec_in, <3 x i32> <i32 0, i32 1, i32 2> + call spir_func void @bar2(ptr addrspace(1) %a, <3 x i32> %extractVec) + ret void +} + +declare spir_func void @bar1(ptr addrspace(1)) +declare spir_func void @bar2(ptr addrspace(1), <3 x i32>) + +!opencl.ocl.version = !{!7} +!7 = !{i32 2, i32 0} diff --git a/llvm/test/CodeGen/SPIRV/lit.local.cfg b/llvm/test/CodeGen/SPIRV/lit.local.cfg index f139d13..5179542 100644 --- a/llvm/test/CodeGen/SPIRV/lit.local.cfg +++ b/llvm/test/CodeGen/SPIRV/lit.local.cfg @@ -1,16 +1,6 @@ if not "SPIRV" in config.root.targets: config.unsupported = True -spirv_sim_root = os.path.join(config.llvm_src_root, "utils", "spirv-sim") - -config.substitutions.append( - ( - "spirv-sim", - "'%s' %s" % (config.python_executable, - os.path.join(spirv_sim_root, "spirv-sim.py")), - ) -) - if config.spirv_tools_tests: config.available_features.add("spirv-tools") config.substitutions.append(("spirv-dis", os.path.join(config.llvm_tools_dir, "spirv-dis"))) diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll new file mode 100644 index 0000000..438fff6 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll @@ -0,0 +1,92 @@ +;; Check that llvm.bitreverse.* intrinsics are lowered for +;; 2/4-bit scalar and vector types. + +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - | FileCheck %s +; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpCapability ArbitraryPrecisionIntegersINTEL +; CHECK: OpExtension "SPV_INTEL_arbitrary_precision_integers" + +; CHECK: %[[#I4:]] = OpTypeInt 4 0 +; CHECK: %[[#I2:]] = OpTypeInt 2 0 +; CHECK: %[[#Z4:]] = OpConstantNull %[[#I4]] +; CHECK: %[[#Z2:]] = OpConstantNull %[[#I2]] +; CHECK: %[[#V2I2:]] = OpTypeVector %[[#I2]] 2 +; CHECK: %[[#V2I4:]] = OpTypeVector %[[#I4]] 2 +; CHECK: %[[#V3I2:]] = OpTypeVector %[[#I2]] 3 +; CHECK: %[[#V3I4:]] = OpTypeVector %[[#I4]] 3 +; CHECK: %[[#V4I2:]] = OpTypeVector %[[#I2]] 4 +; CHECK: %[[#V4I4:]] = OpTypeVector %[[#I4]] 4 +; CHECK: %[[#V8I2:]] = OpTypeVector %[[#I2]] 8 +; CHECK: %[[#V8I4:]] = OpTypeVector %[[#I4]] 8 +; CHECK: %[[#V16I2:]] = OpTypeVector %[[#I2]] 16 +; CHECK: %[[#V16I4:]] = OpTypeVector %[[#I4]] 16 + + +; CHECK: %[[#]] = OpBitReverse %[[#I2]] %[[#Z2]] +; CHECK: %[[#]] = OpBitReverse %[[#I4]] %[[#Z4]] +; CHECK: %[[#]] = OpBitReverse %[[#V2I2]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V2I4]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V3I2]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V3I4]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V4I2]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V4I4]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V8I2]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V8I4]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V16I2]] %[[#]] +; CHECK: %[[#]] = OpBitReverse %[[#V16I4]] %[[#]] + +define spir_kernel void @testBitRev() { +entry: + %call2 = call i2 @llvm.bitreverse.i2(i2 0) + %call4 = call i4 @llvm.bitreverse.i4(i4 0) + ret void +} + +define spir_kernel void @testBitRevV2(<2 x i2> %a, <2 x i4> %b) { +entry: + %call2 = call <2 x i2> @llvm.bitreverse.v2i2(<2 x i2> %a) + %call4 = call <2 x i4> @llvm.bitreverse.v2i4(<2 x i4> %b) + ret void +} + +define spir_kernel void @testBitRevV3(<3 x i2> %a, <3 x i4> %b) { +entry: + %call2 = call <3 x i2> @llvm.bitreverse.v3i2(<3 x i2> %a) + %call4 = call <3 x i4> @llvm.bitreverse.v3i4(<3 x i4> %b) + ret void +} + +define spir_kernel void @testBitRevV4(<4 x i2> %a, <4 x i4> %b) { +entry: + %call2 = call <4 x i2> @llvm.bitreverse.v4i2(<4 x i2> %a) + %call4 = call <4 x i4> @llvm.bitreverse.v4i4(<4 x i4> %b) + ret void +} + +define spir_kernel void @testBitRevV8(<8 x i2> %a, <8 x i4> %b) { +entry: + %call2 = call <8 x i2> @llvm.bitreverse.v8i2(<8 x i2> %a) + %call4 = call <8 x i4> @llvm.bitreverse.v8i4(<8 x i4> %b) + ret void +} + +define spir_kernel void @testBitRevV16(<16 x i2> %a, <16 x i4> %b) { +entry: + %call2 = call <16 x i2> @llvm.bitreverse.v16i2(<16 x i2> %a) + %call4 = call <16 x i4> @llvm.bitreverse.v16i4(<16 x i4> %b) + ret void +} + +declare i2 @llvm.bitreverse.i2(i2) +declare i4 @llvm.bitreverse.i4(i4) +declare <2 x i2> @llvm.bitreverse.v2i2(<2 x i2>) +declare <2 x i4> @llvm.bitreverse.v2i4(<2 x i4>) +declare <3 x i2> @llvm.bitreverse.v3i2(<3 x i2>) +declare <3 x i4> @llvm.bitreverse.v3i4(<3 x i4>) +declare <4 x i2> @llvm.bitreverse.v4i2(<4 x i2>) +declare <4 x i4> @llvm.bitreverse.v4i4(<4 x i4>) +declare <8 x i2> @llvm.bitreverse.v8i2(<8 x i2>) +declare <8 x i4> @llvm.bitreverse.v8i4(<8 x i4>) +declare <16 x i2> @llvm.bitreverse.v16i2(<16 x i2>) +declare <16 x i4> @llvm.bitreverse.v16i4(<16 x i4>) diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fp-to-int-intrinsics.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fp-to-int-intrinsics.ll new file mode 100644 index 0000000..66c744f --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fp-to-int-intrinsics.ll @@ -0,0 +1,196 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unkown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unkown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpDecorate %[[#SAT1:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT2:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT3:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT4:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT5:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT6:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT7:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT8:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT9:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT10:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT11:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT12:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT13:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT14:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT15:]] SaturatedConversion +; CHECK: OpDecorate %[[#SAT16:]] SaturatedConversion + + +; CHECK: %[[#SAT1]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_signed_i8(float %input) { +entry: + %ptr = alloca i8 + %signed_int = call i8 @llvm.fptosi.sat.i8.f32(float %input) + store i8 %signed_int, i8* %ptr + ret void + +} +declare i8 @llvm.fptosi.sat.i8.f32(float) + + +; CHECK: %[[#SAT2]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_signed_i16(float %input) { +entry: + %ptr = alloca i16 + %signed_int = call i16 @llvm.fptosi.sat.i16.f32(float %input) + store i16 %signed_int, i16* %ptr + ret void + +} +declare i16 @llvm.fptosi.sat.i16.f32(float) + +; CHECK: %[[#SAT3]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_signed_i32(float %input) { +entry: + %ptr = alloca i32 + %signed_int = call i32 @llvm.fptosi.sat.i32.f32(float %input) + store i32 %signed_int, i32* %ptr + ret void + +} +declare i32 @llvm.fptosi.sat.i32.f32(float) + + +; CHECK: %[[#SAT4]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_signed_i64(float %input) { +entry: + %ptr = alloca i64 + %signed_int = call i64 @llvm.fptosi.sat.i64.f32(float %input) + store i64 %signed_int, i64* %ptr + ret void +} +declare i64 @llvm.fptosi.sat.i64.f32(float) + + +; CHECK: %[[#SAT5]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_signed_i8(double %input) { +entry: + %ptr = alloca i8 + %signed_int = call i8 @llvm.fptosi.sat.i8.f64(double %input) + store i8 %signed_int, i8* %ptr + ret void +} +declare i8 @llvm.fptosi.sat.i8.f64(double) + + +; CHECK: %[[#SAT6]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_signed_i16(double %input) { +entry: + %ptr = alloca i16 + %signed_int = call i16 @llvm.fptosi.sat.i16.f64(double %input) + store i16 %signed_int, i16* %ptr + ret void +} +declare i16 @llvm.fptosi.sat.i16.f64(double) + + +; CHECK: %[[#SAT7]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_signed_i32(double %input) { +entry: + %ptr = alloca i32 + %signed_int = call i32 @llvm.fptosi.sat.i32.f64(double %input) + store i32 %signed_int, i32* %ptr + ret void +} +declare i32 @llvm.fptosi.sat.i32.f64(double) + + +; CHECK: %[[#SAT8]] = OpConvertFToS %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_signed_i64(double %input) { +entry: + %ptr = alloca i64 + %signed_int = call i64 @llvm.fptosi.sat.i64.f64(double %input) + store i64 %signed_int, i64* %ptr + ret void +} +declare i64 @llvm.fptosi.sat.i64.f64(double) + +; CHECK: %[[#SAT9]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_unsigned_i8(float %input) { +entry: + %ptr = alloca i8 + %unsigned_int = call i8 @llvm.fptoui.sat.i8.f32(float %input) + store i8 %unsigned_int, i8* %ptr + ret void +} +declare i8 @llvm.fptoui.sat.i8.f32(float) + + +; CHECK: %[[#SAT10]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_unsigned_i16(float %input) { +entry: + %ptr = alloca i16 + %unsigned_int = call i16 @llvm.fptoui.sat.i16.f32(float %input) + store i16 %unsigned_int, i16* %ptr + ret void +} +declare i16 @llvm.fptoui.sat.i16.f32(float) + + +; CHECK: %[[#SAT11]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_unsigned_i32(float %input) { +entry: + %ptr = alloca i32 + %unsigned_int = call i32 @llvm.fptoui.sat.i32.f32(float %input) + store i32 %unsigned_int, i32* %ptr + ret void +} +declare i32 @llvm.fptoui.sat.i32.f32(float) + + +; CHECK: %[[#SAT12]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_float_to_unsigned_i64(float %input) { +entry: + %ptr = alloca i64 + %unsigned_int = call i64 @llvm.fptoui.sat.i64.f32(float %input) + store i64 %unsigned_int, i64* %ptr + ret void +} +declare i64 @llvm.fptoui.sat.i64.f32(float) + + +; CHECK: %[[#SAT13]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_unsigned_i8(double %input) { +entry: + %ptr = alloca i8 + %unsigned_int = call i8 @llvm.fptoui.sat.i8.f64(double %input) + store i8 %unsigned_int, i8* %ptr + ret void +} +declare i8 @llvm.fptoui.sat.i8.f64(double) + + +; CHECK: %[[#SAT14]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_unsigned_i16(double %input) { +entry: + %ptr = alloca i16 + %unsigned_int = call i16 @llvm.fptoui.sat.i16.f64(double %input) + store i16 %unsigned_int, i16* %ptr + ret void +} +declare i16 @llvm.fptoui.sat.i16.f64(double) + + +; CHECK: %[[#SAT15]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_unsigned_i32(double %input) { +entry: + %ptr = alloca i32 + %unsigned_int = call i32 @llvm.fptoui.sat.i32.f64(double %input) + store i32 %unsigned_int, i32* %ptr + ret void +} +declare i32 @llvm.fptoui.sat.i32.f64(double) + + +; CHECK: %[[#SAT16]] = OpConvertFToU %[[#]] %[[#]] +define spir_kernel void @testfunction_double_to_unsigned_i64(double %input) { +entry: + %ptr = alloca i64 + %unsigned_int = call i64 @llvm.fptoui.sat.i64.f64(double %input) + store i64 %unsigned_int, i64* %ptr + ret void +} +declare i64 @llvm.fptoui.sat.i64.f64(double) diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll index 9d07b63..3d93eca 100644 --- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll @@ -17,70 +17,63 @@ ; CL: OpFunction ; CL: %[[#FooVar:]] = OpVariable ; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]] -; CL-NEXT: OpLifetimeStart %[[#Casted1]], 72 -; CL-NEXT: OpCopyMemorySized +; CL-NEXT: OpLifetimeStart %[[#Casted1]], 16 ; CL-NEXT: OpBitcast ; CL-NEXT: OpInBoundsPtrAccessChain ; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]] -; CL-NEXT: OpLifetimeStop %[[#Casted2]], 72 +; CL-NEXT: OpLifetimeStop %[[#Casted2]], 16 ; VK: OpFunction ; VK: %[[#FooVar:]] = OpVariable -; VK-NEXT: OpCopyMemorySized ; VK-NEXT: OpInBoundsAccessChain ; VK-NEXT: OpReturn define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) { %RoundedRangeKernel = alloca %tprange, align 8 - call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %RoundedRangeKernel) - call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false) + call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel) %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8 - call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel) + call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel) ret void } ; CL: OpFunction ; CL: %[[#BarVar:]] = OpVariable -; CL-NEXT: OpLifetimeStart %[[#BarVar]], 0 -; CL-NEXT: OpCopyMemorySized +; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]] +; CL-NEXT: OpLifetimeStart %[[#Casted1]], 16 ; CL-NEXT: OpBitcast ; CL-NEXT: OpInBoundsPtrAccessChain -; CL-NEXT: OpLifetimeStop %[[#BarVar]], 0 +; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]] +; CL-NEXT: OpLifetimeStop %[[#Casted2]], 16 ; VK: OpFunction ; VK: %[[#BarVar:]] = OpVariable -; VK-NEXT: OpCopyMemorySized ; VK-NEXT: OpInBoundsAccessChain ; VK-NEXT: OpReturn define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange) { %RoundedRangeKernel = alloca %tprange, align 8 - call void @llvm.lifetime.start.p0(i64 -1, ptr nonnull %RoundedRangeKernel) - call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false) + call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel) %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8 - call void @llvm.lifetime.end.p0(i64 -1, ptr nonnull %RoundedRangeKernel) + call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel) ret void } ; CL: OpFunction ; CL: %[[#TestVar:]] = OpVariable ; CL-NEXT: OpLifetimeStart %[[#TestVar]], 1 -; CL-NEXT: OpCopyMemorySized ; CL-NEXT: OpInBoundsPtrAccessChain ; CL-NEXT: OpLifetimeStop %[[#TestVar]], 1 ; VK: OpFunction ; VK: %[[#Test:]] = OpVariable -; VK-NEXT: OpCopyMemorySized ; VK-NEXT: OpInBoundsAccessChain ; VK-NEXT: OpReturn define spir_func void @test(ptr noundef align 8 %_arg) { %var = alloca i8, align 8 - call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %var) - call void @llvm.memcpy.p0.p0.i64(ptr align 8 %var, ptr align 8 %_arg, i64 1, i1 false) - %KernelFunc = getelementptr inbounds i8, ptr %var, i64 0 - call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %var) + call void @llvm.lifetime.start.p0(ptr nonnull %var) + %KernelFunc = getelementptr inbounds i8, ptr %var, i64 1 + call void @llvm.lifetime.end.p0(ptr nonnull %var) ret void } -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.start.p0(ptr nocapture) declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.end.p0(ptr nocapture) diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memcpy.align.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memcpy.align.ll new file mode 100644 index 0000000..66a12b1 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memcpy.align.ll @@ -0,0 +1,54 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +%struct.B = type { [2 x i32] } +%struct.A = type { i64, %struct.B } + +@__const.foo.b = private unnamed_addr addrspace(2) constant %struct.B { [2 x i32] [i32 1, i32 2] }, align 4 +@__const.bar.a = private unnamed_addr addrspace(2) constant %struct.A { i64 0, %struct.B { [2 x i32] [i32 1, i32 2] } }, align 8 + +define spir_func void @foo(%struct.A* noalias sret(%struct.A) %agg.result) { +entry: + %b = alloca %struct.B, align 4 + %0 = bitcast %struct.B* %b to i8* + call void @llvm.lifetime.start.p0i8(i64 8, i8* %0) + %1 = bitcast %struct.B* %b to i8* + call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %1, i8 addrspace(2)* align 4 bitcast (%struct.B addrspace(2)* @__const.foo.b to i8 addrspace(2)*), i32 8, i1 false) +; CHECK: OpCopyMemorySized %[[#]] %[[#]] %[[#]] Aligned 4 + %b1 = getelementptr inbounds %struct.A, %struct.A* %agg.result, i32 0, i32 1 + %2 = bitcast %struct.B* %b1 to i8* + %3 = bitcast %struct.B* %b to i8* + call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 8 %2, i8* align 4 %3, i32 8, i1 false) +; CHECK: %[[#PTR1:]] = OpInBoundsPtrAccessChain %[[#]] %[[#]] %[[#]] %[[#]] +; CHECK: OpCopyMemorySized %[[#PTR1]] %[[#]] %[[#]] Aligned 8 + %4 = bitcast %struct.B* %b to i8* + call void @llvm.lifetime.end.p0i8(i64 8, i8* %4) + ret void +} + +declare void @llvm.lifetime.start.p0i8(i64, i8* captures(none)) + +declare void @llvm.memcpy.p0i8.p2i8.i32(i8* captures(none) writeonly, i8 addrspace(2)* captures(none) readonly, i32, i1) + +declare void @llvm.memcpy.p0i8.p0i8.i32(i8* captures(none) writeonly, i8* captures(none) readonly, i32, i1) + +declare void @llvm.lifetime.end.p0i8(i64, i8* captures(none)) + +define spir_func void @bar(%struct.B* noalias sret(%struct.B) %agg.result) { +entry: + %a = alloca %struct.A, align 8 + %0 = bitcast %struct.A* %a to i8* + call void @llvm.lifetime.start.p0i8(i64 16, i8* %0) + %1 = bitcast %struct.A* %a to i8* + call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 8 %1, i8 addrspace(2)* align 8 bitcast (%struct.A addrspace(2)* @__const.bar.a to i8 addrspace(2)*), i32 16, i1 false) +; CHECK: OpCopyMemorySized %[[#]] %[[#]] %[[#]] Aligned 8 + %b = getelementptr inbounds %struct.A, %struct.A* %a, i32 0, i32 1 + %2 = bitcast %struct.B* %agg.result to i8* + %3 = bitcast %struct.B* %b to i8* + call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 %2, i8* align 8 %3, i32 8, i1 false) +; CHECK: %[[#PTR2:]] = OpInBoundsPtrAccessChain %[[#]] %[[#]] %[[#]] %[[#]] +; CHECK: OpCopyMemorySized %[[#]] %[[#PTR2]] %[[#]] Aligned 4 + %4 = bitcast %struct.A* %a to i8* + call void @llvm.lifetime.end.p0i8(i64 16, i8* %4) + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/tan.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/tan.ll new file mode 100644 index 0000000..dfb185da --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/tan.ll @@ -0,0 +1,21 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-DAG: %[[#ext:]] = OpExtInstImport "OpenCL.std" +; CHECK-DAG: %[[#type_f32:]] = OpTypeFloat 32 +; CHECK-DAG: %[[#type_f64:]] = OpTypeFloat 64 +; CHECK: %[[#extinst_f32:]] = OpExtInst %[[#type_f32]] %[[#ext]] tan %[[#]] +; CHECK: %[[#extinst_f64:]] = OpExtInst %[[#type_f64]] %[[#ext]] tan %[[#]] + +define float @test_tan_f32(float %x) { + %res = call float @llvm.tan.f32(float %x) + ret float %res +} + +define double @test_tan_f64(double %x) { + %res = call double @llvm.tan.f64(double %x) + ret double %res +} + +declare float @llvm.tan.f32(float) +declare double @llvm.tan.f64(double) diff --git a/llvm/test/CodeGen/SPIRV/memory_model_md.ll b/llvm/test/CodeGen/SPIRV/memory_model_md.ll index e52343c..684a163 100644 --- a/llvm/test/CodeGen/SPIRV/memory_model_md.ll +++ b/llvm/test/CodeGen/SPIRV/memory_model_md.ll @@ -1,6 +1,7 @@ ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %} -; SPV: OpMemoryModel Physical32 Simple +; SPV: OpMemoryModel Physical32 OpenCL define dso_local dllexport void @k_no_fc(i32 %ibuf, i32 %obuf) local_unnamed_addr { entry: ret void @@ -8,4 +9,4 @@ entry: !spirv.MemoryModel = !{!0} -!0 = !{i32 1, i32 0} +!0 = !{i32 1, i32 2} diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll index 544c657..19451d2 100644 --- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll +++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll @@ -1,5 +1,5 @@ -; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s -; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %} @PrivInternal = internal addrspace(10) global i32 456 ; CHECK-DAG: %[[#type:]] = OpTypeInt 32 0 @@ -7,7 +7,7 @@ ; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456 ; CHECK-DAG: %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]] -define spir_kernel void @Foo() { +define hidden spir_func void @Foo() { %p = addrspacecast ptr addrspace(10) @PrivInternal to ptr %v = load i32, ptr %p, align 4 ret void @@ -15,3 +15,9 @@ define spir_kernel void @Foo() { ; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4 ; CHECK-Next: OpReturn } + +define void @main() #1 { + ret void +} + +attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll new file mode 100644 index 0000000..51db120 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll @@ -0,0 +1,21 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3%} + +; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0 + +; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456 +; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]] +; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]] +@PrivInternal = internal addrspace(10) global i32 456 + +define hidden spir_func void @Foo() { + %tmp = load i32, ptr addrspace(10) @PrivInternal + ret void +} + +define void @main() #1 { + ret void +} + +declare void @llvm.memcpy.p1.p2.i64(ptr addrspace(1) noalias nocapture writeonly, ptr addrspace(2) noalias nocapture readonly, i64, i1 immarg) +attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll index a1ded05..6914f4f 100644 --- a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll +++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll @@ -1,5 +1,5 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s -; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %} ; CHECK-DAG: %[[#U8:]] = OpTypeInt 8 0 ; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0 @@ -15,12 +15,7 @@ ; CHECK-DAG: %[[#INIT:]] = OpVariable %[[#VTYPE]] UniformConstant %[[#VAL]] @Init = private addrspace(2) constant i32 123 -; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456 -; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]] -; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]] -@PrivInternal = internal addrspace(10) global i32 456 - -define spir_kernel void @Foo() { +define internal spir_func void @Foo() { ; CHECK: %[[#]] = OpLoad %[[#]] %[[#PTR]] Aligned 8 %l = load ptr addrspace(1), ptr addrspace(1) @Ptr, align 8 ; CHECK: OpCopyMemorySized %[[#]] %[[#INIT]] %[[#]] Aligned 4 diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll index f396b5a..838c551 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll @@ -1,13 +1,19 @@ -; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %} ; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 ; CHECK-SPIRV: OpBitReverse %[[#int]] -define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr { +define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) { entry: %call = tail call i32 @llvm.bitreverse.i32(i32 %b) - store i32 %call, i32 addrspace(1)* %res, align 4 + store i32 %call, ptr %res, align 4 + ret void +} + +define void @main() #1 { ret void } declare i32 @llvm.bitreverse.i32(i32) +attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll deleted file mode 100644 index 8f04929..0000000 --- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll +++ /dev/null @@ -1,14 +0,0 @@ -; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV - -; CHECK-SPIRV: %[[#short:]] = OpTypeInt 16 -; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2 -; CHECK-SPIRV: OpBitReverse %[[#short2]] - -define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> addrspace(1)* nocapture %res) local_unnamed_addr { -entry: - %call = tail call <2 x i16> @llvm.bitreverse.v2i16(<2 x i16> %b) - store <2 x i16> %call, <2 x i16> addrspace(1)* %res, align 4 - ret void -} - -declare <2 x i16> @llvm.bitreverse.v2i16(<2 x i16>) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll new file mode 100644 index 0000000..3e2ed8b --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll @@ -0,0 +1,20 @@ +; RUN: llc -O0 -mtriple=spirv-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %} + +; CHECK-SPIRV: %[[#short:]] = OpTypeInt 32 +; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2 +; CHECK-SPIRV: OpBitReverse %[[#short2]] + +define hidden spir_func void @testBitRev(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c, ptr %res) { +entry: + %call = tail call <2 x i32> @llvm.bitreverse.v2i32(<2 x i32> %b) + store <2 x i32> %call, ptr %res, align 4 + ret void +} + +define void @main() #1 { + ret void +} + +declare <2 x i32> @llvm.bitreverse.v2i32(<2 x i32>) +attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" } diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/convert_functions.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/convert_functions.ll new file mode 100644 index 0000000..13a61b0 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/convert_functions.ll @@ -0,0 +1,56 @@ +; This test checks that functions with `convert_` prefix are translated as +; OpenCL builtins only in case they match the specification. Otherwise, we +; expect such functions to be translated to SPIR-V FunctionCall. + +; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-SPIRV: OpName %[[#Func:]] "_Z18convert_float_func" +; CHECK-SPIRV: OpName %[[#Func1:]] "_Z20convert_uint_satfunc" +; CHECK-SPIRV: OpName %[[#Func2:]] "_Z21convert_float_rtzfunc" +; CHECK-SPIRV-DAG: %[[#VoidTy:]] = OpTypeVoid +; CHECK-SPIRV-DAG: %[[#CharTy:]] = OpTypeInt 8 +; CHECK-SPIRV-DAG: %[[#FloatTy:]] = OpTypeFloat 32 + +; CHECK-SPIRV: %[[#Func]] = OpFunction %[[#VoidTy]] None %[[#]] +; CHECK-SPIRV: %[[#ConvertId1:]] = OpUConvert %[[#CharTy]] %[[#]] +; CHECK-SPIRV: %[[#ConvertId2:]] = OpConvertSToF %[[#FloatTy]] %[[#]] +; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#VoidTy]] %[[#Func]] %[[#ConvertId2]] +; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#VoidTy]] %[[#Func1]] %[[#]] +; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#VoidTy]] %[[#Func2]] %[[#ConvertId2]] +; CHECK-SPIRV-NOT: OpFConvert +; CHECK-SPIRV-NOT: OpConvertUToF + +define dso_local spir_func void @_Z18convert_float_func(float noundef %x) { +entry: + %x.addr = alloca float, align 4 + store float %x, ptr %x.addr, align 4 + ret void +} + +define dso_local spir_func void @_Z20convert_uint_satfunc(i32 noundef %x) { +entry: + ret void +} + +define dso_local spir_func void @_Z21convert_float_rtzfunc(float noundef %x) { +entry: + ret void +} + +define dso_local spir_func void @convert_int_bf16(i32 noundef %x) { +entry: + %x.addr = alloca i32, align 4 + store i32 %x, ptr %x.addr, align 4 + %0 = load i32, ptr %x.addr, align 4 + call spir_func signext i8 @_Z16convert_char_rtei(i32 noundef %0) + %call = call spir_func float @_Z13convert_floati(i32 noundef %0) + call spir_func void @_Z18convert_float_func(float noundef %call) + call spir_func void @_Z20convert_uint_satfunc(i32 noundef %0) + call spir_func void @_Z21convert_float_rtzfunc(float noundef %call) + ret void +} + +declare spir_func signext i8 @_Z16convert_char_rtei(i32 noundef) + +declare spir_func float @_Z13convert_floati(i32 noundef) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/nan.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/nan.ll new file mode 100644 index 0000000..1072f07 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/nan.ll @@ -0,0 +1,15 @@ +; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; Check OpenCL built-in nan translation. + +; CHECK-SPIRV: %[[#]] = OpExtInst %[[#]] %[[#]] nan %[[#]] + +define dso_local spir_kernel void @test(ptr addrspace(1) align 4 %a, i32 %b) { +entry: + %call = tail call spir_func float @_Z3nanj(i32 %b) + store float %call, ptr addrspace(1) %a, align 4 + ret void +} + +declare spir_func float @_Z3nanj(i32) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/shuffle.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/shuffle.ll new file mode 100644 index 0000000..aeca431 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/shuffle.ll @@ -0,0 +1,23 @@ +; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; Check OpenCL built-in shuffle and shuffle2 translation. + +; CHECK-SPIRV: %[[#]] = OpExtInst %[[#]] %[[#]] shuffle %[[#]] %[[#]] +; CHECK-SPIRV: %[[#]] = OpExtInst %[[#]] %[[#]] shuffle2 %[[#]] %[[#]] %[[#]] + +define spir_kernel void @test() { +entry: + %call = call spir_func <2 x float> @_Z7shuffleDv2_fDv2_j(<2 x float> zeroinitializer, <2 x i32> zeroinitializer) + ret void +} + +declare spir_func <2 x float> @_Z7shuffleDv2_fDv2_j(<2 x float>, <2 x i32>) + +define spir_kernel void @test2() { +entry: + %call = call spir_func <4 x float> @_Z8shuffle2Dv2_fS_Dv4_j(<2 x float> zeroinitializer, <2 x float> zeroinitializer, <4 x i32> zeroinitializer) + ret void +} + +declare spir_func <4 x float> @_Z8shuffle2Dv2_fS_Dv4_j(<2 x float>, <2 x float>, <4 x i32>) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/printf.ll b/llvm/test/CodeGen/SPIRV/transcoding/printf.ll new file mode 100644 index 0000000..338f0a5 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/printf.ll @@ -0,0 +1,14 @@ +; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +@.str = private unnamed_addr addrspace(2) constant [12 x i8] c"Hello World\00", align 1 + +; CHECK-SPIRV: %[[#]] = OpExtInst %[[#]] %[[#]] printf %[[#]] + +define dso_local spir_kernel void @BuiltinPrintf() { +entry: + %call = tail call i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) noundef @.str) + ret void +} + +declare noundef i32 @printf(ptr addrspace(2) nocapture noundef readonly, ...) diff --git a/llvm/test/CodeGen/SPIRV/zero-length-array.ll b/llvm/test/CodeGen/SPIRV/zero-length-array.ll new file mode 100644 index 0000000..666176c --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/zero-length-array.ll @@ -0,0 +1,11 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val %} + +; CHECK: %[[#type:]] = OpTypeInt 32 0 +; CHECK: %[[#ext:]] = OpConstant %[[#type]] 0 + +define spir_func void @_Z3foov() { +entry: + %i = alloca [0 x i32], align 4 + ret void +} |