diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt index 53d77ee5..d95d29aa 100644 --- a/test_conformance/extensions/CMakeLists.txt +++ b/test_conformance/extensions/CMakeLists.txt @@ -1,2 +1,3 @@ add_subdirectory( cl_ext_cxx_for_opencl ) +add_subdirectory( cl_khr_command_buffer ) add_subdirectory( cl_khr_dx9_media_sharing ) diff --git a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt new file mode 100644 index 00000000..ac259f6d --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt @@ -0,0 +1,8 @@ +set(MODULE_NAME CL_KHR_COMMAND_BUFFER) + +set(${MODULE_NAME}_SOURCES + main.cpp + basic_command_buffer.cpp +) + +include(../../CMakeCommon.txt) diff --git a/test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.cpp b/test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.cpp new file mode 100644 index 00000000..62a02d83 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.cpp @@ -0,0 +1,588 @@ +// +// 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 "command_buffer_test_base.h" +#include "procs.h" +#include "harness/typeWrappers.h" + +#include +#include +#include + +#define CHECK_VERIFICATION_ERROR(reference, result, index) \ + { \ + if (reference != result) \ + { \ + log_error("Expected %d was %d at index %u\n", reference, result, \ + index); \ + return TEST_FAIL; \ + } \ + } + +namespace { + +// Helper test fixture for constructing OpenCL objects used in testing +// a variety of simple command-buffer enqueue scenarios. +struct BasicCommandBufferTest : CommandBufferTestBase +{ + + BasicCommandBufferTest(cl_device_id device, cl_context context, + cl_command_queue queue) + : CommandBufferTestBase(device), context(context), queue(queue), + command_buffer(this), simultaneous_use(false), + out_of_order_support(false), num_elements(0) + {} + + virtual bool Skip() + { + cl_command_queue_properties required_properties; + cl_int error = clGetDeviceInfo( + device, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR, + sizeof(required_properties), &required_properties, NULL); + test_error(error, + "Unable to query " + "CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR"); + + cl_command_queue_properties queue_properties; + + error = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES, + sizeof(queue_properties), + &queue_properties, NULL); + test_error(error, "Unable to query CL_QUEUE_PROPERTIES"); + + // Skip if queue properties don't contain those required + return required_properties != (required_properties & queue_properties); + } + + virtual cl_int SetUp(int elements) + { + cl_int error = init_extension_functions(); + if (error != CL_SUCCESS) + { + return error; + } + + // Query if device supports simultaneous use + cl_device_command_buffer_capabilities_khr capabilities; + error = + clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR, + sizeof(capabilities), &capabilities, NULL); + test_error(error, + "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR"); + simultaneous_use = + capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR; + out_of_order_support = + capabilities & CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR; + + if (elements <= 0) + { + return CL_INVALID_VALUE; + } + num_elements = static_cast(elements); + + // Kernel performs a parallel copy from an input buffer to output buffer + // is created. + const char *kernel_str = + R"( + __kernel void copy(__global int* in, __global int* out) { + size_t id = get_global_id(0); + out[id] = in[id]; + })"; + + 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"); + + in_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, + sizeof(cl_int) * num_elements, nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + out_mem = + clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(cl_int) * num_elements, nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + kernel = clCreateKernel(program, "copy", &error); + test_error(error, "Failed to create copy kernel"); + + error = clSetKernelArg(kernel, 0, sizeof(in_mem), &in_mem); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem); + test_error(error, "clSetKernelArg failed"); + + if (simultaneous_use) + { + cl_command_buffer_properties_khr properties[3] = { + CL_COMMAND_BUFFER_FLAGS_KHR, + CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, 0 + }; + command_buffer = + clCreateCommandBufferKHR(1, &queue, properties, &error); + } + else + { + command_buffer = + clCreateCommandBufferKHR(1, &queue, nullptr, &error); + } + test_error(error, "clCreateCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + // Test body returning an OpenCL error code + virtual cl_int Run() = 0; + + +protected: + size_t data_size() const { return num_elements * sizeof(cl_int); } + + cl_context context; + cl_command_queue queue; + clCommandBufferWrapper command_buffer; + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper in_mem, out_mem; + size_t num_elements; + + // Device support query results + bool simultaneous_use; + bool out_of_order_support; +}; + +// Test enqueuing a command-buffer containing a single NDRange command once +struct BasicEnqueueTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int 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"); + + const cl_int pattern = 42; + error = clEnqueueFillBuffer(queue, in_mem, &pattern, 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"); + + std::vector output_data(num_elements); + error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern, output_data[i], i); + } + + return CL_SUCCESS; + } +}; + +// Test enqueuing a command-buffer containing multiple command, including +// operations other than NDRange kernel execution. +struct MixedCommandsTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int error; + const size_t iterations = 4; + clMemWrapper result_mem = + clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_int) * iterations, nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + const cl_int pattern_base = 42; + for (size_t i = 0; i < iterations; i++) + { + const cl_int pattern = pattern_base + i; + cl_int error = clCommandFillBufferKHR( + command_buffer, nullptr, in_mem, &pattern, sizeof(cl_int), 0, + data_size(), 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandFillBufferKHR failed"); + + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &num_elements, nullptr, 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + const size_t result_offset = i * sizeof(cl_int); + error = clCommandCopyBufferKHR( + command_buffer, nullptr, out_mem, result_mem, 0, result_offset, + sizeof(cl_int), 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandCopyBufferKHR failed"); + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + std::vector result_data(num_elements); + error = clEnqueueReadBuffer(queue, result_mem, CL_TRUE, 0, + iterations * sizeof(cl_int), + result_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + for (size_t i = 0; i < iterations; i++) + { + const cl_int ref = pattern_base + i; + CHECK_VERIFICATION_ERROR(ref, result_data[i], i); + } + + return CL_SUCCESS; + } +}; + +// Test enqueueing a command-buffer blocked on a user-event +struct UserEventTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int 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"); + + clEventWrapper user_event = clCreateUserEvent(context, &error); + test_error(error, "clCreateUserEvent failed"); + + const cl_int pattern = 42; + error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0, + data_size(), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1, + &user_event, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + std::vector output_data(num_elements); + error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + error = clSetUserEventStatus(user_event, CL_COMPLETE); + test_error(error, "clSetUserEventStatus failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern, output_data[i], i); + } + + return CL_SUCCESS; + } +}; + +// Test flushing the command-queue between command-buffer enqueues +struct ExplicitFlushTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int 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"); + + const cl_int pattern_A = 42; + error = clEnqueueFillBuffer(queue, in_mem, &pattern_A, sizeof(cl_int), + 0, data_size(), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + error = clFlush(queue); + test_error(error, "clFlush failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + std::vector output_data_A(num_elements); + error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(), + output_data_A.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + const cl_int pattern_B = 0xA; + error = clEnqueueFillBuffer(queue, in_mem, &pattern_B, sizeof(cl_int), + 0, data_size(), 0, nullptr, nullptr); + test_error(error, "clEnqueueFillBuffer failed"); + + error = clFlush(queue); + test_error(error, "clFlush failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clFlush(queue); + test_error(error, "clFlush failed"); + + std::vector output_data_B(num_elements); + error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(), + output_data_B.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern_A, output_data_A[i], i); + + CHECK_VERIFICATION_ERROR(pattern_B, output_data_B[i], i); + } + return CL_SUCCESS; + } + + bool Skip() override + { + return !simultaneous_use || BasicCommandBufferTest::Skip(); + } +}; + +// Test enqueueing a command-buffer twice separated by another enqueue operation +struct InterleavedEnqueueTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int 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"); + + cl_int pattern = 42; + error = clEnqueueFillBuffer(queue, in_mem, &pattern, 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"); + + pattern = 0xABCD; + error = clEnqueueFillBuffer(queue, in_mem, &pattern, 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 = clEnqueueCopyBuffer(queue, in_mem, out_mem, 0, 0, data_size(), + 0, nullptr, nullptr); + test_error(error, "clEnqueueCopyBuffer failed"); + + std::vector output_data(num_elements); + error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern, output_data[i], i); + } + + return CL_SUCCESS; + } + + bool Skip() override + { + return !simultaneous_use || BasicCommandBufferTest::Skip(); + } +}; + +// Test sync-points with an out-of-order command-buffer +struct OutOfOrderTest : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + OutOfOrderTest(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue), + out_of_order_command_buffer(this), out_of_order_queue(nullptr), + event(nullptr) + {} + + cl_int Run() override + { + cl_sync_point_khr sync_points[2]; + + const cl_int pattern = 42; + cl_int error = + clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, in_mem, + &pattern, sizeof(cl_int), 0, data_size(), 0, + nullptr, &sync_points[0], nullptr); + test_error(error, "clCommandFillBufferKHR failed"); + + const cl_int overwritten_pattern = 0xACDC; + error = clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, + out_mem, &overwritten_pattern, + sizeof(cl_int), 0, data_size(), 0, + nullptr, &sync_points[1], nullptr); + test_error(error, "clCommandFillBufferKHR failed"); + + error = clCommandNDRangeKernelKHR( + out_of_order_command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &num_elements, nullptr, 2, sync_points, nullptr, nullptr); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clFinalizeCommandBufferKHR(out_of_order_command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + error = clEnqueueCommandBufferKHR( + 0, nullptr, out_of_order_command_buffer, 0, nullptr, &event); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + std::vector output_data(num_elements); + error = clEnqueueReadBuffer(out_of_order_queue, out_mem, CL_TRUE, 0, + data_size(), output_data.data(), 1, &event, + nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern, output_data[i], i); + } + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + if (!out_of_order_support) + { + // Test will skip as device doesn't support out-of-order + // command-buffers + return CL_SUCCESS; + } + + out_of_order_queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error); + test_error(error, "Unable to create command queue to test with"); + + out_of_order_command_buffer = + clCreateCommandBufferKHR(1, &out_of_order_queue, nullptr, &error); + test_error(error, "clCreateCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + bool Skip() override + { + return !out_of_order_support || BasicCommandBufferTest::Skip(); + } + + clCommandQueueWrapper out_of_order_queue; + clCommandBufferWrapper out_of_order_command_buffer; + clEventWrapper event; +}; + +#undef CHECK_VERIFICATION_ERROR + +template +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + CHECK_COMMAND_BUFFER_EXTENSION_AVAILABLE(device); + + auto test_fixture = T(device, context, queue); + cl_int error = test_fixture.SetUp(num_elements); + test_error_ret(error, "Error in test initialization", TEST_FAIL); + + if (test_fixture.Skip()) + { + return TEST_SKIPPED_ITSELF; + } + + error = test_fixture.Run(); + test_error_ret(error, "Test Failed", TEST_FAIL); + + return TEST_PASS; +} +} // anonymous namespace + +int test_single_ndrange(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_interleaved_enqueue(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_mixed_commands(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_explicit_flush(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_user_events(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_out_of_order(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/command_buffer_test_base.h b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_test_base.h new file mode 100644 index 00000000..0fd2e4ec --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_test_base.h @@ -0,0 +1,177 @@ +// +// 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. + +#ifndef _CL_KHR_COMMAND_BUFFER_TEST_BASE_H +#define _CL_KHR_COMMAND_BUFFER_TEST_BASE_H + +#include +#include "harness/deviceInfo.h" +#include "harness/testHarness.h" + + +// Base class for setting function pointers to new extension entry points +struct CommandBufferTestBase +{ + CommandBufferTestBase(cl_device_id device): device(device) {} + + cl_int init_extension_functions() + { + cl_platform_id platform; + cl_int error = + clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), + &platform, nullptr); + test_error(error, "clGetDeviceInfo for CL_DEVICE_PLATFORM failed"); + + // If it is supported get the addresses of all the APIs here. +#define GET_EXTENSION_ADDRESS(FUNC) \ + FUNC = reinterpret_cast( \ + clGetExtensionFunctionAddressForPlatform(platform, #FUNC)); \ + if (FUNC == nullptr) \ + { \ + log_error("ERROR: clGetExtensionFunctionAddressForPlatform failed" \ + " with " #FUNC "\n"); \ + return TEST_FAIL; \ + } + + GET_EXTENSION_ADDRESS(clCreateCommandBufferKHR); + GET_EXTENSION_ADDRESS(clReleaseCommandBufferKHR); + GET_EXTENSION_ADDRESS(clRetainCommandBufferKHR); + GET_EXTENSION_ADDRESS(clFinalizeCommandBufferKHR); + GET_EXTENSION_ADDRESS(clEnqueueCommandBufferKHR); + GET_EXTENSION_ADDRESS(clCommandBarrierWithWaitListKHR); + GET_EXTENSION_ADDRESS(clCommandCopyBufferKHR); + GET_EXTENSION_ADDRESS(clCommandCopyBufferRectKHR); + GET_EXTENSION_ADDRESS(clCommandCopyBufferToImageKHR); + GET_EXTENSION_ADDRESS(clCommandCopyImageKHR); + GET_EXTENSION_ADDRESS(clCommandCopyImageToBufferKHR); + GET_EXTENSION_ADDRESS(clCommandFillBufferKHR); + GET_EXTENSION_ADDRESS(clCommandFillImageKHR); + GET_EXTENSION_ADDRESS(clCommandNDRangeKernelKHR); + GET_EXTENSION_ADDRESS(clGetCommandBufferInfoKHR); +#undef GET_EXTENSION_ADDRESS + return CL_SUCCESS; + } + + clCreateCommandBufferKHR_fn clCreateCommandBufferKHR = nullptr; + clReleaseCommandBufferKHR_fn clReleaseCommandBufferKHR = nullptr; + clRetainCommandBufferKHR_fn clRetainCommandBufferKHR = nullptr; + clFinalizeCommandBufferKHR_fn clFinalizeCommandBufferKHR = nullptr; + clEnqueueCommandBufferKHR_fn clEnqueueCommandBufferKHR = nullptr; + clCommandBarrierWithWaitListKHR_fn clCommandBarrierWithWaitListKHR = + nullptr; + clCommandCopyBufferKHR_fn clCommandCopyBufferKHR = nullptr; + clCommandCopyBufferRectKHR_fn clCommandCopyBufferRectKHR = nullptr; + clCommandCopyBufferToImageKHR_fn clCommandCopyBufferToImageKHR = nullptr; + clCommandCopyImageKHR_fn clCommandCopyImageKHR = nullptr; + clCommandCopyImageToBufferKHR_fn clCommandCopyImageToBufferKHR = nullptr; + clCommandFillBufferKHR_fn clCommandFillBufferKHR = nullptr; + clCommandFillImageKHR_fn clCommandFillImageKHR = nullptr; + clCommandNDRangeKernelKHR_fn clCommandNDRangeKernelKHR = nullptr; + clGetCommandBufferInfoKHR_fn clGetCommandBufferInfoKHR = nullptr; + + cl_device_id device = nullptr; +}; + +// Wrapper class based off generic typeWrappers.h wrappers. However, because +// the release/retain functions are queried at runtime from the platform, +// rather than known at compile time we cannot link the instantiated template. +// Instead, pass an instance of `CommandBufferTestBase` on wrapper construction +// to access the release/retain functions. +class clCommandBufferWrapper { + cl_command_buffer_khr object = nullptr; + + void retain() + { + if (!object) return; + + auto err = base->clRetainCommandBufferKHR(object); + if (err != CL_SUCCESS) + { + print_error(err, "clRetainCommandBufferKHR() failed"); + std::abort(); + } + } + + void release() + { + if (!object) return; + + auto err = base->clReleaseCommandBufferKHR(object); + if (err != CL_SUCCESS) + { + print_error(err, "clReleaseCommandBufferKHR() failed"); + std::abort(); + } + } + + // Used to access release/retain functions + CommandBufferTestBase *base; + +public: + // We always want to have base available to dereference + clCommandBufferWrapper() = delete; + + clCommandBufferWrapper(CommandBufferTestBase *base): base(base) {} + + // On assignment, assume the object has a refcount of one. + clCommandBufferWrapper &operator=(cl_command_buffer_khr rhs) + { + reset(rhs); + return *this; + } + + // Copy semantics, increase retain count. + clCommandBufferWrapper(clCommandBufferWrapper const &w) { *this = w; } + clCommandBufferWrapper &operator=(clCommandBufferWrapper const &w) + { + reset(w.object); + retain(); + return *this; + } + + // Move semantics, directly take ownership. + clCommandBufferWrapper(clCommandBufferWrapper &&w) { *this = std::move(w); } + clCommandBufferWrapper &operator=(clCommandBufferWrapper &&w) + { + reset(w.object); + w.object = nullptr; + return *this; + } + + ~clCommandBufferWrapper() { reset(); } + + // Release the existing object, if any, and own the new one, if any. + void reset(cl_command_buffer_khr new_object = nullptr) + { + release(); + object = new_object; + } + + operator cl_command_buffer_khr() const { return object; } +}; + +#define CHECK_COMMAND_BUFFER_EXTENSION_AVAILABLE(device) \ + { \ + if (!is_extension_available(device, "cl_khr_command_buffer")) \ + { \ + log_info( \ + "Device does not support 'cl_khr_command_buffer'. Skipping " \ + "the test.\n"); \ + return TEST_SKIPPED_ITSELF; \ + } \ + } + + +#endif // _CL_KHR_COMMAND_BUFFER_TEST_BASE_H diff --git a/test_conformance/extensions/cl_khr_command_buffer/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/main.cpp new file mode 100644 index 00000000..4dece455 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/main.cpp @@ -0,0 +1,35 @@ +// 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 "procs.h" +#include "harness/testHarness.h" + +test_definition test_list[] = { + ADD_TEST(single_ndrange), ADD_TEST(interleaved_enqueue), + ADD_TEST(mixed_commands), ADD_TEST(explicit_flush), + ADD_TEST(user_events), ADD_TEST(out_of_order) +}; + + +int main(int argc, const char *argv[]) +{ + // A device may report the required properties of a queue that + // is compatible with command-buffers via the query + // CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR. We account + // for this in the tests themselves, rather than here, where we have a + // device to query. + const cl_command_queue_properties queue_properties = 0; + return runTestHarnessWithCheck(argc, argv, ARRAY_SIZE(test_list), test_list, + false, queue_properties, nullptr); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/procs.h b/test_conformance/extensions/cl_khr_command_buffer/procs.h new file mode 100644 index 00000000..58fd228f --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/procs.h @@ -0,0 +1,35 @@ +// +// 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. +// +#ifndef _CL_KHR_COMMAND_BUFFER_PROCS_H +#define _CL_KHR_COMMAND_BUFFER_PROCS_H + +#include + +// Basic command-buffer tests +extern int test_single_ndrange(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_interleaved_enqueue(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_mixed_commands(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_explicit_flush(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_user_events(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); + +#endif /*_CL_KHR_COMMAND_BUFFER_PROCS_H*/