From b63ef2d8f6b30b00c0b0da4b534ecc2d03aa6b72 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 5 Aug 2025 10:17:00 -0700 Subject: [PATCH] add tests for unified SVM corner cases (#2436) This PR adds tests for a number of interesting unified SVM corner cases. Not all of these may be valid tests! If we decide that some of these tests are invalid, I will remove them. Added tests include: * Calling clSVMAllocWithPropertiesKHR to allocate zero bytes for each unified SVM type. * Calling clSVMFreeWithPropertiesKHR to free a NULL pointer. * Calling clEnqueueSVMFree to asynchronously free an empty set of SVM pointers. * Calling clEnqueueSVMFree to asynchronously free a NULL pointer. * Calling clSetKernelArgSVMPointer to set a NULL pointer kernel argument. * Calling clSetKernelArgSVMPointer to set a bogus pointer kernel argument. * Calling clSetKernelExecInfo with CL_KERNEL_EXEC_INFO_SVM_PTRS with an empty set of SVM pointers. * Calling clSetKernelExecInfo with CL_KERNEL_EXEC_INFO_SVM_PTRS with a NULL pointer. * Calling clSetKernelExecInfo with CL_KERNEL_EXEC_INFO_SVM_PTRS with a bogus pointer. * Calling clEnqueueSVMMemcpy with a size of zero and a NULL source or destination pointer. * Calling clEnqueueSVMMemcpy with a size of zero and a bogus source or destination pointer. * Calling clEnqueueSVMMemcpy with a size of zero and a valid source or destination pointer. * Calling clEnqueueSVMMemFill with a size of zero and a NULL destination pointer. * Calling clEnqueueSVMMemFill with a size of zero and a bogus destination pointer. * Calling clEnqueueSVMMemFill with a size of zero and a valid destination pointer. * Calling clEnqueueSVMMigrateMem with a size of zero and a NULL pointer. * Calling clEnqueueSVMMigrateMem with a size of zero and a valid pointer. --- test_conformance/SVM/CMakeLists.txt | 1 + .../SVM/test_unified_svm_corner_cases.cpp | 837 ++++++++++++++++++ test_conformance/SVM/unified_svm_fixture.h | 18 +- 3 files changed, 851 insertions(+), 5 deletions(-) create mode 100644 test_conformance/SVM/test_unified_svm_corner_cases.cpp diff --git a/test_conformance/SVM/CMakeLists.txt b/test_conformance/SVM/CMakeLists.txt index d35730db..d9da8703 100644 --- a/test_conformance/SVM/CMakeLists.txt +++ b/test_conformance/SVM/CMakeLists.txt @@ -17,6 +17,7 @@ set(${MODULE_NAME}_SOURCES test_shared_sub_buffers.cpp test_migrate.cpp test_unified_svm_consistency.cpp + test_unified_svm_corner_cases.cpp test_unified_svm_capabilities.cpp test_unified_svm_apis.cpp test_unified_svm_api_query_defaults.cpp diff --git a/test_conformance/SVM/test_unified_svm_corner_cases.cpp b/test_conformance/SVM/test_unified_svm_corner_cases.cpp new file mode 100644 index 00000000..2ec3be93 --- /dev/null +++ b/test_conformance/SVM/test_unified_svm_corner_cases.cpp @@ -0,0 +1,837 @@ +// +// Copyright (c) 2025 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 "unified_svm_fixture.h" +#include +#include + +struct UnifiedSVMCornerCaseAllocFree : UnifiedSVMBase +{ + UnifiedSVMCornerCaseAllocFree(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int run() override + { + cl_int err; + for (cl_uint ti = 0; ti < static_cast(deviceUSVMCaps.size()); + ti++) + { + log_info(" testing SVM type %u\n", ti); + + auto mem = get_usvm_wrapper(ti); + + log_info(" testing zero-byte allocation\n"); + err = mem->allocate(0); + test_error(err, "zero-byte SVM allocation failed"); + test_assert_error( + mem->get_ptr() == nullptr, + "zero-byte SVM allocation did not return a null pointer"); + } + + log_info(" testing NULL pointer free\n"); + err = clSVMFreeWithPropertiesKHR(context, nullptr, 0, nullptr); + test_error(err, "clSVMFreeWithPropertiesKHR with NULL pointer failed"); + + log_info(" testing asynchronous empty set free\n"); + clEventWrapper event; + err = clEnqueueSVMFree(queue, 0, nullptr, nullptr, nullptr, 0, nullptr, + &event); + test_error(err, "clEnqueueSVMFree with empty set failed"); + + err = clFinish(queue); + test_error(err, + "clFinish after clEnqueueSVMFree with empty set failed"); + + err = check_event_type(event, CL_COMMAND_SVM_FREE); + test_error(err, + "clEnqueueSVMFree did not return a " + "CL_COMMAND_SVM_FREE event"); + + log_info(" testing asynchronous NULL pointer free\n"); + event = nullptr; + void* svm_pointers[] = { nullptr }; + err = clEnqueueSVMFree(queue, 1, svm_pointers, nullptr, nullptr, 0, + nullptr, &event); + test_error(err, "clEnqueueSVMFree with NULL pointer failed"); + + err = clFinish(queue); + test_error(err, + "clFinish after clEnqueueSVMFree with NULL pointer failed"); + + err = check_event_type(event, CL_COMMAND_SVM_FREE); + test_error(err, + "clEnqueueSVMFree did not return a " + "CL_COMMAND_SVM_FREE event"); + + return CL_SUCCESS; + } +}; + +REGISTER_TEST(unified_svm_corner_case_alloc_free) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseAllocFree Test(context, device, queue, num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} + + +struct UnifiedSVMCornerCaseSetKernelArg : UnifiedSVMBase +{ + UnifiedSVMCornerCaseSetKernelArg(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_PointerKernelArg(const void* test) + { + cl_int err = clSetKernelArgSVMPointer(kernel_StorePointer, 0, test); + test_error(err, "clSetKernelArgSVMPointer failed"); + + err = clSetKernelArg(kernel_StorePointer, 1, sizeof(ptr_dst), &ptr_dst); + test_error(err, "clSetKernelArg failed"); + + size_t global_work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel_StorePointer, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, + nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + void* check = &err; + err = clEnqueueReadBuffer(queue, ptr_dst, CL_TRUE, 0, sizeof(cl_int*), + &check, 0, nullptr, nullptr); + test_error(err, "could not read output buffer"); + + test_assert_error(check == test, + "stored pointer does not match input pointer"); + + return CL_SUCCESS; + } + + cl_int setup() override + { + cl_int err = UnifiedSVMBase::setup(); + test_error(err, "UnifiedSVMBase setup failed"); + + const char* programString = R"( + // workaround for error: kernel parameter cannot be declared as a pointer to a pointer + struct s { const global int* ptr; }; + kernel void test_StorePointer(const global int* ptr, global struct s* dst) + { + dst[get_global_id(0)].ptr = ptr; + } + )"; + + clProgramWrapper program; + err = + create_single_kernel_helper(context, &program, &kernel_StorePointer, + 1, &programString, "test_StorePointer"); + test_error(err, "could not create StorePointer kernel"); + + ptr_dst = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int*), + nullptr, &err); + test_error(err, "could not create destination buffer"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + + log_info(" testing clSetKernelArgSVMPointer with a NULL pointer\n"); + err = test_PointerKernelArg(nullptr); + test_error(err, "clSetKernelArgSVMPointer with a NULL pointer failed"); + + log_info(" testing clSetKernelArgSVMPointer with a bogus pointer\n"); + err = test_PointerKernelArg((const void*)0xDEADBEEF); + test_error(err, "clSetKernelArgSVMPointer with a bogus pointer failed"); + + return CL_SUCCESS; + } + + clKernelWrapper kernel_StorePointer; + clMemWrapper ptr_dst; +}; + +REGISTER_TEST(unified_svm_corner_case_set_kernel_arg) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseSetKernelArg Test(context, device, queue, num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} + +struct UnifiedSVMCornerCaseSetKernelExecInfo : UnifiedSVMBase +{ + UnifiedSVMCornerCaseSetKernelExecInfo(cl_context context, + cl_device_id device, + cl_command_queue queue, + int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_EmptySet() + { + cl_int err = clSetKernelExecInfo( + kernel_OneArg, CL_KERNEL_EXEC_INFO_SVM_PTRS, 0, nullptr); + test_error(err, + "clSetKernelExecInfo with an empty set returned an error"); + + return CL_SUCCESS; + } + + cl_int test_NullPointer() + { + const void* svm_ptrs[] = { nullptr }; + cl_int err = + clSetKernelExecInfo(kernel_OneArg, CL_KERNEL_EXEC_INFO_SVM_PTRS, + sizeof(svm_ptrs), svm_ptrs); + test_error(err, + "clSetKernelExecInfo with a NULL pointer returned an error"); + + return CL_SUCCESS; + } + + cl_int test_BogusPointer() + { + const void* bogus = (const void*)0xDEADBEEF; + const void* svm_ptrs[] = { bogus }; + cl_int err = + clSetKernelExecInfo(kernel_OneArg, CL_KERNEL_EXEC_INFO_SVM_PTRS, + sizeof(svm_ptrs), svm_ptrs); + test_error( + err, "clSetKernelExecInfo with a bogus pointer returned an error"); + + return CL_SUCCESS; + } + + cl_int setup() override + { + cl_int err = UnifiedSVMBase::setup(); + test_error(err, "UnifiedSVMBase setup failed"); + + const char* programString = R"( + kernel void test_OneArg(global int* dst) + { + dst[get_global_id(0)] = -1; + } + )"; + + clProgramWrapper program; + err = create_single_kernel_helper(context, &program, &kernel_OneArg, 1, + &programString, "test_OneArg"); + test_error(err, "could not create OneArg kernel"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + + log_info(" testing clSetKernelExecInfo with an empty set\n"); + err = test_EmptySet(); + test_error(err, "clSetKernelExecInfo with an empty set failed"); + + log_info(" testing clSetKernelExecInfo with a NULL pointer\n"); + err = test_NullPointer(); + test_error(err, "clSetKernelExecInfo with a NULL pointer failed"); + + log_info(" testing clSetKernelExecInfo with a bogus pointer\n"); + err = test_BogusPointer(); + test_error(err, "clSetKernelExecInfo with a bogus pointer failed"); + + return CL_SUCCESS; + } + + clKernelWrapper kernel_OneArg; +}; + +REGISTER_TEST(unified_svm_corner_case_set_kernel_exec_info) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseSetKernelExecInfo Test(context, device, queue, + num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} + +struct UnifiedSVMCornerCaseMemcpy : UnifiedSVMBase +{ + UnifiedSVMCornerCaseMemcpy(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_NullPointer() + { + cl_int value = 0; + + clEventWrapper event; + cl_int err = clEnqueueSVMMemcpy(queue, CL_TRUE, nullptr, &value, 0, 0, + nullptr, &event); + test_error(err, + "clEnqueueSVMMemcpy with a NULL destination pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + event = nullptr; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &value, nullptr, 0, 0, nullptr, + &event); + test_error( + err, + "clEnqueueSVMMemcpy with a NULL source pointer returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + + return CL_SUCCESS; + } + + cl_int test_BogusPointer() + { + void* bogus = (void*)0xDEADBEEF; + cl_int value = 0; + + clEventWrapper event; + cl_int err = clEnqueueSVMMemcpy(queue, CL_TRUE, bogus, &value, 0, 0, + nullptr, &event); + test_error(err, + "clEnqueueSVMMemcpy with a bogus destination pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + event = nullptr; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &value, bogus, 0, 0, nullptr, + &event); + test_error( + err, + "clEnqueueSVMMemcpy with a bogus source pointer returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + + return CL_SUCCESS; + } + + cl_int test_ValidPointer(cl_uint typeIndex) + { + cl_int err; + cl_int value = 0; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + + clEventWrapper event; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, mem->get_ptr(), &value, 0, 0, + nullptr, &event); + test_error( + err, + "clEnqueueSVMMemcpy with valid SVM dst pointer returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + event = nullptr; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &value, mem->get_ptr(), 0, 0, + nullptr, &event); + test_error( + err, + "clEnqueueSVMMemcpy with valid SVM src pointer returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMCPY); + test_error(err, + "clEnqueueSVMMemcpy did not return a " + "CL_COMMAND_SVM_MEMCPY event"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + + log_info(" testing clEnqueueSVMMemcpy with a NULL pointer and a " + "size of zero\n"); + err = test_NullPointer(); + test_error( + err, + "clEnqueueSVMMemcpy with a NULL pointer and a size of zero failed"); + + log_info(" testing clEnqueueSVMMemcpy with a bogus pointer and a " + "size of zero\n"); + err = test_BogusPointer(); + test_error(err, + "clEnqueueSVMMemcpy with a bogus pointer and a size of zero " + "failed"); + + for (cl_uint ti = 0; ti < static_cast(deviceUSVMCaps.size()); + ti++) + { + log_info(" testing SVM type %u\n", ti); + + log_info(" testing clEnqueueSVMMemcpy with a valid pointer and " + "a size of zero\n"); + err = test_ValidPointer(ti); + test_error(err, + "clEnqueueSVMMemcpy with a valid pointer and a size of " + "zero failed"); + } + + return CL_SUCCESS; + } +}; + +REGISTER_TEST(unified_svm_corner_case_memcpy) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseMemcpy Test(context, device, queue, num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} + +struct UnifiedSVMCornerCaseMemFill : UnifiedSVMBase +{ + UnifiedSVMCornerCaseMemFill(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_NullPointer() + { + const cl_int pattern = 0; + + clEventWrapper event; + cl_int err = clEnqueueSVMMemFill( + queue, nullptr, &pattern, sizeof(pattern), 0, 0, nullptr, &event); + test_error(err, + "clEnqueueSVMMemFill with a NULL destination pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMFILL); + test_error(err, + "clEnqueueSVMMemFill did not return a " + "CL_COMMAND_SVM_MEMFILL event"); + + return CL_SUCCESS; + } + + cl_int test_BogusPointer() + { + void* bogus = (void*)0xDEADBEEF; + const cl_int pattern = 0; + + clEventWrapper event; + cl_int err = clEnqueueSVMMemFill( + queue, bogus, &pattern, sizeof(pattern), 0, 0, nullptr, &event); + test_error(err, + "clEnqueueSVMMemFill with a bogus destination pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMFILL); + test_error(err, + "clEnqueueSVMMemFill did not return a " + "CL_COMMAND_SVM_MEMFILL event"); + + return CL_SUCCESS; + } + + cl_int test_ValidPointer(cl_uint typeIndex) + { + cl_int err; + const cl_int pattern = 0; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + + clEventWrapper event; + err = clEnqueueSVMMemFill(queue, mem->get_ptr(), &pattern, + sizeof(pattern), 0, 0, nullptr, &event); + test_error(err, + "clEnqueueSVMMemFill with a valid destination pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MEMFILL); + test_error(err, + "clEnqueueSVMMemFill did not return a " + "CL_COMMAND_SVM_MEMFILL event"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + + log_info(" testing clEnqueueSVMMemFill with a NULL pointer and a " + "size of zero\n"); + err = test_NullPointer(); + test_error(err, + "clEnqueueSVMMemFill with a NULL pointer and a size of zero " + "failed"); + + log_info(" testing clEnqueueSVMMemFill with a bogus pointer and a " + "size of zero\n"); + err = test_BogusPointer(); + test_error( + err, + "clEnqueueSVMMemFill with a bogus pointer and a size of zero " + "failed"); + + for (cl_uint ti = 0; ti < static_cast(deviceUSVMCaps.size()); + ti++) + { + log_info(" testing SVM type %u\n", ti); + + log_info( + " testing clEnqueueSVMMemFill with a valid pointer and " + "a size of zero\n"); + err = test_ValidPointer(ti); + test_error(err, + "clEnqueueSVMMemFill with a valid pointer and a size of " + "zero failed"); + } + + return CL_SUCCESS; + } +}; + +REGISTER_TEST(unified_svm_corner_case_mem_fill) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseMemFill Test(context, device, queue, num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} + +struct UnifiedSVMCornerCaseMigrateMem : UnifiedSVMBase +{ + UnifiedSVMCornerCaseMigrateMem(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_NullPointer() + { + cl_int err; + + const void* svm_pointers[] = { nullptr }; + const size_t sizes[] = { 0 }; + clEventWrapper event; + err = clEnqueueSVMMigrateMem(queue, 1, svm_pointers, sizes, + CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED, 0, + nullptr, &event); + test_error( + err, + "clEnqueueSVMMigrateMem with a NULL pointer returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MIGRATE_MEM); + test_error(err, + "clEnqueueSVMMigrateMem did not return a " + "CL_COMMAND_SVM_MIGRATE_MEM event"); + + return CL_SUCCESS; + } + + cl_int test_ValidPointer(cl_uint typeIndex) + { + cl_int err; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + + const void* svm_pointers[] = { mem->get_ptr() }; + const size_t sizes[] = { 0 }; + clEventWrapper event; + err = clEnqueueSVMMigrateMem(queue, 1, svm_pointers, sizes, + CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED, 0, + nullptr, &event); + test_error(err, + "clEnqueueSVMMigrateMem with a valid pointer " + "returned an error"); + + err = check_event_type(event, CL_COMMAND_SVM_MIGRATE_MEM); + test_error(err, + "clEnqueueSVMMigrateMem did not return a " + "CL_COMMAND_SVM_MIGRATE_MEM event"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + + log_info(" testing clEnqueueSVMMigrateMem with a NULL pointer and a " + "size of zero\n"); + err = test_NullPointer(); + test_error( + err, + "clEnqueueSVMMigrateMem with a NULL pointer and a size of zero " + "failed"); + + for (cl_uint ti = 0; ti < static_cast(deviceUSVMCaps.size()); + ti++) + { + log_info(" testing SVM type %u\n", ti); + + log_info( + " testing clEnqueueSVMMigrateMem with a valid pointer and " + "a size of zero\n"); + err = test_ValidPointer(ti); + test_error( + err, + "clEnqueueSVMMigrateMem with a valid pointer and a size of " + "zero failed"); + } + + return CL_SUCCESS; + } +}; + +REGISTER_TEST(unified_svm_corner_case_migrate_mem) +{ + if (!is_extension_available(device, "cl_khr_unified_svm")) + { + log_info("cl_khr_unified_svm is not supported, skipping test.\n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int err; + + clContextWrapper contextWrapper; + clCommandQueueWrapper queueWrapper; + + // For now: create a new context and queue. + // If we switch to a new test executable and run the tests without + // forceNoContextCreation then this can be removed, and we can just use the + // context and the queue from the harness. + if (context == nullptr) + { + contextWrapper = + clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err); + test_error(err, "clCreateContext failed"); + context = contextWrapper; + } + + if (queue == nullptr) + { + queueWrapper = clCreateCommandQueue(context, device, 0, &err); + test_error(err, "clCreateCommandQueue failed"); + queue = queueWrapper; + } + + UnifiedSVMCornerCaseMigrateMem Test(context, device, queue, num_elements); + err = Test.setup(); + test_error(err, "test setup failed"); + + err = Test.run(); + test_error(err, "test failed"); + + return TEST_PASS; +} diff --git a/test_conformance/SVM/unified_svm_fixture.h b/test_conformance/SVM/unified_svm_fixture.h index 8797868d..a5b72330 100644 --- a/test_conformance/SVM/unified_svm_fixture.h +++ b/test_conformance/SVM/unified_svm_fixture.h @@ -84,11 +84,19 @@ public: if (caps & CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR) { - // For now, just unconditionally align to the device maximum - data = static_cast( - align_malloc(count * sizeof(T), deviceMaxAlignment)); - test_assert_error_ret(data != nullptr, "Failed to allocate memory", - CL_OUT_OF_RESOURCES); + if (count == 0) + { + data = nullptr; + } + else + { + // For now, just unconditionally align to the device maximum + data = static_cast( + align_malloc(count * sizeof(T), deviceMaxAlignment)); + test_assert_error_ret(data != nullptr, + "Failed to allocate memory", + CL_OUT_OF_RESOURCES); + } } else {