From 50ee8412b2b7c356755d9678f80a02ddbf368cc2 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Thu, 12 Dec 2024 17:23:02 +0200 Subject: [PATCH 1/9] add tests for cl_ext_buffer_device_address version 0.9.0 --- test_conformance/extensions/CMakeLists.txt | 1 + .../CMakeLists.txt | 7 + .../buffer_device_address.cpp | 413 ++++++++++++++++++ .../cl_ext_buffer_device_address/main.cpp | 28 ++ .../cl_ext_buffer_device_address/procs.h | 25 ++ 5 files changed, 474 insertions(+) create mode 100644 test_conformance/extensions/cl_ext_buffer_device_address/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp create mode 100644 test_conformance/extensions/cl_ext_buffer_device_address/main.cpp create mode 100644 test_conformance/extensions/cl_ext_buffer_device_address/procs.h diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt index d064e8a9bb..aa57990b38 100644 --- a/test_conformance/extensions/CMakeLists.txt +++ b/test_conformance/extensions/CMakeLists.txt @@ -8,6 +8,7 @@ add_subdirectory( cl_khr_dx9_media_sharing ) add_subdirectory( cl_khr_external_memory_dma_buf ) add_subdirectory( cl_khr_semaphore ) add_subdirectory( cl_khr_kernel_clock ) +add_subdirectory( cl_ext_buffer_device_address ) if(VULKAN_IS_SUPPORTED) add_subdirectory( cl_khr_external_semaphore ) endif() diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/CMakeLists.txt b/test_conformance/extensions/cl_ext_buffer_device_address/CMakeLists.txt new file mode 100644 index 0000000000..c4478693cf --- /dev/null +++ b/test_conformance/extensions/cl_ext_buffer_device_address/CMakeLists.txt @@ -0,0 +1,7 @@ +set(MODULE_NAME CL_EXT_BUFFER_DEVICE_ADDRESS) + +set(${MODULE_NAME}_SOURCES + main.cpp buffer_device_address.cpp +) + +include(../../CMakeCommon.txt) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp new file mode 100644 index 0000000000..1332062529 --- /dev/null +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -0,0 +1,413 @@ +// 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 "procs.h" +#include "harness/typeWrappers.h" + +#define BUF_SIZE 1024 +#define BUF_SIZE_STR "1024" + +namespace { + +static const char *program_source = + R"raw( + // A kernel that gets the device-seen address of the buffer. + __kernel void get_addr (__global int *buffer, + __global ulong* addr) { + for (int i = 0; i < BUF_SIZE; ++i) + buffer[i] += 1; + *addr = (ulong)buffer; + } + + // A kernel that accesses another buffer indirectly. + __kernel void indirect_access (__global long* in_addr, + __global int* out) { + *out = **(int __global* __global*)in_addr; + } + + // A kernel that gets passed a pointer to a middle of a buffer, + // with the data _before_ the passed pointer. Tests the property + // of sub-buffers to synchronize the whole parent buffer when + // using the CL_MEM_BUFFER_DEVICE_ADDRESS flag. + __kernel void ptr_arith (__global int* in_addr, + __global int* out) { + *out = *(in_addr - 2); + } + )raw"; + +class BufferDeviceAddressTest { + +public: + BufferDeviceAddressTest(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_mem_properties address_type) + : device(device), context(context), queue(queue), + address_type(address_type) + {} + + bool Skip() + { + cl_int error = 0; + clMemWrapper TempBuffer = clCreateBuffer( + context, (cl_mem_flags)(CL_MEM_READ_WRITE | address_type), + (size_t)BUF_SIZE * sizeof(cl_int), BufferHost, &error); + return (error != CL_SUCCESS); + } + + cl_int RunTest() + { + cl_int error; + + clProgramWrapper program; + clKernelWrapper get_addr_kernel, indirect_access_kernel, + ptr_arith_kernel; + clMemWrapper dev_addr_buffer, buffer_long, buffer_int, + dev_addr_no_host_buffer; + + error = create_single_kernel_helper(context, &program, &get_addr_kernel, + 1, &program_source, "get_addr", + "-DBUF_SIZE=" BUF_SIZE_STR); + test_error(error, "Failed to create program with source\n"); + + indirect_access_kernel = + clCreateKernel(program, "indirect_access", &error); + test_error(error, "Failed to create kernel indirect_access\n"); + + ptr_arith_kernel = clCreateKernel(program, "ptr_arith", &error); + test_error(error, "Failed to create kernel ptr_arith\n"); + + buffer_long = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_ulong), nullptr, &error); + test_error(error, "clCreateBuffer failed\n"); + + buffer_int = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), + nullptr, &error); + test_error(error, "clCreateBuffer failed\n"); + + // Test a buffer with hostptr copied data + dev_addr_buffer = clCreateBuffer( + context, CL_MEM_READ_WRITE | address_type | CL_MEM_COPY_HOST_PTR, + sizeof(cl_int) * BUF_SIZE, BufferHost, &error); + test_error(error, "clCreateBuffer with device address 1 failed\n"); + + if (test_buffer(dev_addr_buffer, buffer_long, get_addr_kernel) + != TEST_PASS) + test_fail("test_buffer_host failed\n"); + + // Test a buffer which doesn't have any hostptr associated with it. + dev_addr_no_host_buffer = + clCreateBuffer(context, CL_MEM_READ_WRITE | address_type, + sizeof(cl_int) * BUF_SIZE, nullptr, &error); + test_error(error, "clCreateBuffer with device address 2 failed\n"); + + if (test_buffer(dev_addr_no_host_buffer, buffer_long, get_addr_kernel) + != TEST_PASS) + test_fail("test_buffer_no_host failed\n"); + + // Test a buffer passed indirectly + if (test_indirect_buffer(dev_addr_buffer, buffer_long, buffer_int, + indirect_access_kernel) + != TEST_PASS) + test_fail("test_indirect_buffer failed\n"); + + if (test_set_kernel_arg(dev_addr_buffer, buffer_int, ptr_arith_kernel) + != TEST_PASS) + test_fail("test_set_kernel_arg failed\n"); + + return TEST_PASS; + } + +private: + int BufferHost[BUF_SIZE]; + size_t global_size_one[3] = { 1, 1, 1 }; + cl_device_id device; + cl_context context; + cl_command_queue queue; + cl_mem_properties address_type; + + int check_device_address_from_api(cl_mem buf, + cl_mem_device_address_EXT &Addr) + { + Addr = 0; + cl_int error = clGetMemObjectInfo(buf, CL_MEM_DEVICE_ADDRESS_EXT, + sizeof(Addr), &Addr, NULL); + if (error) + { + print_error( + error, + "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) failed\n"); + return error; + } + if (Addr == 0) + { + print_error(error, + "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) " + "returned 0 as address\n"); + return CL_INVALID_VALUE; + } + return CL_SUCCESS; + } + + int test_buffer(clMemWrapper &dev_addr_buffer, clMemWrapper &plain_buffer, + clKernelWrapper &get_addr_kernel) + { + cl_int error = 0; + cl_ulong DeviceAddrFromKernel = 0; + cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + + for (int i = 0; i < BUF_SIZE; ++i) + { + BufferHost[i] = i; + } + + error = + check_device_address_from_api(dev_addr_buffer, DeviceAddrFromAPI); + test_error_fail(error, + "device address buffer does not have device address") + + error = clEnqueueWriteBuffer(queue, dev_addr_buffer, + CL_FALSE, // block + 0, BUF_SIZE * sizeof(cl_int), + BufferHost, 0, NULL, NULL); + test_error_fail(error, + "clEnqueueWriteBuffer of dev_addr_buffer failed\n"); + + error = clSetKernelArg(get_addr_kernel, 0, sizeof(cl_mem), + &dev_addr_buffer); + test_error_fail(error, "clSetKernelArg 0 failed\n"); + error = + clSetKernelArg(get_addr_kernel, 1, sizeof(cl_mem), &plain_buffer); + test_error_fail(error, "clSetKernelArg 1 failed\n"); + + error = clEnqueueNDRangeKernel(queue, get_addr_kernel, 1, NULL, + global_size_one, NULL, 0, NULL, NULL); + test_error_fail(error, "clNDRangeKernel failed\n"); + + error = clEnqueueReadBuffer(queue, dev_addr_buffer, CL_FALSE, 0, + sizeof(cl_int) * BUF_SIZE, BufferHost, 0, + NULL, NULL); + test_error_fail(error, + "clEnqueueReadBuffer of dev_addr_buffer failed\n"); + + error = clEnqueueReadBuffer(queue, plain_buffer, CL_FALSE, 0, + sizeof(cl_ulong), &DeviceAddrFromKernel, 0, + NULL, NULL); + test_error_fail(error, "clEnqueueReadBuffer of buffer failed\n"); + + error = clFinish(queue); + test_error_fail(error, "clFinish failed\n"); + + for (int i = 0; i < BUF_SIZE; ++i) + { + if (BufferHost[i] != (i + 1)) + { + test_fail("BufferHost[%i] expected " + "to be: %i, but is: %i\n", + i, i + 1, BufferHost[i]); + } + } + + if (DeviceAddrFromAPI != DeviceAddrFromKernel) + { + test_fail("DeviceAddrFromAPI(%lu) != DeviceAddrFromKernel(%lu)\n", + DeviceAddrFromAPI, DeviceAddrFromKernel); + } + return TEST_PASS; + } + + int test_indirect_buffer(clMemWrapper &dev_addr_buffer, + clMemWrapper &buffer_in_long, + clMemWrapper &buffer_out_int, + clKernelWrapper &ind_access_kernel) + { + cl_int error = 0; + cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + + int DataIn = 0x12348765; + int DataOut = -1; + + // A devaddr buffer with the payload data. + error = clEnqueueWriteBuffer(queue, dev_addr_buffer, + CL_TRUE, // block + 0, sizeof(cl_int), &DataIn, 0, NULL, NULL); + test_error_fail(error, + "clEnqueueWriteBuffer of dev_addr_buffer failed\n"); + + error = + check_device_address_from_api(dev_addr_buffer, DeviceAddrFromAPI); + test_error_fail(error, + "device address buffer does not have device address") + + // A basic buffer used to pass the other buffer's address. + error = clEnqueueWriteBuffer(queue, buffer_in_long, + CL_TRUE, // block + 0, sizeof(cl_long), &DeviceAddrFromAPI, + 0, NULL, NULL); + test_error_fail(error, + "clEnqueueWriteBuffer of dev_addr_buffer failed\n"); + + error = clSetKernelArg(ind_access_kernel, 0, sizeof(cl_mem), + &buffer_in_long); + test_error_fail(error, "clSetKernelArg 0 failed\n"); + error = clSetKernelArg(ind_access_kernel, 1, sizeof(cl_mem), + &buffer_out_int); + test_error_fail(error, "clSetKernelArg 1 failed\n"); + + error = clSetKernelExecInfo(ind_access_kernel, + CL_KERNEL_EXEC_INFO_DEVICE_PTRS_EXT, + sizeof(void *), &DeviceAddrFromAPI); + test_error_fail(error, + "Setting indirect access for " + "device ptrs failed!\n"); + + error = clEnqueueNDRangeKernel(queue, ind_access_kernel, 1, NULL, + global_size_one, NULL, 0, NULL, NULL); + test_error_fail(error, "clNDRangeKernel failed\n"); + + error = clEnqueueReadBuffer(queue, buffer_out_int, CL_FALSE, 0, + sizeof(cl_int), &DataOut, 0, NULL, NULL); + test_error_fail(error, "clEnqueueReadBuffer of buffer failed\n"); + + error = clFinish(queue); + test_error_fail(error, "clFinish failed\n"); + + for (int i = 0; i < BUF_SIZE; ++i) + { + if (BufferHost[i] != i + 1) + { + test_fail("PinnedBufferHost[%i] expected " + "to be: %i, but is: %i\n", + i, i + 1, BufferHost[i]); + } + } + + if (DataIn != DataOut) + { + test_fail("Passing data via indirect buffers failed. " + "Got: %i expected: %i\n", + DataOut, DataIn); + } + return TEST_PASS; + } + + int test_set_kernel_arg(clMemWrapper &dev_addr_buffer, + clMemWrapper &buffer_out_int, + clKernelWrapper &ptr_arith_kernel) + { + cl_int error = 0; + cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + int DataOut = -1; + int DataIn = 0x12348765; + + clSetKernelArgDevicePointerEXT_fn clSetKernelArgDevicePointer = + (clSetKernelArgDevicePointerEXT_fn) + clGetExtensionFunctionAddressForPlatform( + getPlatformFromDevice(device), + "clSetKernelArgDevicePointerEXT"); + if (clSetKernelArgDevicePointer == nullptr) + test_error_fail( + error, "Cannot get address of clSetKernelArgDevicePointerEXT"); + + error = clEnqueueWriteBuffer(queue, dev_addr_buffer, + CL_FALSE, // block + 0, sizeof(cl_int), &DataIn, 0, NULL, NULL); + test_error_fail(error, + "clEnqueueWriteBuffer of dev_addr_buffer failed\n"); + + error = + check_device_address_from_api(dev_addr_buffer, DeviceAddrFromAPI); + test_error_fail(error, "dev_addr_buffer does not have device address") + + error = clSetKernelArgDevicePointer( + ptr_arith_kernel, 0, + (cl_mem_device_address_EXT)(((cl_uint *)DeviceAddrFromAPI) + + 2)); + test_error_fail(error, "clSetKernelArgDevicePointer failed\n"); + error = clSetKernelArg(ptr_arith_kernel, 1, sizeof(cl_mem), + &buffer_out_int); + test_error_fail(error, "clSetKernelArg 1 failed\n"); + + error = clEnqueueNDRangeKernel(queue, ptr_arith_kernel, 1, NULL, + global_size_one, NULL, 0, NULL, NULL); + test_error_fail(error, "clNDRangeKernel failed\n"); + + error = clEnqueueReadBuffer(queue, buffer_out_int, CL_FALSE, 0, + sizeof(cl_int), &DataOut, 0, NULL, NULL); + test_error_fail(error, "clEnqueueReadBuffer of buffer failed\n"); + + error = clFinish(queue); + test_error_fail(error, "clFinish failed\n"); + + if (DataIn != DataOut) + { + test_fail("Negative offsetting from passed in pointer failed: " + "got: %i expected: %i", + DataOut, DataIn); + } + return TEST_PASS; + } +}; + +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_mem_properties address_type) +{ + if (!is_extension_available(device, "cl_ext_buffer_device_address")) + { + log_info("The device does not support the " + "cl_ext_buffer_device_address extension.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_version ext_version = + get_extension_version(device, "cl_ext_buffer_device_address"); + if (ext_version != CL_MAKE_VERSION(0, 9, 0)) + { + log_info("The test is written against cl_ext_buffer_device_address " + "extension version 0.9.0, device supports version: %u.%u.%u\n", + CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), + CL_VERSION_PATCH(ext_version)); + return TEST_SKIPPED_ITSELF; + } + + BufferDeviceAddressTest test_fixture = + BufferDeviceAddressTest(device, context, queue, address_type); + + if (test_fixture.Skip()) + { + log_info("TEST FIXTURE SKIP\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int error = test_fixture.RunTest(); + test_error_ret(error, "Test Failed", TEST_FAIL); + + return TEST_PASS; +} + +} + +int test_shared_address(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + CL_MEM_DEVICE_SHARED_ADDRESS_EXT); +} + +int test_private_address(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT); +} diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp new file mode 100644 index 0000000000..a1ad19f550 --- /dev/null +++ b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp @@ -0,0 +1,28 @@ +// 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 "procs.h" +#include "harness/testHarness.h" + +test_definition test_list[] = { + ADD_TEST(private_address), + ADD_TEST(shared_address), +}; + + +int main(int argc, const char *argv[]) +{ + return runTestHarness(argc, argv, ARRAY_SIZE(test_list), test_list, false, + 0); +} diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/procs.h b/test_conformance/extensions/cl_ext_buffer_device_address/procs.h new file mode 100644 index 0000000000..b682cc3b13 --- /dev/null +++ b/test_conformance/extensions/cl_ext_buffer_device_address/procs.h @@ -0,0 +1,25 @@ +// 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. +// +#ifndef CL_EXT_BUFFER_DEVICE_ADDRESS_H +#define CL_EXT_BUFFER_DEVICE_ADDRESS_H + +#include + +int test_private_address(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +int test_shared_address(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); + +#endif /* CL_EXT_BUFFER_DEVICE_ADDRESS_H */ From 51a2140fc1c901587eac23ddd8b19361a94db8f9 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Sat, 14 Dec 2024 22:34:42 +0200 Subject: [PATCH 2/9] update tests of cl_ext_buffer_device_address to 0.9.1 --- .../buffer_device_address.cpp | 11 ++--------- .../extensions/cl_ext_buffer_device_address/main.cpp | 1 - 2 files changed, 2 insertions(+), 10 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index 1332062529..d5882dd42b 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -372,10 +372,10 @@ int MakeAndRunTest(cl_device_id device, cl_context context, cl_version ext_version = get_extension_version(device, "cl_ext_buffer_device_address"); - if (ext_version != CL_MAKE_VERSION(0, 9, 0)) + if (ext_version != CL_MAKE_VERSION(0, 9, 1)) { log_info("The test is written against cl_ext_buffer_device_address " - "extension version 0.9.0, device supports version: %u.%u.%u\n", + "extension version 0.9.1, device supports version: %u.%u.%u\n", CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), CL_VERSION_PATCH(ext_version)); return TEST_SKIPPED_ITSELF; @@ -398,13 +398,6 @@ int MakeAndRunTest(cl_device_id device, cl_context context, } -int test_shared_address(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements) -{ - return MakeAndRunTest(device, context, queue, - CL_MEM_DEVICE_SHARED_ADDRESS_EXT); -} - int test_private_address(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp index a1ad19f550..84cd302dae 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp @@ -17,7 +17,6 @@ test_definition test_list[] = { ADD_TEST(private_address), - ADD_TEST(shared_address), }; From 19d99a0b0cb5bbe6c29ac6ee31d39aeae4b70de8 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Thu, 19 Dec 2024 16:47:39 +0200 Subject: [PATCH 3/9] test_cl_ext_buffer_device_address: use the new registration framework --- .../buffer_device_address.cpp | 41 +++--------------- .../cl_ext_buffer_device_address/main.cpp | 43 ++++++++++++++++--- .../cl_ext_buffer_device_address/procs.h | 25 ----------- 3 files changed, 44 insertions(+), 65 deletions(-) delete mode 100644 test_conformance/extensions/cl_ext_buffer_device_address/procs.h diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index d5882dd42b..0b40ff9fdb 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -13,7 +13,6 @@ // limitations under the License. // -#include "procs.h" #include "harness/typeWrappers.h" #define BUF_SIZE 1024 @@ -62,7 +61,7 @@ class BufferDeviceAddressTest { cl_int error = 0; clMemWrapper TempBuffer = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE | address_type), - (size_t)BUF_SIZE * sizeof(cl_int), BufferHost, &error); + (size_t)BUF_SIZE * sizeof(cl_int), nullptr, &error); return (error != CL_SUCCESS); } @@ -360,33 +359,16 @@ class BufferDeviceAddressTest { } }; -int MakeAndRunTest(cl_device_id device, cl_context context, - cl_command_queue queue, cl_mem_properties address_type) -{ - if (!is_extension_available(device, "cl_ext_buffer_device_address")) - { - log_info("The device does not support the " - "cl_ext_buffer_device_address extension.\n"); - return TEST_SKIPPED_ITSELF; - } - - cl_version ext_version = - get_extension_version(device, "cl_ext_buffer_device_address"); - if (ext_version != CL_MAKE_VERSION(0, 9, 1)) - { - log_info("The test is written against cl_ext_buffer_device_address " - "extension version 0.9.1, device supports version: %u.%u.%u\n", - CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), - CL_VERSION_PATCH(ext_version)); - return TEST_SKIPPED_ITSELF; - } +} - BufferDeviceAddressTest test_fixture = - BufferDeviceAddressTest(device, context, queue, address_type); +REGISTER_TEST(private_address) +{ + BufferDeviceAddressTest test_fixture = BufferDeviceAddressTest( + device, context, queue, CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT); if (test_fixture.Skip()) { - log_info("TEST FIXTURE SKIP\n"); + log_info("Test fixture skip\n"); return TEST_SKIPPED_ITSELF; } @@ -395,12 +377,3 @@ int MakeAndRunTest(cl_device_id device, cl_context context, return TEST_PASS; } - -} - -int test_private_address(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements) -{ - return MakeAndRunTest(device, context, queue, - CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT); -} diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp index 84cd302dae..1ce3000803 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp @@ -12,16 +12,47 @@ // See the License for the specific language governing permissions and // limitations under the License. // -#include "procs.h" + #include "harness/testHarness.h" +#include "harness/deviceInfo.h" + +test_status InitCL(cl_device_id device) +{ + auto version = get_device_cl_version(device); + auto expected_min_version = Version(3, 0); -test_definition test_list[] = { - ADD_TEST(private_address), -}; + if (version < expected_min_version) + { + version_expected_info("Test", "OpenCL", + expected_min_version.to_string().c_str(), + version.to_string().c_str()); + return TEST_SKIP; + } + if (!is_extension_available(device, "cl_ext_buffer_device_address")) + { + log_info("The device does not support the " + "cl_ext_buffer_device_address extension.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_version ext_version = + get_extension_version(device, "cl_ext_buffer_device_address"); + if (ext_version != CL_MAKE_VERSION(0, 9, 1)) + { + log_info("The test is written against cl_ext_buffer_device_address " + "extension version 0.9.1, device supports version: %u.%u.%u\n", + CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), + CL_VERSION_PATCH(ext_version)); + return TEST_SKIPPED_ITSELF; + } + + return TEST_PASS; +} int main(int argc, const char *argv[]) { - return runTestHarness(argc, argv, ARRAY_SIZE(test_list), test_list, false, - 0); + return runTestHarnessWithCheck( + argc, argv, test_registry::getInstance().num_tests(), + test_registry::getInstance().definitions(), false, 0, InitCL); } diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/procs.h b/test_conformance/extensions/cl_ext_buffer_device_address/procs.h deleted file mode 100644 index b682cc3b13..0000000000 --- a/test_conformance/extensions/cl_ext_buffer_device_address/procs.h +++ /dev/null @@ -1,25 +0,0 @@ -// 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. -// -#ifndef CL_EXT_BUFFER_DEVICE_ADDRESS_H -#define CL_EXT_BUFFER_DEVICE_ADDRESS_H - -#include - -int test_private_address(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements); -int test_shared_address(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements); - -#endif /* CL_EXT_BUFFER_DEVICE_ADDRESS_H */ From 1715f0507e21aea73ee1cac7de400bd1b24e3442 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Fri, 17 Jan 2025 21:16:24 +0200 Subject: [PATCH 4/9] fix review comments --- .../buffer_device_address.cpp | 75 +++++++++++++++---- 1 file changed, 61 insertions(+), 14 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index 0b40ff9fdb..5bc1d271b1 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -96,8 +96,9 @@ class BufferDeviceAddressTest { test_error(error, "clCreateBuffer failed\n"); // Test a buffer with hostptr copied data - dev_addr_buffer = clCreateBuffer( - context, CL_MEM_READ_WRITE | address_type | CL_MEM_COPY_HOST_PTR, + cl_mem_properties buf_props[] = { address_type, CL_TRUE, 0 }; + dev_addr_buffer = clCreateBufferWithProperties( + context, buf_props, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int) * BUF_SIZE, BufferHost, &error); test_error(error, "clCreateBuffer with device address 1 failed\n"); @@ -142,13 +143,8 @@ class BufferDeviceAddressTest { Addr = 0; cl_int error = clGetMemObjectInfo(buf, CL_MEM_DEVICE_ADDRESS_EXT, sizeof(Addr), &Addr, NULL); - if (error) - { - print_error( - error, - "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) failed\n"); - return error; - } + test_error(error, + "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) failed\n"); if (Addr == 0) { print_error(error, @@ -159,6 +155,57 @@ class BufferDeviceAddressTest { return CL_SUCCESS; } + int check_svm_buffer() + { + clSVMWrapper svm_buffer; + clMemWrapper buffer; + cl_int error = 0; + + cl_device_svm_capabilities svm_caps = 0; + error = clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, + sizeof(svm_caps), &svm_caps, NULL); + if (error != CL_SUCCESS) + { + print_error(error, "Unable to get SVM capabilities, skipping"); + return 0; + } + if (svm_caps == 0) + { + print_error(error, "Device has no SVM capabilities, skipping"); + return 0; + } + + svm_buffer = + clSVMWrapper(context, sizeof(cl_int) * BUF_SIZE, + CL_DEVICE_SVM_COARSE_GRAIN_BUFFER | CL_MEM_READ_WRITE); + if (svm_buffer() == nullptr) + { + test_error(CL_OUT_OF_RESOURCES, "SVM allocation failed"); + } + + cl_mem_properties buf_props[] = { address_type, CL_TRUE, 0 }; + buffer = clCreateBufferWithProperties( + context, buf_props, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, + sizeof(cl_int) * BUF_SIZE, svm_buffer(), &error); + test_error(error, "clCreateBuffer with device address 1 failed\n"); + + cl_mem_device_address_EXT Addr = 0; + error = clGetMemObjectInfo(buffer, CL_MEM_DEVICE_ADDRESS_EXT, + sizeof(Addr), &Addr, NULL); + test_error(error, + "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) failed\n"); + + if ((void *)Addr != svm_buffer()) + { + print_error(error, + "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) " + "returned different address than clSVMAlloc\n"); + return CL_INVALID_VALUE; + } + return CL_SUCCESS; + } + + int test_buffer(clMemWrapper &dev_addr_buffer, clMemWrapper &plain_buffer, clKernelWrapper &get_addr_kernel) { @@ -174,12 +221,12 @@ class BufferDeviceAddressTest { error = check_device_address_from_api(dev_addr_buffer, DeviceAddrFromAPI); test_error_fail(error, - "device address buffer does not have device address") + "device address buffer does not have device address"); - error = clEnqueueWriteBuffer(queue, dev_addr_buffer, - CL_FALSE, // block - 0, BUF_SIZE * sizeof(cl_int), - BufferHost, 0, NULL, NULL); + error = clEnqueueWriteBuffer(queue, dev_addr_buffer, + CL_FALSE, // block + 0, BUF_SIZE * sizeof(cl_int), BufferHost, + 0, NULL, NULL); test_error_fail(error, "clEnqueueWriteBuffer of dev_addr_buffer failed\n"); From a84735dea5990a818f481c7ff887989a5fb5aa88 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Mon, 20 Jan 2025 11:56:54 +0200 Subject: [PATCH 5/9] add fixes by Karol Herbst --- .../buffer_device_address.cpp | 24 ++++++++++--------- 1 file changed, 13 insertions(+), 11 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index 5bc1d271b1..d91e3e6d5c 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -59,8 +59,10 @@ class BufferDeviceAddressTest { bool Skip() { cl_int error = 0; - clMemWrapper TempBuffer = clCreateBuffer( - context, (cl_mem_flags)(CL_MEM_READ_WRITE | address_type), + + cl_mem_properties buf_props[] = { address_type, CL_TRUE, 0 }; + clMemWrapper TempBuffer = clCreateBufferWithProperties( + context, buf_props, CL_MEM_READ_WRITE, (size_t)BUF_SIZE * sizeof(cl_int), nullptr, &error); return (error != CL_SUCCESS); } @@ -107,9 +109,9 @@ class BufferDeviceAddressTest { test_fail("test_buffer_host failed\n"); // Test a buffer which doesn't have any hostptr associated with it. - dev_addr_no_host_buffer = - clCreateBuffer(context, CL_MEM_READ_WRITE | address_type, - sizeof(cl_int) * BUF_SIZE, nullptr, &error); + dev_addr_no_host_buffer = clCreateBufferWithProperties( + context, buf_props, CL_MEM_READ_WRITE, sizeof(cl_int) * BUF_SIZE, + nullptr, &error); test_error(error, "clCreateBuffer with device address 2 failed\n"); if (test_buffer(dev_addr_no_host_buffer, buffer_long, get_addr_kernel) @@ -189,7 +191,7 @@ class BufferDeviceAddressTest { sizeof(cl_int) * BUF_SIZE, svm_buffer(), &error); test_error(error, "clCreateBuffer with device address 1 failed\n"); - cl_mem_device_address_EXT Addr = 0; + cl_mem_device_address_ext Addr = 0; error = clGetMemObjectInfo(buffer, CL_MEM_DEVICE_ADDRESS_EXT, sizeof(Addr), &Addr, NULL); test_error(error, @@ -374,12 +376,12 @@ class BufferDeviceAddressTest { error = check_device_address_from_api(dev_addr_buffer, DeviceAddrFromAPI); - test_error_fail(error, "dev_addr_buffer does not have device address") + test_error_fail(error, "dev_addr_buffer does not have device address"); - error = clSetKernelArgDevicePointer( - ptr_arith_kernel, 0, - (cl_mem_device_address_EXT)(((cl_uint *)DeviceAddrFromAPI) - + 2)); + cl_mem_device_address_ext DeviceAddrFromAPIP2 = + (cl_mem_device_address_ext)(((cl_uint *)DeviceAddrFromAPI) + 2); + error = clSetKernelArgDevicePointer(ptr_arith_kernel, 0, + DeviceAddrFromAPIP2); test_error_fail(error, "clSetKernelArgDevicePointer failed\n"); error = clSetKernelArg(ptr_arith_kernel, 1, sizeof(cl_mem), &buffer_out_int); From 2f748ebd56f2841129ccda6553f591dfd5f93089 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Mon, 20 Jan 2025 19:39:39 +0200 Subject: [PATCH 6/9] cl_ext_buffer_device_address: update to 1.0.0 --- .../extensions/cl_ext_buffer_device_address/main.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp index 1ce3000803..1c0e221aeb 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp @@ -38,10 +38,10 @@ test_status InitCL(cl_device_id device) cl_version ext_version = get_extension_version(device, "cl_ext_buffer_device_address"); - if (ext_version != CL_MAKE_VERSION(0, 9, 1)) + if (ext_version != CL_MAKE_VERSION(1, 0, 0)) { log_info("The test is written against cl_ext_buffer_device_address " - "extension version 0.9.1, device supports version: %u.%u.%u\n", + "extension version 1.0.0, device supports version: %u.%u.%u\n", CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), CL_VERSION_PATCH(ext_version)); return TEST_SKIPPED_ITSELF; From 769f4a840a50e7545096fa241a92f3011f94379a Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Thu, 27 Feb 2025 16:18:01 +0200 Subject: [PATCH 7/9] cl_ext_buffer_device_address: update to 1.0.2 --- .../buffer_device_address.cpp | 8 ++++---- .../extensions/cl_ext_buffer_device_address/main.cpp | 4 ++-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index d91e3e6d5c..7fe4522b8f 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -140,7 +140,7 @@ class BufferDeviceAddressTest { cl_mem_properties address_type; int check_device_address_from_api(cl_mem buf, - cl_mem_device_address_EXT &Addr) + cl_mem_device_address_ext &Addr) { Addr = 0; cl_int error = clGetMemObjectInfo(buf, CL_MEM_DEVICE_ADDRESS_EXT, @@ -213,7 +213,7 @@ class BufferDeviceAddressTest { { cl_int error = 0; cl_ulong DeviceAddrFromKernel = 0; - cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + cl_mem_device_address_ext DeviceAddrFromAPI = 0; for (int i = 0; i < BUF_SIZE; ++i) { @@ -281,7 +281,7 @@ class BufferDeviceAddressTest { clKernelWrapper &ind_access_kernel) { cl_int error = 0; - cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + cl_mem_device_address_ext DeviceAddrFromAPI = 0; int DataIn = 0x12348765; int DataOut = -1; @@ -355,7 +355,7 @@ class BufferDeviceAddressTest { clKernelWrapper &ptr_arith_kernel) { cl_int error = 0; - cl_mem_device_address_EXT DeviceAddrFromAPI = 0; + cl_mem_device_address_ext DeviceAddrFromAPI = 0; int DataOut = -1; int DataIn = 0x12348765; diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp index 1c0e221aeb..31d01fa97b 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/main.cpp @@ -38,10 +38,10 @@ test_status InitCL(cl_device_id device) cl_version ext_version = get_extension_version(device, "cl_ext_buffer_device_address"); - if (ext_version != CL_MAKE_VERSION(1, 0, 0)) + if (ext_version != CL_MAKE_VERSION(1, 0, 2)) { log_info("The test is written against cl_ext_buffer_device_address " - "extension version 1.0.0, device supports version: %u.%u.%u\n", + "extension version 1.0.2, device supports version: %u.%u.%u\n", CL_VERSION_MAJOR(ext_version), CL_VERSION_MINOR(ext_version), CL_VERSION_PATCH(ext_version)); return TEST_SKIPPED_ITSELF; From 3a88039b66e11cc2df1a95bdc32c41c665eef6b5 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Thu, 27 Feb 2025 16:34:49 +0200 Subject: [PATCH 8/9] buffer_device_address.cpp: fix printf format specifiers --- .../cl_ext_buffer_device_address/buffer_device_address.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index 7fe4522b8f..9d5b775f39 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -14,6 +14,7 @@ // #include "harness/typeWrappers.h" +#include #define BUF_SIZE 1024 #define BUF_SIZE_STR "1024" @@ -269,8 +270,10 @@ class BufferDeviceAddressTest { if (DeviceAddrFromAPI != DeviceAddrFromKernel) { - test_fail("DeviceAddrFromAPI(%lu) != DeviceAddrFromKernel(%lu)\n", - DeviceAddrFromAPI, DeviceAddrFromKernel); + test_fail("DeviceAddrFromAPI(%" PRIu64 + ") != DeviceAddrFromKernel(%" PRIu64 ")\n", + (uint64_t)DeviceAddrFromAPI, + (uint64_t)DeviceAddrFromKernel); } return TEST_PASS; } From bbbb807e7677daa94751babe086453e383892a05 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Wed, 19 Mar 2025 20:21:44 +0200 Subject: [PATCH 9/9] cl_ext_buffer_device_address : enable also testing on SVM memory --- .../buffer_device_address.cpp | 28 +++++++++++-------- 1 file changed, 17 insertions(+), 11 deletions(-) diff --git a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp index 9d5b775f39..97872613c7 100644 --- a/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp +++ b/test_conformance/extensions/cl_ext_buffer_device_address/buffer_device_address.cpp @@ -129,6 +129,9 @@ class BufferDeviceAddressTest { != TEST_PASS) test_fail("test_set_kernel_arg failed\n"); + if (test_svm_buffer() == TEST_FAIL) + test_fail("test_svm_buffer failed\n"); + return TEST_PASS; } @@ -158,7 +161,7 @@ class BufferDeviceAddressTest { return CL_SUCCESS; } - int check_svm_buffer() + int test_svm_buffer() { clSVMWrapper svm_buffer; clMemWrapper buffer; @@ -169,13 +172,18 @@ class BufferDeviceAddressTest { sizeof(svm_caps), &svm_caps, NULL); if (error != CL_SUCCESS) { - print_error(error, "Unable to get SVM capabilities, skipping"); - return 0; + print_missing_feature(error, + "Unable to get SVM capabilities, " + "skipping"); + return TEST_SKIP; } - if (svm_caps == 0) + if ((svm_caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) == 0) { - print_error(error, "Device has no SVM capabilities, skipping"); - return 0; + print_missing_feature(error, + "Device doesn't support " + "CL_DEVICE_SVM_COARSE_" + "GRAIN_BUFFER, skipping"); + return TEST_SKIP; } svm_buffer = @@ -200,12 +208,10 @@ class BufferDeviceAddressTest { if ((void *)Addr != svm_buffer()) { - print_error(error, - "clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) " - "returned different address than clSVMAlloc\n"); - return CL_INVALID_VALUE; + test_fail("clGetMemObjectInfo(CL_MEM_DEVICE_ADDRESS_EXT) " + "returned different address than clSVMAlloc\n"); } - return CL_SUCCESS; + return TEST_PASS; }