diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 16a61b4075..a9aa0432e1 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/test_conformance/spirv_new/CMakeLists.txt @@ -27,6 +27,7 @@ set(${MODULE_NAME}_SOURCES test_op_vector_extract.cpp test_op_vector_insert.cpp test_op_vector_times_scalar.cpp + test_spirv_14.cpp ) set(TEST_HARNESS_SOURCES diff --git a/test_conformance/spirv_new/spirvInfo.hpp b/test_conformance/spirv_new/spirvInfo.hpp new file mode 100644 index 0000000000..ed4d6c794c --- /dev/null +++ b/test_conformance/spirv_new/spirvInfo.hpp @@ -0,0 +1,41 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#pragma once + +#include "harness/compat.h" + +#include + +extern bool gVersionSkip; + +static bool is_spirv_version_supported(cl_device_id deviceID, + const char* version) +{ + std::string ilVersions = get_device_il_version_string(deviceID); + + if (gVersionSkip) + { + log_info(" Skipping version check for %s.\n", version); + return true; + } + else if (ilVersions.find(version) == std::string::npos) + { + return false; + } + + return true; +} diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 new file mode 100644 index 0000000000..a00e66ebe3 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 31 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability ImageBasic + OpCapability LiteralSampler + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %read_image_test "read_image_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %index_0 = OpConstant %uint 0 + %void = OpTypeVoid + %v4uint = OpTypeVector %uint 4 +%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint + %image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly + %7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t + %sampler_t = OpTypeSampler +%sampledimage_t = OpTypeSampledImage %image2d_t + %v2uint = OpTypeVector %uint 2 + %float = OpTypeFloat 32 + %sampler = OpConstantSampler %sampler_t None 0 Nearest + %coord_0_0 = OpConstantNull %v2uint + %float_0 = OpConstant %float 0 +%read_image_test = OpFunction %void None %7 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint + %image = OpFunctionParameter %image2d_t + %entry = OpLabel +%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler + %call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|SignExtend %float_0 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0 + OpStore %arrayidx %call Aligned 16 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 new file mode 100644 index 0000000000..316eb17a50 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 31 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpCapability ImageBasic + OpCapability LiteralSampler + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %read_image_test "read_image_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %index_0 = OpConstant %ulong 0 + %void = OpTypeVoid + %v4uint = OpTypeVector %uint 4 +%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint + %image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly + %7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t + %sampler_t = OpTypeSampler +%sampledimage_t = OpTypeSampledImage %image2d_t + %v2uint = OpTypeVector %uint 2 + %float = OpTypeFloat 32 + %sampler = OpConstantSampler %sampler_t None 0 Nearest + %coord_0_0 = OpConstantNull %v2uint + %float_0 = OpConstant %float 0 +%read_image_test = OpFunction %void None %7 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint + %image = OpFunctionParameter %image2d_t + %entry = OpLabel +%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler + %call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|SignExtend %float_0 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0 + OpStore %arrayidx %call Aligned 16 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 new file mode 100644 index 0000000000..c763af6bf9 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 31 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability ImageBasic + OpCapability LiteralSampler + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %read_image_test "read_image_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %index_0 = OpConstant %uint 0 + %void = OpTypeVoid + %v4uint = OpTypeVector %uint 4 +%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint + %image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly + %7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t + %sampler_t = OpTypeSampler +%sampledimage_t = OpTypeSampledImage %image2d_t + %v2uint = OpTypeVector %uint 2 + %float = OpTypeFloat 32 + %sampler = OpConstantSampler %sampler_t None 0 Nearest + %coord_0_0 = OpConstantNull %v2uint + %float_0 = OpConstant %float 0 +%read_image_test = OpFunction %void None %7 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint + %image = OpFunctionParameter %image2d_t + %entry = OpLabel +%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler + %call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|ZeroExtend %float_0 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0 + OpStore %arrayidx %call Aligned 16 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 new file mode 100644 index 0000000000..64910cdff8 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 31 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpCapability ImageBasic + OpCapability LiteralSampler + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %read_image_test "read_image_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %index_0 = OpConstant %ulong 0 + %void = OpTypeVoid + %v4uint = OpTypeVector %uint 4 +%_ptr_CrossWorkgroup_v4uint = OpTypePointer CrossWorkgroup %v4uint + %image2d_t = OpTypeImage %void 2D 0 0 0 0 Unknown ReadOnly + %7 = OpTypeFunction %void %_ptr_CrossWorkgroup_v4uint %image2d_t + %sampler_t = OpTypeSampler +%sampledimage_t = OpTypeSampledImage %image2d_t + %v2uint = OpTypeVector %uint 2 + %float = OpTypeFloat 32 + %sampler = OpConstantSampler %sampler_t None 0 Nearest + %coord_0_0 = OpConstantNull %v2uint + %float_0 = OpConstant %float 0 +%read_image_test = OpFunction %void None %7 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_v4uint + %image = OpFunctionParameter %image2d_t + %entry = OpLabel +%TempSampledImage = OpSampledImage %sampledimage_t %image %sampler + %call = OpImageSampleExplicitLod %v4uint %TempSampledImage %coord_0_0 Lod|ZeroExtend %float_0 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_v4uint %dst %index_0 + OpStore %arrayidx %call Aligned 16 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 new file mode 100644 index 0000000000..2755ee446a --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc IterationMultiple 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 new file mode 100644 index 0000000000..80ffd20b48 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ulong_0 = OpConstant %ulong 0 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc IterationMultiple 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 new file mode 100644 index 0000000000..c9220dce81 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc MaxIterations 16 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 new file mode 100644 index 0000000000..a9c4933f6e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ulong_0 = OpConstant %ulong 0 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc MaxIterations 16 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 new file mode 100644 index 0000000000..e2ec34d2a0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc MinIterations 4 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 new file mode 100644 index 0000000000..a19336a66a --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ulong_0 = OpConstant %ulong 0 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc MinIterations 4 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 new file mode 100644 index 0000000000..cbe8cceb29 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc PartialCount 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 new file mode 100644 index 0000000000..51216f63e9 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ulong_0 = OpConstant %ulong 0 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc PartialCount 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 new file mode 100644 index 0000000000..ee3639ff6e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc PeelCount 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %uint_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 new file mode 100644 index 0000000000..2d94190a52 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %loop_control_test "loop_control_test" + OpSource OpenCL_C 102000 + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ulong_0 = OpConstant %ulong 0 + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %uint %uint + %bool = OpTypeBool +%loop_control_test = OpFunction %void None %5 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %count = OpFunctionParameter %uint + %value = OpFunctionParameter %uint + %entry = OpLabel + OpBranch %for_cond + %for_cond = OpLabel + %sum_0 = OpPhi %uint %uint_0 %entry %add %for_inc + %i_0 = OpPhi %uint %uint_0 %entry %inc %for_inc + %cmp = OpSLessThan %bool %i_0 %count + OpLoopMerge %for_cond_cleanup %for_inc PeelCount 2 + OpBranchConditional %cmp %for_body %for_cond_cleanup +%for_cond_cleanup = OpLabel + OpBranch %for_end + %for_body = OpLabel + %add = OpIAdd %uint %sum_0 %value + OpBranch %for_inc + %for_inc = OpLabel + %inc = OpIAdd %uint %i_0 %uint_1 + OpBranch %for_cond + %for_end = OpLabel + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %ulong_0 + OpStore %arrayidx %sum_0 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp new file mode 100644 index 0000000000..83a72d44d7 --- /dev/null +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -0,0 +1,211 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include "testBase.h" +#include "spirvInfo.hpp" +#include "types.hpp" + +#include +#include + +static int test_image_operand_helper(cl_device_id deviceID, cl_context context, + cl_command_queue queue, bool signExtend) +{ + const char* filename = signExtend ? "spv1.4/image_operand_signextend" + : "spv1.4/image_operand_zeroextend"; + cl_image_format image_format = { + CL_RGBA, + signExtend ? CL_SIGNED_INT8 : CL_UNSIGNED_INT8, + }; + + cl_int error = CL_SUCCESS; + + std::vector imgData({ 0x1, 0x80, 0xFF, 0x0 }); + std::vector expected; + for (auto v : imgData) + { + if (signExtend) + { + expected.push_back((cl_int)(cl_char)v); + } + else + { + expected.push_back(v); + } + } + + clProgramWrapper prog; + error = get_program_with_il(prog, deviceID, context, filename); + SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + + clKernelWrapper kernel = clCreateKernel(prog, "read_image_test", &error); + SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + + std::vector h_dst({ 0, 0, 0, 0 }); + clMemWrapper dst = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + h_dst.size() * sizeof(cl_uint), h_dst.data(), &error); + SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); + + clMemWrapper src = + clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &image_format, 1, 1, 0, imgData.data(), &error); + SPIRV_CHECK_ERROR(error, "Failed to create src image"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(src), &src); + SPIRV_CHECK_ERROR(error, "Failed to set kernel args"); + + size_t global = 1; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel"); + + error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, + h_dst.size() * sizeof(cl_uint), h_dst.data(), 0, + NULL, NULL); + SPIRV_CHECK_ERROR(error, "Unable to read destination buffer"); + + if (h_dst != expected) + { + log_error("Mismatch! Got: %u, %u, %u, %u, Wanted: %u, %u, %u, %u\n", + h_dst[0], h_dst[1], h_dst[2], h_dst[3], expected[0], + expected[1], expected[2], expected[3]); + return TEST_FAIL; + } + + return TEST_PASS; +} + +TEST_SPIRV_FUNC(spirv14_image_operand_signextend) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_image_operand_helper(deviceID, context, queue, true); +} + +TEST_SPIRV_FUNC(spirv14_image_operand_zeroextend) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_image_operand_helper(deviceID, context, queue, false); +} + +static int test_loop_control_helper(cl_device_id deviceID, cl_context context, + cl_command_queue queue, + const char* filename) +{ + const int count = 10; + const int value = 5; + + cl_int error = CL_SUCCESS; + + clProgramWrapper prog; + std::string full_filename = "spv1.4/" + std::string(filename); + error = get_program_with_il(prog, deviceID, context, full_filename.c_str()); + SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + + clKernelWrapper kernel = clCreateKernel(prog, "loop_control_test", &error); + SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + + int h_dst = 0; + clMemWrapper dst = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + sizeof(h_dst), &h_dst, &error); + SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(count), &count); + error |= clSetKernelArg(kernel, 2, sizeof(value), &value); + SPIRV_CHECK_ERROR(error, "Failed to set kernel args"); + + size_t global = 1; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel"); + + error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(h_dst), &h_dst, + 0, NULL, NULL); + SPIRV_CHECK_ERROR(error, "Unable to read destination buffer"); + + if (h_dst != count * value) + { + log_error("Mismatch! Got: %i, Wanted: %i\n", h_dst, count * value); + return TEST_FAIL; + } + + return TEST_PASS; +} + +TEST_SPIRV_FUNC(spirv14_loop_control_miniterations) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_loop_control_helper(deviceID, context, queue, + "loop_control_miniterations"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_maxiterations) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_loop_control_helper(deviceID, context, queue, + "loop_control_maxiterations"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_iterationmultiple) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_loop_control_helper(deviceID, context, queue, + "loop_control_iterationmultiple"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_peelcount) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_loop_control_helper(deviceID, context, queue, + "loop_control_peelcount"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_partialcount) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + return test_loop_control_helper(deviceID, context, queue, + "loop_control_partialcount"); +}