diff --git a/test_conformance/SVM/CMakeLists.txt b/test_conformance/SVM/CMakeLists.txt index d16310ec34..2bda3dfe3d 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_capabilities.cpp ) set_gnulike_module_compile_flags("-Wno-sometimes-uninitialized -Wno-sign-compare") diff --git a/test_conformance/SVM/test_unified_svm_capabilities.cpp b/test_conformance/SVM/test_unified_svm_capabilities.cpp new file mode 100644 index 0000000000..9f1640ba51 --- /dev/null +++ b/test_conformance/SVM/test_unified_svm_capabilities.cpp @@ -0,0 +1,751 @@ +// +// 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 UnifiedSVMCapabilities : UnifiedSVMBase +{ + UnifiedSVMCapabilities(cl_context context, cl_device_id device, + cl_command_queue queue, int num_elements) + : UnifiedSVMBase(context, device, queue, num_elements) + {} + + cl_int test_CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE_KHR(cl_uint typeIndex) + { + cl_int err; + + if (!kernel_StorePointer) + { + err = createStorePointerKernel(); + test_error(err, "could not create StorePointer kernel"); + } + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate source memory"); + + clMemWrapper out = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_int*), nullptr, &err); + test_error(err, "could not create destination buffer"); + + err |= clSetKernelArgSVMPointer(kernel_StorePointer, 0, mem->get_ptr()); + err |= clSetKernelArg(kernel_StorePointer, 1, sizeof(out), &out); + test_error(err, "could not set kernel arguments"); + + 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"); + + cl_int* check = nullptr; + err = clEnqueueReadBuffer(queue, out, CL_TRUE, 0, sizeof(cl_int*), + &check, 0, nullptr, nullptr); + test_error(err, "could not read output buffer"); + + test_assert_error(check == mem->get_ptr(), + "stored pointer does not match input pointer"); + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR(cl_uint typeIndex) + { + const auto caps = deviceUSVMCaps[typeIndex]; + if (caps & CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR) + { + return CL_SUCCESS; + } + + cl_int err; + + void* ptr; + + ptr = clSVMAllocWithPropertiesKHR(context, nullptr, typeIndex, 1, &err); + test_error(err, "allocating without associated device failed"); + + err = clSVMFreeWithPropertiesKHR(context, nullptr, 0, ptr); + test_error(err, "freeing without associated device failed"); + + cl_svm_alloc_properties_khr props[] = { + CL_SVM_ALLOC_ASSOCIATED_DEVICE_HANDLE_KHR, + reinterpret_cast(device), 0 + }; + ptr = clSVMAllocWithPropertiesKHR(context, props, typeIndex, 1, &err); + test_error(err, "allocating with associated device failed"); + + err = clSVMFreeWithPropertiesKHR(context, nullptr, 0, ptr); + test_error(err, "freeing with associated device failed"); + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_HOST_READ_KHR(cl_uint typeIndex) + { + const auto caps = deviceUSVMCaps[typeIndex]; + cl_int err; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + cl_int value = genrand_int32(d); + err = mem->write(value); + test_error(err, "could not write to usvm memory"); + + cl_int check = mem->get_ptr()[0]; + test_assert_error(check == value, "read value does not match"); + + if (caps & CL_SVM_CAPABILITY_DEVICE_WRITE_KHR) + { + value = genrand_int32(d); + err = clEnqueueSVMMemcpy(queue, CL_TRUE, mem->get_ptr(), &value, + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not write to usvm memory on the device"); + + check = mem->get_ptr()[0]; + test_assert_error(check == value, "read value does not match"); + } + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_HOST_WRITE_KHR(cl_uint typeIndex) + { + const auto caps = deviceUSVMCaps[typeIndex]; + cl_int err; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + cl_int value = genrand_int32(d); + mem->get_ptr()[0] = value; + + cl_int check; + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + test_assert_error(check == value, "read value does not match"); + + if (caps & CL_SVM_CAPABILITY_DEVICE_READ_KHR) + { + value = genrand_int32(d); + mem->get_ptr()[0] = value; + + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &check, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not read from usvm memory on the device"); + test_assert_error(check == value, "read value does not match"); + } + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_HOST_MAP_KHR(cl_uint typeIndex) + { + const auto caps = deviceUSVMCaps[typeIndex]; + cl_int err; + + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + // map for writing, then map for reading + cl_int value = genrand_int32(d); + err = + clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, + mem->get_ptr(), sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for writing"); + + mem->get_ptr()[0] = value; + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for reading"); + + cl_int check = mem->get_ptr()[0]; + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + test_assert_error(check == value, "read value does not match"); + + // write directly on the host, map for reading on the host + if (caps & CL_SVM_CAPABILITY_HOST_WRITE_KHR) + { + value = genrand_int32(d); + mem->get_ptr()[0] = value; + + err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for reading"); + + check = mem->get_ptr()[0]; + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + test_assert_error(check == value, "read value does not match"); + } + + // map for writing on the host, read directly on the host + if (caps & CL_SVM_CAPABILITY_HOST_READ_KHR) + { + value = genrand_int32(d); + err = clEnqueueSVMMap( + queue, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for writing"); + + mem->get_ptr()[0] = value; + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + check = mem->get_ptr()[0]; + test_assert_error(check == value, "read value does not match"); + } + + // write on the device, map for reading on the host + if (caps & CL_SVM_CAPABILITY_DEVICE_WRITE_KHR) + { + value = genrand_int32(d); + err = clEnqueueSVMMemcpy(queue, CL_TRUE, mem->get_ptr(), &value, + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not write to usvm memory on the device"); + + err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for reading"); + + check = mem->get_ptr()[0]; + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + test_assert_error(check == value, "read value does not match"); + } + + // map for writing on the host, read on the device + if (caps & CL_SVM_CAPABILITY_DEVICE_READ_KHR) + { + cl_int value = genrand_int32(d); + err = clEnqueueSVMMap( + queue, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not map usvm memory for writing"); + + mem->get_ptr()[0] = value; + + err = clEnqueueSVMUnmap(queue, mem->get_ptr(), 0, nullptr, nullptr); + test_error(err, "could not unmap usvm memory"); + + cl_int check; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &check, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not read from usvm memory on the device"); + + test_assert_error(check == value, "read value does not match"); + } + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_DEVICE_READ_KHR(cl_uint typeIndex) + { + cl_int err; + + // setup + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + if (!kernel_CopyMemory) + { + err = createCopyMemoryKernel(); + test_error(err, "could not create CopyMemory kernel"); + } + + // test reading via memcpy: + cl_int value = genrand_int32(d); + err = mem->write(value); + test_error(err, "could not write to usvm memory"); + + cl_int check; + err = clEnqueueSVMMemcpy(queue, CL_TRUE, &check, mem->get_ptr(), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not read from usvm memory with memcpy"); + + test_assert_error(check == value, + "read value with memcpy does not match"); + + // test reading via kernel + value = genrand_int32(d); + err = mem->write(value); + test_error(err, "could not write to usvm memory"); + + clMemWrapper out = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_int), nullptr, &err); + test_error(err, "could not create output buffer"); + + err |= clSetKernelArgSVMPointer(kernel_CopyMemory, 0, mem->get_ptr()); + err |= clSetKernelArg(kernel_CopyMemory, 1, sizeof(out), &out); + test_error(err, "could not set kernel arguments"); + + size_t global_work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel_CopyMemory, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, + nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + err = clEnqueueReadBuffer(queue, out, CL_TRUE, 0, sizeof(cl_int), + &check, 0, nullptr, nullptr); + test_error(err, "could not read output buffer"); + + test_assert_error(check == value, + "read value with kernel does not match"); + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_DEVICE_WRITE_KHR(cl_uint typeIndex) + { + cl_int err; + + // setup + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + if (!kernel_CopyMemory) + { + err = createCopyMemoryKernel(); + test_error(err, "could not create CopyMemory kernel"); + } + + // test writing via memfill + cl_int value = genrand_int32(d); + err = clEnqueueSVMMemFill(queue, mem->get_ptr(), &value, sizeof(value), + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not write to usvm memory with memfill"); + + cl_int check; + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + + test_assert_error(check == value, + "read value with memfill does not match"); + + // test writing via memcpy + value = genrand_int32(d); + err = clEnqueueSVMMemcpy(queue, CL_TRUE, mem->get_ptr(), &value, + sizeof(value), 0, nullptr, nullptr); + test_error(err, "could not write to usvm memory with memcpy"); + + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + + test_assert_error(check == value, + "read value with memcpy does not match"); + + // test writing via kernel + value = genrand_int32(d); + clMemWrapper in = + clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + sizeof(cl_int), &value, &err); + test_error(err, "could not create input buffer"); + + err |= clSetKernelArg(kernel_CopyMemory, 0, sizeof(in), &in); + err |= clSetKernelArgSVMPointer(kernel_CopyMemory, 1, mem->get_ptr()); + test_error(err, "could not set kernel arguments"); + + size_t global_work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel_CopyMemory, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, + nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + + test_assert_error(check == value, + "read value with kernel does not match"); + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS_KHR(cl_uint typeIndex) + { + cl_int err; + + // setup + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + if (!kernel_AtomicIncrement) + { + err = createAtomicIncrementKernel(); + test_error(err, "could not create AtomicIncrement kernel"); + } + + err = mem->write(0); + test_error(err, "could not write to usvm memory"); + + err = + clSetKernelArgSVMPointer(kernel_AtomicIncrement, 0, mem->get_ptr()); + test_error(err, "could not set kernel arguments"); + + size_t global_work_size = num_elements; + err = clEnqueueNDRangeKernel(queue, kernel_AtomicIncrement, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, + nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + cl_int check; + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + + test_assert_error(check == num_elements, + "read value does not match expected value"); + + return CL_SUCCESS; + } + + cl_int test_CL_SVM_CAPABILITY_INDIRECT_ACCESS_KHR(cl_uint typeIndex) + { + cl_int err; + + // setup + auto mem = get_usvm_wrapper(typeIndex); + err = mem->allocate(1); + test_error(err, "could not allocate usvm memory"); + + if (!kernel_IndirectAccessRead) + { + err = createIndirectAccessKernel(); + test_error(err, "could not create IndirectAccess kernel"); + } + + // test reading indirectly + cl_int value = genrand_int32(d); + err = mem->write(value); + test_error(err, "could not write to usvm memory"); + + auto ptr = mem->get_ptr(); + clMemWrapper indirect = + clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + sizeof(ptr), &ptr, &err); + test_error(err, "could not create indirect buffer"); + + clMemWrapper direct = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_int), nullptr, &err); + test_error(err, "could not create direct buffer"); + + err |= clSetKernelArg(kernel_IndirectAccessRead, 0, sizeof(indirect), + &indirect); + err |= clSetKernelArg(kernel_IndirectAccessRead, 1, sizeof(direct), + &direct); + test_error(err, "could not set kernel arguments"); + + cl_bool enable = CL_TRUE; + err = clSetKernelExecInfo(kernel_IndirectAccessRead, + CL_KERNEL_EXEC_INFO_SVM_INDIRECT_ACCESS_KHR, + sizeof(enable), &enable); + test_error(err, "could not enable indirect access"); + + size_t global_work_size = 1; + err = clEnqueueNDRangeKernel(queue, kernel_IndirectAccessRead, 1, + nullptr, &global_work_size, nullptr, 0, + nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + cl_int check; + err = clEnqueueReadBuffer(queue, direct, CL_TRUE, 0, sizeof(cl_int), + &check, 0, nullptr, nullptr); + test_error(err, "could not read direct buffer"); + + test_assert_error(check == value, "read value does not match"); + + // test writing indirectly + value = genrand_int32(d); + err = clEnqueueWriteBuffer(queue, direct, CL_TRUE, 0, sizeof(cl_int), + &value, 0, nullptr, nullptr); + test_error(err, "could not write to direct buffer"); + + err |= clSetKernelArg(kernel_IndirectAccessWrite, 0, sizeof(indirect), + &indirect); + err |= clSetKernelArg(kernel_IndirectAccessWrite, 1, sizeof(direct), + &direct); + test_error(err, "could not set kernel arguments"); + + err = clSetKernelExecInfo(kernel_IndirectAccessWrite, + CL_KERNEL_EXEC_INFO_SVM_INDIRECT_ACCESS_KHR, + sizeof(enable), &enable); + test_error(err, "could not enable indirect access"); + + err = clEnqueueNDRangeKernel(queue, kernel_IndirectAccessWrite, 1, + nullptr, &global_work_size, nullptr, 0, + nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clFinish(queue); + test_error(err, "clFinish failed"); + + err = mem->read(check); + test_error(err, "could not read from usvm memory"); + + test_assert_error(check == value, "read value does not match"); + + return CL_SUCCESS; + } + + cl_int run() override + { + cl_int err; + for (cl_uint ti = 0; ti < static_cast(deviceUSVMCaps.size()); + ti++) + { + const auto caps = deviceUSVMCaps[ti]; + log_info(" testing SVM type %u\n", ti); + + if (caps & CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE_KHR) + { + log_info( + " testing CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE\n"); + err = test_CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE_KHR(ti); + test_error(err, + "CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE failed"); + } + // CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR + // CL_SVM_CAPABILITY_DEVICE_OWNED_KHR + if (caps & CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR) + { + log_info( + " testing CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED\n"); + err = test_CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED failed"); + } + // CL_SVM_CAPABILITY_CONTEXT_ACCESS_KHR + // CL_SVM_CAPABILITY_HOST_OWNED_KHR + if (caps & CL_SVM_CAPABILITY_HOST_READ_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_HOST_READ\n"); + err = test_CL_SVM_CAPABILITY_HOST_READ_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_HOST_READ failed"); + } + if (caps & CL_SVM_CAPABILITY_HOST_WRITE_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_HOST_WRITE\n"); + err = test_CL_SVM_CAPABILITY_HOST_WRITE_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_HOST_WRITE failed"); + } + if (caps & CL_SVM_CAPABILITY_HOST_MAP_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_HOST_MAP\n"); + err = test_CL_SVM_CAPABILITY_HOST_MAP_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_HOST_MAP failed"); + } + if (caps & CL_SVM_CAPABILITY_DEVICE_READ_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_DEVICE_READ\n"); + err = test_CL_SVM_CAPABILITY_DEVICE_READ_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_DEVICE_READ failed"); + } + if (caps & CL_SVM_CAPABILITY_DEVICE_WRITE_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_DEVICE_WRITE\n"); + err = test_CL_SVM_CAPABILITY_DEVICE_READ_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_DEVICE_READ failed"); + } + if (caps & CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS_KHR) + { + log_info( + " testing CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS\n"); + err = test_CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS_KHR(ti); + test_error(err, + "CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS failed"); + } + // CL_SVM_CAPABILITY_CONCURRENT_ACCESS_KHR + // CL_SVM_CAPABILITY_CONCURRENT_ATOMIC_ACCESS_KHR + if (caps & CL_SVM_CAPABILITY_INDIRECT_ACCESS_KHR) + { + log_info(" testing CL_SVM_CAPABILITY_INDIRECT_ACCESS\n"); + err = test_CL_SVM_CAPABILITY_INDIRECT_ACCESS_KHR(ti); + test_error(err, "CL_SVM_CAPABILITY_INDIRECT_ACCESS failed"); + } + } + return CL_SUCCESS; + } + + cl_int createStorePointerKernel() + { + cl_int err; + + 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"); + + return CL_SUCCESS; + } + + cl_int createCopyMemoryKernel() + { + cl_int err; + + const char* programString = R"( + kernel void test_CopyMemory(const global int* src, global int* dst) + { + dst[get_global_id(0)] = src[get_global_id(0)]; + } + )"; + + clProgramWrapper program; + err = create_single_kernel_helper(context, &program, &kernel_CopyMemory, + 1, &programString, "test_CopyMemory"); + test_error(err, "could not create CopyMemory kernel"); + + return CL_SUCCESS; + } + + cl_int createAtomicIncrementKernel() + { + cl_int err; + + const char* programString = R"( + kernel void test_AtomicIncrement(global int* ptr) + { + atomic_inc(ptr); + } + )"; + + clProgramWrapper program; + err = create_single_kernel_helper( + context, &program, &kernel_AtomicIncrement, 1, &programString, + "test_AtomicIncrement"); + test_error(err, "could not create AtomicIncrement kernel"); + + return CL_SUCCESS; + } + + cl_int createIndirectAccessKernel() + { + cl_int err; + + const char* programString = R"( + struct s { const global int* ptr; }; + kernel void test_IndirectAccessRead(const global struct s* src, global int* dst) + { + dst[get_global_id(0)] = src->ptr[get_global_id(0)]; + } + + struct d { global int* ptr; }; + kernel void test_IndirectAccessWrite(global struct d* dst, const global int* src) + { + dst->ptr[get_global_id(0)] = src[get_global_id(0)]; + } + )"; + + clProgramWrapper program; + err = create_single_kernel_helper( + context, &program, &kernel_IndirectAccessRead, 1, &programString, + "test_IndirectAccessRead"); + test_error(err, "could not create IndirectAccessRead kernel"); + + kernel_IndirectAccessWrite = + clCreateKernel(program, "test_IndirectAccessWrite", &err); + test_error(err, "could not create IndirectAccessWrite kernel"); + + return CL_SUCCESS; + } + + clKernelWrapper kernel_StorePointer; + clKernelWrapper kernel_CopyMemory; + clKernelWrapper kernel_AtomicIncrement; + clKernelWrapper kernel_IndirectAccessRead; + clKernelWrapper kernel_IndirectAccessWrite; +}; + +REGISTER_TEST(unified_svm_capabilities) +{ + 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; + } + + UnifiedSVMCapabilities 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/test_unified_svm_consistency.cpp b/test_conformance/SVM/test_unified_svm_consistency.cpp index 40bde9147e..000147919b 100644 --- a/test_conformance/SVM/test_unified_svm_consistency.cpp +++ b/test_conformance/SVM/test_unified_svm_consistency.cpp @@ -19,7 +19,7 @@ REGISTER_TEST(unified_svm_consistency) { - if (!is_extension_available(deviceID, "cl_khr_unified_svm")) + 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; @@ -28,7 +28,7 @@ REGISTER_TEST(unified_svm_consistency) cl_int err; cl_platform_id platformID; - err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), + err = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), (void *)(&platformID), nullptr); test_error(err, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM"); @@ -122,7 +122,7 @@ REGISTER_TEST(unified_svm_consistency) } if (platformCapabilities[i] != check) { - test_fail("Platform SVM type capabilities at index %zu: 0x%" PRIx64 + test_fail("Platform SVM type capabilities at index %d: 0x%" PRIx64 " do not match the intersection of device capabilities " "0x%" PRIx64 ".\n", i, platformCapabilities[i], check); @@ -135,7 +135,7 @@ REGISTER_TEST(unified_svm_consistency) // supported. std::vector deviceCapabilities(capabilityCount); - err = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_TYPE_CAPABILITIES_KHR, + err = clGetDeviceInfo(device, CL_DEVICE_SVM_TYPE_CAPABILITIES_KHR, platformSize, deviceCapabilities.data(), nullptr); test_error(err, "clGetDeviceInfo failed for CL_DEVICE_SVM_CAPABILITIES_KHR"); @@ -148,7 +148,7 @@ REGISTER_TEST(unified_svm_consistency) if (!consistent) { test_fail( - "Device SVM type capabilities at index %zu: 0x%" PRIx64 + "Device SVM type capabilities at index %d: 0x%" PRIx64 " are not consistent with platform SVM type capabilities: " "0x%" PRIx64 ".\n", i, deviceCapabilities[i], platformCapabilities[i]); diff --git a/test_conformance/SVM/unified_svm_fixture.h b/test_conformance/SVM/unified_svm_fixture.h new file mode 100644 index 0000000000..feaa54db62 --- /dev/null +++ b/test_conformance/SVM/unified_svm_fixture.h @@ -0,0 +1,372 @@ +// +// 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 "common.h" + +#include +#include + +static inline void parseSVMAllocProperties( + std::vector props, cl_device_id& device, + cl_svm_alloc_access_flags_khr& accessFlags, size_t& alignment) +{ + device = nullptr; + accessFlags = 0; + alignment = 0; + + if (!props.empty()) + { + size_t i = 0; + while (props[i]) + { + switch (props[i]) + { + case CL_SVM_ALLOC_ASSOCIATED_DEVICE_HANDLE_KHR: + device = reinterpret_cast(props[++i]); + break; + case CL_SVM_ALLOC_ACCESS_FLAGS_KHR: + accessFlags = + static_cast(props[++i]); + break; + case CL_SVM_ALLOC_ALIGNMENT_KHR: + alignment = static_cast(props[++i]); + break; + default: + log_error("Unknown SVM property: %X\n", + static_cast(props[i])); + return; + } + ++i; + } + } +} + +template class USVMWrapper { +public: + USVMWrapper(cl_context context_, cl_device_id device_, + cl_command_queue queue_, cl_uint typeIndex_, + cl_svm_capabilities_khr caps_, size_t deviceMaxAlignment_, + clSVMAllocWithPropertiesKHR_fn clSVMAllocWithPropertiesKHR_, + clSVMFreeWithPropertiesKHR_fn clSVMFreeWithPropertiesKHR_, + clGetSVMPointerInfoKHR_fn clGetSVMPointerInfoKHR_, + clGetSVMSuggestedTypeIndexKHR_fn clGetSVMSuggestedTypeIndexKHR_) + : context(context_), device(device_), queue(queue_), + typeIndex(typeIndex_), caps(caps_), + deviceMaxAlignment(deviceMaxAlignment_), + clSVMAllocWithPropertiesKHR(clSVMAllocWithPropertiesKHR_), + clSVMFreeWithPropertiesKHR(clSVMFreeWithPropertiesKHR_), + clGetSVMPointerInfoKHR(clGetSVMPointerInfoKHR_), + clGetSVMSuggestedTypeIndexKHR(clGetSVMSuggestedTypeIndexKHR_) + {} + + ~USVMWrapper() { free(); } + + cl_int allocate(const size_t count, + const std::vector props_ = {}) + { + if (data != nullptr) + { + free(); + } + + 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); + } + else + { + std::vector props = props_; + if (!props.empty()) + { + props.pop_back(); + } + if (!(caps & CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR) + && std::find(props.begin(), props.end(), + CL_SVM_ALLOC_ASSOCIATED_DEVICE_HANDLE_KHR) + == props.end()) + { + props.push_back(CL_SVM_ALLOC_ASSOCIATED_DEVICE_HANDLE_KHR); + props.push_back( + reinterpret_cast(device)); + } + if (!props.empty() || !props_.empty()) + { + props.push_back(0); + } + + cl_int err; + data = (T*)clSVMAllocWithPropertiesKHR( + context, props.empty() ? nullptr : props.data(), typeIndex, + count * sizeof(T), &err); + test_error(err, "clSVMAllocWithPropertiesKHR failed"); + } + + return CL_SUCCESS; + } + + cl_int free() + { + if (data) + { + if (caps & CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR) + { + align_free(data); + } + else + { + cl_int err; + err = clSVMFreeWithPropertiesKHR(context, nullptr, 0, data); + test_error(err, "clSVMFreeWithPropertiesKHR failed"); + } + + data = nullptr; + } + + return CL_SUCCESS; + } + + cl_int write(const T* source, size_t count, size_t offset = 0) + { + if (data == nullptr) + { + return CL_INVALID_OPERATION; + } + + cl_int err; + + if (caps & CL_SVM_CAPABILITY_HOST_WRITE_KHR) + { + std::copy(source, source + count, data + offset); + } + else if (caps & CL_SVM_CAPABILITY_HOST_MAP_KHR) + { + err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, data, + count * sizeof(T), 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMMap failed"); + + std::copy(source, source + count, data + offset); + + err = clEnqueueSVMUnmap(queue, data, 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMUnmap failed"); + } + else if (caps & CL_SVM_CAPABILITY_DEVICE_WRITE_KHR) + { + err = clEnqueueSVMMemcpy(queue, CL_TRUE, data + offset, source, + count * sizeof(T), 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMMemcpy failed"); + } + else + { + log_error("Not sure how to write to SVM type index %u!\n", + typeIndex); + return CL_INVALID_OPERATION; + } + + return CL_SUCCESS; + } + + cl_int write(const std::vector& source, size_t offset = 0) + { + return write(source.data(), source.size(), offset); + } + + cl_int write(T source, size_t offset = 0) + { + return write(&source, 1, offset); + } + + cl_int read(T* dst, size_t count, size_t offset = 0) + { + if (data == nullptr) + { + return CL_INVALID_OPERATION; + } + + cl_int err; + + if (caps & CL_SVM_CAPABILITY_HOST_READ_KHR) + { + std::copy(data + offset, data + offset + count, dst); + } + else if (caps & CL_SVM_CAPABILITY_HOST_MAP_KHR) + { + err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, data, + count * sizeof(T), 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMMap failed"); + + std::copy(data + offset, data + offset + count, dst); + + err = clEnqueueSVMUnmap(queue, data, 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMUnmap failed"); + } + else if (caps & CL_SVM_CAPABILITY_DEVICE_READ_KHR) + { + err = clEnqueueSVMMemcpy(queue, CL_TRUE, dst, data + offset, + count * sizeof(T), 0, nullptr, nullptr); + test_error(err, "clEnqueueSVMMemcpy failed"); + } + else + { + log_error("Not sure how to read from SVM type index %u!\n", + typeIndex); + return CL_INVALID_OPERATION; + } + + return CL_SUCCESS; + } + + cl_int read(std::vector& dst, size_t offset = 0) + { + return read(dst.data(), dst.size(), offset); + } + + cl_int read(T& dst, size_t offset = 0) { return read(&dst, 1, offset); } + + T* get_ptr() { return data; } + +private: + cl_context context = nullptr; + cl_device_id device = nullptr; + cl_command_queue queue = nullptr; + cl_uint typeIndex = 0; + cl_svm_capabilities_khr caps = 0; + size_t deviceMaxAlignment = 0; + + clSVMAllocWithPropertiesKHR_fn clSVMAllocWithPropertiesKHR = nullptr; + clSVMFreeWithPropertiesKHR_fn clSVMFreeWithPropertiesKHR = nullptr; + clGetSVMPointerInfoKHR_fn clGetSVMPointerInfoKHR = nullptr; + clGetSVMSuggestedTypeIndexKHR_fn clGetSVMSuggestedTypeIndexKHR = nullptr; + + T* data = nullptr; +}; + +struct UnifiedSVMBase +{ + UnifiedSVMBase(cl_context context_, cl_device_id device_, + cl_command_queue queue_, int num_elements_) + : d(gRandomSeed), context(context_), device(device_), queue(queue_), + num_elements(num_elements_) + {} + + virtual cl_int setup() + { + cl_int err; + + cl_platform_id platform{}; + err = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, + sizeof(cl_platform_id), &platform, nullptr); + test_error(err, "clGetDeviceInfo failed for CL_DEVICE_PLATFORM"); + + size_t sz{}; + err = clGetPlatformInfo(platform, CL_PLATFORM_SVM_TYPE_CAPABILITIES_KHR, + 0, nullptr, &sz); + test_error(err, + "clGetPlatformInfo failed for " + "CL_PLATFORM_SVM_TYPE_CAPABILITIES_KHR size"); + + platformUSVMCaps.resize(sz / sizeof(cl_svm_capabilities_khr)); + err = clGetPlatformInfo(platform, CL_PLATFORM_SVM_TYPE_CAPABILITIES_KHR, + sz, platformUSVMCaps.data(), nullptr); + test_error(err, + "clGetPlatformInfo failed for " + "CL_PLATFORM_SVM_TYPE_CAPABILITIES_KHR data"); + + err = clGetDeviceInfo(device, CL_DEVICE_SVM_TYPE_CAPABILITIES_KHR, 0, + nullptr, &sz); + test_error( + err, + "clGetDeviceInfo failed for CL_DEVICE_SVM_CAPABILITIES_KHR size"); + + deviceUSVMCaps.resize(sz / sizeof(cl_svm_capabilities_khr)); + err = clGetDeviceInfo(device, CL_DEVICE_SVM_TYPE_CAPABILITIES_KHR, sz, + deviceUSVMCaps.data(), nullptr); + test_error( + err, + "clGetDeviceInfo failed for CL_DEVICE_SVM_CAPABILITIES_KHR data"); + + clSVMAllocWithPropertiesKHR = (clSVMAllocWithPropertiesKHR_fn) + clGetExtensionFunctionAddressForPlatform( + platform, "clSVMAllocWithPropertiesKHR"); + test_assert_error_ret(clSVMAllocWithPropertiesKHR != nullptr, + "clSVMAllocWithPropertiesKHR not found", + CL_INVALID_OPERATION); + + clSVMFreeWithPropertiesKHR = (clSVMFreeWithPropertiesKHR_fn) + clGetExtensionFunctionAddressForPlatform( + platform, "clSVMFreeWithPropertiesKHR"); + test_assert_error_ret(clSVMFreeWithPropertiesKHR != nullptr, + "clSVMFreeWithPropertiesKHR not found", + CL_INVALID_OPERATION); + + clGetSVMPointerInfoKHR = + (clGetSVMPointerInfoKHR_fn)clGetExtensionFunctionAddressForPlatform( + platform, "clGetSVMPointerInfoKHR"); + test_assert_error_ret(clGetSVMPointerInfoKHR != nullptr, + "clGetSVMPointerInfoKHR not found", + CL_INVALID_OPERATION); + + clGetSVMSuggestedTypeIndexKHR = (clGetSVMSuggestedTypeIndexKHR_fn) + clGetExtensionFunctionAddressForPlatform( + platform, "clGetSVMSuggestedTypeIndexKHR"); + test_assert_error_ret(clGetSVMSuggestedTypeIndexKHR != nullptr, + "clGetSVMSuggestedTypeIndexKHR not found", + CL_INVALID_OPERATION); + + // The maximum supported alignment is equal to the size of the largest + // data type supported by the device + if (gHasLong || is_extension_available(device, "cl_khr_fp64")) + { + deviceMaxAlignment = 16 * sizeof(cl_long); + } + else + { + deviceMaxAlignment = 16 * sizeof(cl_int); + } + + return CL_SUCCESS; + } + + virtual cl_int run() = 0; + + template + std::unique_ptr> get_usvm_wrapper(cl_uint typeIndex) + { + return std::unique_ptr>(new USVMWrapper( + context, device, queue, typeIndex, deviceUSVMCaps[typeIndex], + deviceMaxAlignment, clSVMAllocWithPropertiesKHR, + clSVMFreeWithPropertiesKHR, clGetSVMPointerInfoKHR, + clGetSVMSuggestedTypeIndexKHR)); + } + + MTdataHolder d; + cl_context context = nullptr; + cl_device_id device = nullptr; + cl_command_queue queue = nullptr; + int num_elements = 0; + + std::vector platformUSVMCaps; + std::vector deviceUSVMCaps; + size_t deviceMaxAlignment = 0; + + clSVMAllocWithPropertiesKHR_fn clSVMAllocWithPropertiesKHR = nullptr; + clSVMFreeWithPropertiesKHR_fn clSVMFreeWithPropertiesKHR = nullptr; + clGetSVMPointerInfoKHR_fn clGetSVMPointerInfoKHR = nullptr; + clGetSVMSuggestedTypeIndexKHR_fn clGetSVMSuggestedTypeIndexKHR = nullptr; +};