From 1b2708920931ae61f58f9ff2a134b8d05c99c2d7 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Mon, 6 Mar 2023 11:30:27 +0100 Subject: [PATCH] clSetKernelArg test cases for cl_khr_command_buffer extenstion (#1593) * Initial commit for clSetKernelArg test case (issue #1369, p.1.7): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Added simultaneous version of clSetKernelArg test, (#1369, p.1.7): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Added/corrected comments (#1369, p.1.7) * Corrected logic of conditions of command-buffer simultaneous use (#1369, p.1.7): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Added missing comment for consistency with other command-buffer tests * Cleanup related to code review (#1369, p.1.7): https://github.com/KhronosGroup/OpenCL-CTS/issues/1369 * Added clang format correction * Removed redundant code (issue #1369 clSetKernelArg) --- .../cl_khr_command_buffer/CMakeLists.txt | 1 + .../command_buffer_set_kernel_arg.cpp | 321 ++++++++++++++++++ .../extensions/cl_khr_command_buffer/main.cpp | 2 + .../extensions/cl_khr_command_buffer/procs.h | 5 + 4 files changed, 329 insertions(+) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/command_buffer_set_kernel_arg.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt index 5a4f82d5..98b9eb7f 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt @@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_COMMAND_BUFFER) set(${MODULE_NAME}_SOURCES main.cpp basic_command_buffer.cpp + command_buffer_set_kernel_arg.cpp command_buffer_event_sync.cpp command_buffer_out_of_order.cpp command_buffer_profiling.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/command_buffer_set_kernel_arg.cpp b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_set_kernel_arg.cpp new file mode 100644 index 00000000..ef938421 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_set_kernel_arg.cpp @@ -0,0 +1,321 @@ +// +// Copyright (c) 2022 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 "basic_command_buffer.h" +#include "procs.h" + +#include + +namespace { + +//////////////////////////////////////////////////////////////////////////////// +// clSetKernelArg tests for cl_khr_command_buffer which handles below cases: +// -test interactions of clSetKernelArg with command-buffers +// -test interactions of clSetKernelArg on a command-buffer pending execution + +template +struct CommandBufferSetKernelArg : public BasicCommandBufferTest +{ + CommandBufferSetKernelArg(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue), trigger_event(nullptr) + { + simultaneous_use_requested = simul_use; + if (simul_use) buffer_size_multiplier = 2; + } + + //-------------------------------------------------------------------------- + cl_int SetUpKernel() override + { + cl_int error = CL_SUCCESS; + + const char* kernel_str = + R"( + __kernel void copy(int in, __global int* out, __global int* offset) + { + size_t id = get_global_id(0); + size_t ind = offset[0] + id; + out[ind] = in; + })"; + + error = create_single_kernel_helper_create_program(context, &program, 1, + &kernel_str); + test_error(error, "Failed to create program with source"); + + error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); + test_error(error, "Failed to build program"); + + kernel = clCreateKernel(program, "copy", &error); + test_error(error, "Failed to create copy kernel"); + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int SetUpKernelArgs() override + { + cl_int error = CL_SUCCESS; + out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + num_elements * buffer_size_multiplier + * sizeof(cl_int), + nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + // create secondary output buffer to test kernel args substitution + out_mem_k2 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + num_elements * buffer_size_multiplier + * sizeof(cl_int), + nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + cl_int offset = 0; + off_mem = + clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(cl_int), &offset, &error); + test_error(error, "clCreateBuffer failed"); + + cl_int in_arg = pattern_pri; + error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 2, sizeof(off_mem), &off_mem); + test_error(error, "clSetKernelArg failed"); + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int Run() override + { + cl_int error = CL_SUCCESS; + if (simultaneous_use_requested) + { + // enqueue simultaneous command-buffers with clSetKernelArg calls + error = RunSimultaneous(); + test_error(error, "RunSimultaneous failed"); + } + else + { + // enqueue single command-buffer with clSetKernelArg calls + error = RunSingle(); + test_error(error, "RunSingle failed"); + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int RecordCommandBuffer() + { + cl_int error = CL_SUCCESS; + + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + // changing kernel args at this point should have no effect, + // test will verify if clSetKernelArg didn't affect the first command + cl_int in_arg = pattern_sec; + error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 1, sizeof(out_mem_k2), &out_mem_k2); + test_error(error, "clSetKernelArg failed"); + + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int RunSingle() + { + cl_int error = CL_SUCCESS; + std::vector output_data(num_elements); + + // record command buffer + error = RecordCommandBuffer(); + test_error(error, "RecordCommandBuffer failed"); + + const cl_int pattern_base = 0; + error = + clEnqueueFillBuffer(queue, out_mem, &pattern_base, sizeof(cl_int), + 0, data_size(), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + const cl_int pattern_base_k2 = 1; + error = clEnqueueFillBuffer(queue, out_mem_k2, &pattern_base_k2, + sizeof(cl_int), 0, data_size(), 0, nullptr, + nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + // verify the result - result buffer must contain initial pattern + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern_pri, output_data[i], i); + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + struct SimulPassData + { + cl_int pattern; + cl_int offset; + std::vector output_buffer; + }; + + //-------------------------------------------------------------------------- + cl_int RecordSimultaneousCommandBuffer() const + { + cl_int error = CL_SUCCESS; + + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int EnqueueSimultaneousPass(SimulPassData& pd) + { + cl_int error = clEnqueueFillBuffer( + queue, out_mem, &pd.pattern, sizeof(cl_int), + pd.offset * sizeof(cl_int), data_size(), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + error = clEnqueueFillBuffer(queue, off_mem, &pd.offset, sizeof(cl_int), + 0, sizeof(cl_int), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + if (!trigger_event) + { + trigger_event = clCreateUserEvent(context, &error); + test_error(error, "clCreateUserEvent failed"); + } + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1, + &trigger_event, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clEnqueueReadBuffer( + queue, out_mem, CL_FALSE, pd.offset * sizeof(cl_int), data_size(), + pd.output_buffer.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + cl_int RunSimultaneous() + { + cl_int error = CL_SUCCESS; + + // record command buffer with primary queue + error = RecordSimultaneousCommandBuffer(); + test_error(error, "RecordSimultaneousCommandBuffer failed"); + + std::vector simul_passes = { + { 0, 0, std::vector(num_elements) } + }; + + error = EnqueueSimultaneousPass(simul_passes.front()); + test_error(error, "EnqueueSimultaneousPass 1 failed"); + + // changing kernel args at this point should have no effect, + // test will verify if clSetKernelArg didn't affect command-buffer + cl_int in_arg = pattern_sec; + error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 1, sizeof(out_mem_k2), &out_mem_k2); + test_error(error, "clSetKernelArg failed"); + + if (simultaneous_use_support) + { + cl_int offset = static_cast(num_elements); + simul_passes.push_back( + { 1, offset, std::vector(num_elements) }); + + error = EnqueueSimultaneousPass(simul_passes.back()); + test_error(error, "EnqueueSimultaneousPass 2 failed"); + } + + error = clSetUserEventStatus(trigger_event, CL_COMPLETE); + test_error(error, "clSetUserEventStatus failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + // verify the result buffer + for (auto&& pass : simul_passes) + { + auto& res_data = pass.output_buffer; + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern_pri, res_data[i], i); + } + } + + return CL_SUCCESS; + } + + //-------------------------------------------------------------------------- + clEventWrapper trigger_event = nullptr; + + const cl_int pattern_pri = 2; + const cl_int pattern_sec = 3; + + clMemWrapper out_mem_k2 = nullptr; +}; + +} // anonymous namespace + +int test_basic_set_kernel_arg(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest>( + device, context, queue, num_elements); +} + +int test_pending_set_kernel_arg(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest>(device, context, + queue, num_elements); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/main.cpp index 909a1025..ef975889 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/main.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/main.cpp @@ -46,6 +46,8 @@ test_definition test_list[] = { ADD_TEST(copy_image_to_buffer), ADD_TEST(copy_buffer_rect), ADD_TEST(barrier_wait_list), + ADD_TEST(basic_set_kernel_arg), + ADD_TEST(pending_set_kernel_arg), ADD_TEST(event_info_command_type), ADD_TEST(event_info_command_queue), ADD_TEST(event_info_execution_status), diff --git a/test_conformance/extensions/cl_khr_command_buffer/procs.h b/test_conformance/extensions/cl_khr_command_buffer/procs.h index 24e3faae..f7322d93 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/procs.h +++ b/test_conformance/extensions/cl_khr_command_buffer/procs.h @@ -29,6 +29,11 @@ extern int test_explicit_flush(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); extern int test_out_of_order(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_basic_set_kernel_arg(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_pending_set_kernel_arg(cl_device_id device, cl_context context, + cl_command_queue queue, + int num_elements); extern int test_regular_wait_for_command_buffer(cl_device_id device, cl_context context, cl_command_queue queue,