mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
add initial unified SVM capability tests (#2210)
These tests are passing on many devices using the unified SVM emulation layer. Specifically, adds tests for: * CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE_KHR * CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR * CL_SVM_CAPABILITY_HOST_READ_KHR * CL_SVM_CAPABILITY_HOST_WRITE_KHR * CL_SVM_CAPABILITY_HOST_MAP_KHR * CL_SVM_CAPABILITY_DEVICE_READ_KHR * CL_SVM_CAPABILITY_DEVICE_WRITE_KHR * CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS_KHR * CL_SVM_CAPABILITY_INDIRECT_ACCESS_KHR Still TODO: * CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR * CL_SVM_CAPABILITY_DEVICE_OWNED_KHR * CL_SVM_CAPABILITY_CONTEXT_ACCESS_KHR * CL_SVM_CAPABILITY_HOST_OWNED_KHR * CL_SVM_CAPABILITY_CONCURRENT_ACCESS_KHR * CL_SVM_CAPABILITY_CONCURRENT_ATOMIC_ACCESS_KHR
This commit is contained in:
Ben Ashbaugh
@@ -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")
|
||||
|
||||
751
test_conformance/SVM/test_unified_svm_capabilities.cpp
Normal file
751
test_conformance/SVM/test_unified_svm_capabilities.cpp
Normal file
@@ -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 <cinttypes>
|
||||
#include <memory>
|
||||
|
||||
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<cl_int>(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<cl_svm_alloc_properties_khr>(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<cl_int>(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<cl_int>(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<cl_int>(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<cl_int>(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<cl_int>(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<cl_int>(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<cl_int>(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<cl_uint>(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;
|
||||
}
|
||||
@@ -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<cl_svm_capabilities_khr> 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]);
|
||||
|
||||
372
test_conformance/SVM/unified_svm_fixture.h
Normal file
372
test_conformance/SVM/unified_svm_fixture.h
Normal file
@@ -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 <algorithm>
|
||||
#include <memory>
|
||||
|
||||
static inline void parseSVMAllocProperties(
|
||||
std::vector<cl_svm_alloc_properties_khr> 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<cl_device_id>(props[++i]);
|
||||
break;
|
||||
case CL_SVM_ALLOC_ACCESS_FLAGS_KHR:
|
||||
accessFlags =
|
||||
static_cast<cl_svm_alloc_access_flags_khr>(props[++i]);
|
||||
break;
|
||||
case CL_SVM_ALLOC_ALIGNMENT_KHR:
|
||||
alignment = static_cast<size_t>(props[++i]);
|
||||
break;
|
||||
default:
|
||||
log_error("Unknown SVM property: %X\n",
|
||||
static_cast<cl_uint>(props[i]));
|
||||
return;
|
||||
}
|
||||
++i;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T> 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<cl_svm_alloc_properties_khr> 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<T*>(
|
||||
align_malloc(count * sizeof(T), deviceMaxAlignment));
|
||||
test_assert_error_ret(data != nullptr, "Failed to allocate memory",
|
||||
CL_OUT_OF_RESOURCES);
|
||||
}
|
||||
else
|
||||
{
|
||||
std::vector<cl_svm_alloc_properties_khr> 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<cl_svm_alloc_properties_khr>(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<T>& 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<T>& 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 <typename T>
|
||||
std::unique_ptr<USVMWrapper<T>> get_usvm_wrapper(cl_uint typeIndex)
|
||||
{
|
||||
return std::unique_ptr<USVMWrapper<T>>(new USVMWrapper<T>(
|
||||
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<cl_svm_capabilities_khr> platformUSVMCaps;
|
||||
std::vector<cl_svm_capabilities_khr> deviceUSVMCaps;
|
||||
size_t deviceMaxAlignment = 0;
|
||||
|
||||
clSVMAllocWithPropertiesKHR_fn clSVMAllocWithPropertiesKHR = nullptr;
|
||||
clSVMFreeWithPropertiesKHR_fn clSVMFreeWithPropertiesKHR = nullptr;
|
||||
clGetSVMPointerInfoKHR_fn clGetSVMPointerInfoKHR = nullptr;
|
||||
clGetSVMSuggestedTypeIndexKHR_fn clGetSVMSuggestedTypeIndexKHR = nullptr;
|
||||
};
|
||||
Reference in New Issue
Block a user