mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-18 22:09:01 +00:00
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 <guillaume.trebuchet@arm.com>
This commit is contained in:
@@ -14,7 +14,10 @@
|
||||
//
|
||||
|
||||
#include "harness/typeWrappers.h"
|
||||
#include "harness/extensionHelpers.h"
|
||||
#include <cinttypes>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
|
||||
#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<cl_device_id> 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<cl_mem_device_address_ext> 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<clCommandQueueWrapper> 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<cl_ulong> 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;
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user