From 27d825f608394bf461d13e506f8dd89ebe6cef96 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Fri, 18 Oct 2024 15:06:00 -0700 Subject: [PATCH 1/3] initial testing for OpCopyLogical --- test_conformance/spirv_new/CMakeLists.txt | 1 + test_conformance/spirv_new/spirvInfo.hpp | 41 ++++++++++ .../spv1.4/copylogical_struct.spvasm32 | 24 ++++++ .../spv1.4/copylogical_struct.spvasm64 | 24 ++++++ .../spv1.4/copyobject_struct.spvasm64 | 22 ++++++ test_conformance/spirv_new/test_spirv_14.cpp | 79 +++++++++++++++++++ 6 files changed, 191 insertions(+) create mode 100644 test_conformance/spirv_new/spirvInfo.hpp create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.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/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/copylogical_struct.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 new file mode 100644 index 0000000000..3076a06713 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 @@ -0,0 +1,24 @@ +; SPIR-V +; Version: 1.4 + OpCapability Addresses + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %kernel "copylogical_test" + %uint = OpTypeInt 32 0 + %float = OpTypeFloat 32 + %void = OpTypeVoid + %struct_a = OpTypeStruct %uint %float +%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a + %struct_b = OpTypeStruct %uint %float +%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b + %kernel_sig = OpTypeFunction %void %ptr_struct_b + %uint_1024 = OpConstant %uint 1024 + %float_pi = OpConstant %float 3.1415 +%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi + %kernel = OpFunction %void None %kernel_sig + %dst = OpFunctionParameter %ptr_struct_b + %entry = OpLabel +%struct_b_dst = OpCopyLogical %struct_b %struct_a_src + OpStore %dst %struct_b_dst + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 new file mode 100644 index 0000000000..2e6247a825 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 @@ -0,0 +1,24 @@ +; SPIR-V +; Version: 1.4 + OpCapability Addresses + OpCapability Kernel + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %kernel "copylogical_test" + %uint = OpTypeInt 32 0 + %float = OpTypeFloat 32 + %void = OpTypeVoid + %struct_a = OpTypeStruct %uint %float +%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a + %struct_b = OpTypeStruct %uint %float +%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b + %kernel_sig = OpTypeFunction %void %ptr_struct_b + %uint_1024 = OpConstant %uint 1024 + %float_pi = OpConstant %float 3.1415 +%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi + %kernel = OpFunction %void None %kernel_sig + %dst = OpFunctionParameter %ptr_struct_b + %entry = OpLabel +%struct_b_dst = OpCopyLogical %struct_b %struct_a_src + OpStore %dst %struct_b_dst + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 new file mode 100644 index 0000000000..f79ddf2620 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 @@ -0,0 +1,22 @@ +; SPIR-V +; Version: 1.4 + OpCapability Addresses + OpCapability Kernel + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %kernel "copylogical_test" + %uint = OpTypeInt 32 0 + %float = OpTypeFloat 32 + %void = OpTypeVoid + %struct_a = OpTypeStruct %uint %float +%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a + %kernel_sig = OpTypeFunction %void %ptr_struct_a + %uint_1024 = OpConstant %uint 1024 + %float_pi = OpConstant %float 3.1415 +%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi + %kernel = OpFunction %void None %kernel_sig + %dst = OpFunctionParameter %ptr_struct_a + %entry = OpLabel +%struct_a_dst = OpCopyObject %struct_a %struct_a_src + OpStore %dst %struct_a_dst + 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..922b8d1b77 --- /dev/null +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -0,0 +1,79 @@ +// +// 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 + +TEST_SPIRV_FUNC(spirv14_copylogical) +{ + 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; + } + + cl_int error = CL_SUCCESS; + + clProgramWrapper prog; +#if 0 + error = get_program_with_il(prog, deviceID, context, + "spv1.4/copylogical_struct"); +#else + // !!! TODO: Delete the copyobject file also, when this code is removed! + error = get_program_with_il(prog, deviceID, context, + "spv1.4/copyobject_struct"); +#endif + SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + + clKernelWrapper kernel = clCreateKernel(prog, "copylogical_test", &error); + SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + + struct TestStruct + { + cl_int i; + cl_float f; + }; + TestStruct results{ 0, 0.0f }; + + clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(results), NULL, &error); + SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + 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(results), + &results, 0, NULL, NULL); + SPIRV_CHECK_ERROR(error, "Unable to read destination buffer"); + + if (results.i != 1024 || results.f != 3.1415f) + { + log_error( + "Results mismatch with different pointers! Got: { %d, %f }\n", + results.i, results.f); + return TEST_FAIL; + } + + return TEST_PASS; +} From 8732c76b9ae37eb4d9b042ae871d9b966b7357c8 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Sat, 2 Nov 2024 16:53:08 -0700 Subject: [PATCH 2/3] switch to the kernel that uses OpCopyLogical --- .../spv1.4/copyobject_struct.spvasm64 | 22 ------------------- test_conformance/spirv_new/test_spirv_14.cpp | 14 +++++++----- 2 files changed, 8 insertions(+), 28 deletions(-) delete mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 deleted file mode 100644 index f79ddf2620..0000000000 --- a/test_conformance/spirv_new/spirv_asm/spv1.4/copyobject_struct.spvasm64 +++ /dev/null @@ -1,22 +0,0 @@ -; SPIR-V -; Version: 1.4 - OpCapability Addresses - OpCapability Kernel - OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %kernel "copylogical_test" - %uint = OpTypeInt 32 0 - %float = OpTypeFloat 32 - %void = OpTypeVoid - %struct_a = OpTypeStruct %uint %float -%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a - %kernel_sig = OpTypeFunction %void %ptr_struct_a - %uint_1024 = OpConstant %uint 1024 - %float_pi = OpConstant %float 3.1415 -%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi - %kernel = OpFunction %void None %kernel_sig - %dst = OpFunctionParameter %ptr_struct_a - %entry = OpLabel -%struct_a_dst = OpCopyObject %struct_a %struct_a_src - OpStore %dst %struct_a_dst - OpReturn - OpFunctionEnd diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 6c8efd6189..c53bd065de 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -360,19 +360,16 @@ TEST_SPIRV_FUNC(spirv14_copylogical) log_info("SPIR-V 1.4 not supported; skipping tests.\n"); return TEST_SKIPPED_ITSELF; } + cl_int error = CL_SUCCESS; clProgramWrapper prog; -#if 0 error = get_program_with_il(prog, deviceID, context, "spv1.4/copylogical_struct"); -#else - // !!! TODO: Delete the copyobject file also, when this code is removed! - error = get_program_with_il(prog, deviceID, context, - "spv1.4/copyobject_struct"); -#endif SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + clKernelWrapper kernel = clCreateKernel(prog, "copylogical_test", &error); SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + struct TestStruct { cl_int i; @@ -382,15 +379,19 @@ TEST_SPIRV_FUNC(spirv14_copylogical) clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(results), NULL, &error); SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); 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(results), &results, 0, NULL, NULL); SPIRV_CHECK_ERROR(error, "Unable to read destination buffer"); + if (results.i != 1024 || results.f != 3.1415f) { log_error( @@ -398,5 +399,6 @@ TEST_SPIRV_FUNC(spirv14_copylogical) results.i, results.f); return TEST_FAIL; } + return TEST_PASS; } From a40669b8abc92ff6953913c861cdb4a8d6ab351c Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Mon, 11 Nov 2024 12:56:24 -0800 Subject: [PATCH 3/3] fix bad error message --- test_conformance/spirv_new/test_spirv_14.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 8907920225..091bd15b30 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -554,7 +554,7 @@ TEST_SPIRV_FUNC(spirv14_copylogical) if (results.i != 1024 || results.f != 3.1415f) { log_error( - "Results mismatch with different pointers! Got: { %d, %f }\n", + "Results mismatch! Got: { %d, %f }\n", results.i, results.f); return TEST_FAIL; }