From 6e305db0179ca47b793dc655d7f710f9857df5a3 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 5 Aug 2024 15:34:46 -0700 Subject: [PATCH 1/7] add testing for the SPIR-V 1.4 Image Operand additions --- test_conformance/spirv_new/CMakeLists.txt | 1 + .../spv1.4/image_operand_signextend.spvasm32 | 38 +++++ .../spv1.4/image_operand_signextend.spvasm64 | 40 ++++++ .../spv1.4/image_operand_zeroextend.spvasm32 | 38 +++++ .../spv1.4/image_operand_zeroextend.spvasm64 | 40 ++++++ test_conformance/spirv_new/test_spirv_14.cpp | 134 ++++++++++++++++++ 6 files changed, 291 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 create mode 100644 test_conformance/spirv_new/test_spirv_14.cpp 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/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..d4e642a47d --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 @@ -0,0 +1,38 @@ +; 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 + %1 = OpExtInstImport "OpenCL.std" + 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..a4b4e7b671 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 @@ -0,0 +1,40 @@ +; 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 + %1 = OpExtInstImport "OpenCL.std" + 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..3c48755541 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 @@ -0,0 +1,38 @@ +; 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 + %1 = OpExtInstImport "OpenCL.std" + 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..dbce0b2aad --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 @@ -0,0 +1,40 @@ +; 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 + %1 = OpExtInstImport "OpenCL.std" + 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/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp new file mode 100644 index 0000000000..5c4eb79d66 --- /dev/null +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -0,0 +1,134 @@ +// +// 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 "types.hpp" + +#include + +extern bool gVersionSkip; + +static int check_spirv_14_support(cl_device_id deviceID) +{ + const char* cVersionString = "SPIR-V 1.4"; + + std::string ilVersions = get_device_il_version_string(deviceID); + + if (gVersionSkip) + { + log_info(" Skipping version check for %s.\n", cVersionString); + } + else if (ilVersions.find(cVersionString) == std::string::npos) + { + log_info(" Version %s is not supported; skipping test.\n", + cVersionString); + return TEST_SKIPPED_ITSELF; + } + + return TEST_PASS; +} + +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"); + + std::vector h_imgdata({ 0x1, 0x80, 0xFF, 0x0 }); + clMemWrapper src = + clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &image_format, 1, 1, 0, h_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) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_image_operand_helper(deviceID, context, queue, true); +} + +TEST_SPIRV_FUNC(spirv14_image_operand_zeroextend) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_image_operand_helper(deviceID, context, queue, false); +} From d47103546f0f3ce9b36363513142a5b359309644 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 5 Aug 2024 16:10:18 -0700 Subject: [PATCH 2/7] add testing for the SPIR-V 1.4 Loop Control additions --- .../loop_control_iterationmultiple.spvasm32 | 44 ++++++++ .../loop_control_iterationmultiple.spvasm64 | 47 ++++++++ .../loop_control_maxiterations.spvasm32 | 44 ++++++++ .../loop_control_maxiterations.spvasm64 | 47 ++++++++ .../loop_control_miniterations.spvasm32 | 44 ++++++++ .../loop_control_miniterations.spvasm64 | 47 ++++++++ .../spv1.4/loop_control_partialcount.spvasm32 | 44 ++++++++ .../spv1.4/loop_control_partialcount.spvasm64 | 47 ++++++++ .../spv1.4/loop_control_peelcount.spvasm64 | 47 ++++++++ .../spv1.4/loop_control_peelocunt.spvasm32 | 44 ++++++++ test_conformance/spirv_new/test_spirv_14.cpp | 105 ++++++++++++++++++ 11 files changed, 560 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 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..443d09bde1 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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..39254bda1f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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..636af31055 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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..bab7bace5e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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..6c3de3f058 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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..680d39f7e2 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_miniterations.spvasm64 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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..da0154a687 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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..f4cf3f44d0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_partialcount.spvasm64 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 new file mode 100644 index 0000000000..e784a14238 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 new file mode 100644 index 0000000000..9e99414849 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 37 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 5c4eb79d66..debf75732d 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -132,3 +132,108 @@ TEST_SPIRV_FUNC(spirv14_image_operand_zeroextend) 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) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_loop_control_helper(deviceID, context, queue, + "loop_control_miniterations"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_maxiterations) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_loop_control_helper(deviceID, context, queue, + "loop_control_maxiterations"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_iterationmultiple) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_loop_control_helper(deviceID, context, queue, + "loop_control_iterationmultiple"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_peelcount) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_loop_control_helper(deviceID, context, queue, + "loop_control_peelcount"); +} + +TEST_SPIRV_FUNC(spirv14_loop_control_partialcount) +{ + int check = check_spirv_14_support(deviceID); + if (check != TEST_PASS) + { + return check; + } + + return test_loop_control_helper(deviceID, context, queue, + "loop_control_partialcount"); +} From 592a688d28d23f76b914e758a05dd561f3ed37bf Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 7 Aug 2024 14:16:37 -0700 Subject: [PATCH 3/7] factor out the function to test for SPIR-v 1.4 support --- test_conformance/spirv_new/spirvInfo.hpp | 41 ++++++++++++ test_conformance/spirv_new/test_spirv_14.cpp | 70 ++++++-------------- 2 files changed, 63 insertions(+), 48 deletions(-) create mode 100644 test_conformance/spirv_new/spirvInfo.hpp 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/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index debf75732d..438dcd6924 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -15,32 +15,13 @@ // #include "testBase.h" +#include "spirvInfo.hpp" #include "types.hpp" #include extern bool gVersionSkip; -static int check_spirv_14_support(cl_device_id deviceID) -{ - const char* cVersionString = "SPIR-V 1.4"; - - std::string ilVersions = get_device_il_version_string(deviceID); - - if (gVersionSkip) - { - log_info(" Skipping version check for %s.\n", cVersionString); - } - else if (ilVersions.find(cVersionString) == std::string::npos) - { - log_info(" Version %s is not supported; skipping test.\n", - cVersionString); - return TEST_SKIPPED_ITSELF; - } - - return TEST_PASS; -} - static int test_image_operand_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, bool signExtend) { @@ -113,23 +94,21 @@ static int test_image_operand_helper(cl_device_id deviceID, cl_context context, TEST_SPIRV_FUNC(spirv14_image_operand_signextend) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; } - return test_image_operand_helper(deviceID, context, queue, false); } @@ -180,60 +159,55 @@ static int test_loop_control_helper(cl_device_id deviceID, cl_context context, TEST_SPIRV_FUNC(spirv14_loop_control_miniterations) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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) { - int check = check_spirv_14_support(deviceID); - if (check != TEST_PASS) + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { - return check; + 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"); } From 6a1a214a6c18d033f273e1589893f23f338ffa8f Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 12 Aug 2024 07:27:21 -0700 Subject: [PATCH 4/7] fix filename --- ...control_peelocunt.spvasm32 => loop_control_peelcount.spvasm32} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename test_conformance/spirv_new/spirv_asm/spv1.4/{loop_control_peelocunt.spvasm32 => loop_control_peelcount.spvasm32} (100%) diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 similarity index 100% rename from test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelocunt.spvasm32 rename to test_conformance/spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 From c178374111a75f9b648b4fa61a436c5c44c3d8c9 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 12 Aug 2024 07:31:54 -0700 Subject: [PATCH 5/7] remove unneeded extended instruction set import --- .../spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm32 | 1 - .../spirv_new/spirv_asm/spv1.4/image_operand_signextend.spvasm64 | 1 - .../spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm32 | 1 - .../spirv_new/spirv_asm/spv1.4/image_operand_zeroextend.spvasm64 | 1 - .../spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm32 | 1 - .../spirv_asm/spv1.4/loop_control_iterationmultiple.spvasm64 | 1 - .../spirv_asm/spv1.4/loop_control_maxiterations.spvasm32 | 1 - .../spirv_asm/spv1.4/loop_control_maxiterations.spvasm64 | 1 - .../spirv_asm/spv1.4/loop_control_miniterations.spvasm32 | 1 - .../spirv_asm/spv1.4/loop_control_miniterations.spvasm64 | 1 - .../spirv_asm/spv1.4/loop_control_partialcount.spvasm32 | 1 - .../spirv_asm/spv1.4/loop_control_partialcount.spvasm64 | 1 - .../spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm32 | 1 - .../spirv_new/spirv_asm/spv1.4/loop_control_peelcount.spvasm64 | 1 - 14 files changed, 14 deletions(-) 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 index d4e642a47d..a00e66ebe3 100644 --- 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 @@ -8,7 +8,6 @@ OpCapability Kernel OpCapability ImageBasic OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %read_image_test "read_image_test" OpSource OpenCL_C 102000 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 index a4b4e7b671..316eb17a50 100644 --- 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 @@ -9,7 +9,6 @@ OpCapability Int64 OpCapability ImageBasic OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %read_image_test "read_image_test" OpSource OpenCL_C 102000 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 index 3c48755541..c763af6bf9 100644 --- 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 @@ -8,7 +8,6 @@ OpCapability Kernel OpCapability ImageBasic OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %read_image_test "read_image_test" OpSource OpenCL_C 102000 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 index dbce0b2aad..64910cdff8 100644 --- 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 @@ -9,7 +9,6 @@ OpCapability Int64 OpCapability ImageBasic OpCapability LiteralSampler - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %read_image_test "read_image_test" OpSource OpenCL_C 102000 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 index 443d09bde1..2755ee446a 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index 39254bda1f..80ffd20b48 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index 636af31055..c9220dce81 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index bab7bace5e..a9c4933f6e 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index 6c3de3f058..e2ec34d2a0 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index 680d39f7e2..a19336a66a 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index da0154a687..cbe8cceb29 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index f4cf3f44d0..51216f63e9 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index 9e99414849..ee3639ff6e 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 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 index e784a14238..2d94190a52 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %loop_control_test "loop_control_test" OpSource OpenCL_C 102000 From 57316d119f83f46ce7d5fd50e301ca97a05b4245 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 12 Aug 2024 07:33:25 -0700 Subject: [PATCH 6/7] remove unneeded extern --- test_conformance/spirv_new/test_spirv_14.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 438dcd6924..9aad1e7f42 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -20,8 +20,6 @@ #include -extern bool gVersionSkip; - static int test_image_operand_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, bool signExtend) { From 8b6da900c3c5d7cfd7c81538b04489b7b2c0cf1a Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 12 Aug 2024 07:35:48 -0700 Subject: [PATCH 7/7] address review comments --- test_conformance/spirv_new/test_spirv_14.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 9aad1e7f42..83a72d44d7 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -19,6 +19,7 @@ #include "types.hpp" #include +#include static int test_image_operand_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, bool signExtend) @@ -59,10 +60,9 @@ static int test_image_operand_helper(cl_device_id deviceID, cl_context context, h_dst.size() * sizeof(cl_uint), h_dst.data(), &error); SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); - std::vector h_imgdata({ 0x1, 0x80, 0xFF, 0x0 }); clMemWrapper src = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - &image_format, 1, 1, 0, h_imgdata.data(), &error); + &image_format, 1, 1, 0, imgData.data(), &error); SPIRV_CHECK_ERROR(error, "Failed to create src image"); error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);