From bd167754d9690a23d0210d04a9913de386ba31b8 Mon Sep 17 00:00:00 2001 From: gtrebuchet-arm Date: Tue, 2 Dec 2025 17:40:16 +0000 Subject: [PATCH] Add multi device and negative tests for cl_ext_buffer_device_address (#2561) The tests checks that invalid parameters for clSetKernelArgDevicePointerEXT are reported successfully and ensure that a kernel can access a buffer from their respective device address on each device in a multi device context. Signed-off-by: Guillaume Trebuchet --- .../buffer_device_address.cpp | 182 ++++++++++++++++++ 1 file changed, 182 insertions(+) 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 97872613..50eb5011 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,7 +14,10 @@ // #include "harness/typeWrappers.h" +#include "harness/extensionHelpers.h" #include +#include +#include #define BUF_SIZE 1024 #define BUF_SIZE_STR "1024" @@ -421,6 +424,8 @@ private: REGISTER_TEST(private_address) { + REQUIRE_EXTENSION("cl_ext_buffer_device_address"); + BufferDeviceAddressTest test_fixture = BufferDeviceAddressTest( device, context, queue, CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT); @@ -435,3 +440,180 @@ REGISTER_TEST(private_address) return TEST_PASS; } + +REGISTER_TEST(private_address_multi_device) +{ + REQUIRE_EXTENSION("cl_ext_buffer_device_address"); + + cl_platform_id platform = 0; + cl_int error = CL_SUCCESS; + cl_uint numDevices = 0; + + error = clGetPlatformIDs(1, &platform, NULL); + test_error_ret(error, "Unable to get platform\n", TEST_FAIL); + + /* Get some devices */ + error = + clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &numDevices); + test_error_ret(error, "Unable to get multiple devices\n", TEST_FAIL); + + if (numDevices < 2) + { + log_info( + "WARNING: multi device test unable to get multiple devices via " + "CL_DEVICE_TYPE_ALL (got %u devices). Skipping test...\n", + numDevices); + return TEST_SKIPPED_ITSELF; + } + + std::vector devices(numDevices); + error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, numDevices, + devices.data(), &numDevices); + test_error_ret(error, "Unable to get multiple devices\n", TEST_FAIL); + + GET_PFN(devices[0], clSetKernelArgDevicePointerEXT); + + cl_context_properties properties[] = { CL_CONTEXT_PLATFORM, + (cl_context_properties)platform, 0 }; + clContextWrapper ctx = clCreateContext( + properties, numDevices, devices.data(), nullptr, nullptr, &error); + test_error_ret(error, "Unable to create context\n", TEST_FAIL); + + /* Create buffer */ + cl_mem_properties props[] = { CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT, CL_TRUE, + 0 }; + clMemWrapper buffer = clCreateBufferWithProperties( + ctx, props, CL_MEM_READ_WRITE, 16, nullptr, &error); + std::vector addresses(numDevices); + error = + clGetMemObjectInfo(buffer, CL_MEM_DEVICE_ADDRESS_EXT, + sizeof(cl_mem_device_address_ext) * addresses.size(), + addresses.data(), nullptr); + test_error_ret(error, "clGetMemObjectInfo failed\n", TEST_FAIL); + + std::vector queues(numDevices); + for (cl_uint i = 0; i < numDevices; ++i) + { + queues[i] = clCreateCommandQueue(ctx, devices[i], 0, &error); + test_error_ret(error, "Unable to create command queue\n", TEST_FAIL); + } + static std::string source = R"( + void kernel test_device_address( + global ulong* ptr, + ulong value) + { + *ptr = value; + })"; + + clProgramWrapper program; + clKernelWrapper kernel; + const char *source_ptr = source.data(); + error = create_single_kernel_helper(ctx, &program, &kernel, 1, &source_ptr, + "test_device_address"); + test_error(error, "Unable to create test kernel"); + for (cl_uint i = 0; i < numDevices; ++i) + { + cl_command_queue queue = queues[i]; + + error = clSetKernelArgDevicePointerEXT(kernel, 0, 0); + test_error_fail(error, + "clSetKernelArgDevicePointerEXT failed with NULL " + "pointer argument\n"); + + error = clSetKernelArgDevicePointerEXT(kernel, 0, addresses[i] + 8); + test_error_ret(error, "Unable to set kernel arg\n", TEST_FAIL); + + const cl_ulong pattern = 0xAABBCCDDEEFF0011 + i; + error = clSetKernelArg(kernel, 1, sizeof(pattern), &pattern); + test_error_ret(error, "Unable to set kernel arg\n", TEST_FAIL); + + size_t gwo = 0; + size_t gws = 1; + size_t lws = 1; + error = clEnqueueNDRangeKernel(queue, kernel, 1, &gwo, &gws, &lws, 0, + nullptr, nullptr); + test_error_ret(error, "Unable to enqueue kernel\n", TEST_FAIL); + + error = clFinish(queue); + test_error_ret(error, "clFinish failed\n", TEST_FAIL); + + std::vector results(2, 0); + error = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0, + results.size() * sizeof(cl_ulong), + results.data(), 0, nullptr, nullptr); + test_error_ret(error, "clEnqueueReadBuffer failed\n", TEST_FAIL); + + if (results[1] != pattern) + test_fail("Test value doesn't match expected value\n"); + } + return TEST_PASS; +} + +REGISTER_TEST(negative_private_address) +{ + REQUIRE_EXTENSION("cl_ext_buffer_device_address"); + + cl_int error = CL_SUCCESS; + + GET_PFN(device, clSetKernelArgDevicePointerEXT); + + /* Create buffer */ + clMemWrapper buffer = clCreateBufferWithProperties( + context, nullptr, CL_MEM_READ_WRITE, 16, nullptr, &error); + cl_mem_device_address_ext address; + error = clGetMemObjectInfo(buffer, CL_MEM_DEVICE_ADDRESS_EXT, + sizeof(cl_mem_device_address_ext), &address, + nullptr); + test_failure_error_ret( + error, CL_INVALID_OPERATION, + "clGetMemObjectInfo should return CL_INVALID_OPERATION when: " + "\"the buffer was not created with CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT\"", + TEST_FAIL); + + static std::string source = R"( + void kernel test_device_address( + global ulong* ptr, + local ulong* ptr2, + ulong value) + { + *ptr = value; + })"; + + clProgramWrapper program; + clKernelWrapper kernel; + const char *source_ptr = source.data(); + error = create_single_kernel_helper(context, &program, &kernel, 1, + &source_ptr, "test_device_address"); + test_error(error, "Unable to create test kernel"); + + error = clSetKernelArgDevicePointerEXT(nullptr, 0, 0); + test_failure_error_ret( + error, CL_INVALID_KERNEL, + "clSetKernelArgDevicePointerEXT should return CL_INVALID_KERNEL when: " + "\"kernel is not a valid kernel object\"", + TEST_FAIL); + + error = clSetKernelArgDevicePointerEXT(kernel, 1, 0x15465); + test_failure_error_ret( + error, CL_INVALID_ARG_INDEX, + "clSetKernelArgDevicePointerEXT should return " + "CL_INVALID_ARG_INDEX when: " + "\"the expected kernel argument is not a pointer to global memory\"", + TEST_FAIL); + + error = clSetKernelArgDevicePointerEXT(kernel, 2, 0x15465); + test_failure_error_ret(error, CL_INVALID_ARG_INDEX, + "clSetKernelArgDevicePointerEXT should return " + "CL_INVALID_ARG_INDEX when: " + "\"the expected kernel argument is not a pointer\"", + TEST_FAIL); + + error = clSetKernelArgDevicePointerEXT(kernel, 3, 0x15465); + test_failure_error_ret(error, CL_INVALID_ARG_INDEX, + "clSetKernelArgDevicePointerEXT should return " + "CL_INVALID_ARG_INDEX when: " + "\"arg_index is not a valid argument index\"", + TEST_FAIL); + + return TEST_PASS; +}