Google Git
Sign in
android/platform/external/clvk/refs/heads/android17-release/./src/api.cpp
blob: dab5fc90cbe72a5dcc279d83c17e7c845eded84c [file] [edit]
// Copyright 2018 The clvk authors.
//
// 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 "cl_headers.hpp"
#include "icd.hpp"
#include "image_format.hpp"
#include "init.hpp"
#include "kernel.hpp"
#include "log.hpp"
#include "memory.hpp"
#include "objects.hpp"
#include "program.hpp"
#include "queue.hpp"
#include "semaphore.hpp"
#include "tracing.hpp"
#define LOG_API_CALL(fmt, ...) \
cvk_debug_group_fn(loggroup::api, fmt, __VA_ARGS__)
#define CLVK_API_CALL CL_API_CALL
namespace {
// Validation functions
bool is_valid_platform(cl_platform_id platform) {
return platform != nullptr && icd_downcast(platform)->is_valid();
}
bool is_valid_device(cl_device_id device) {
return device != nullptr && icd_downcast(device)->is_valid();
}
bool is_valid_context(cl_context context) {
return context != nullptr && icd_downcast(context)->is_valid();
}
bool is_valid_program(cl_program program) {
return program != nullptr && icd_downcast(program)->is_valid();
}
bool is_valid_kernel(cl_kernel kernel) {
return kernel != nullptr && icd_downcast(kernel)->is_valid();
}
bool is_valid_sampler(cl_sampler sampler) {
return sampler != nullptr && icd_downcast(sampler)->is_valid();
}
bool is_valid_mem_object(cl_mem mem) {
return mem != nullptr && icd_downcast(mem)->is_valid();
}
bool is_valid_buffer(cl_mem mem) {
return is_valid_mem_object(mem) && icd_downcast(mem)->is_buffer_type();
}
bool is_valid_image(cl_mem mem) {
return is_valid_mem_object(mem) && icd_downcast(mem)->is_image_type();
}
bool is_valid_command_queue(cl_command_queue queue) {
return queue != nullptr && icd_downcast(queue)->is_valid();
}
bool is_valid_event(cl_event event) {
return event != nullptr && icd_downcast(event)->is_valid();
}
bool is_valid_semaphore(cl_semaphore_khr sem) {
return sem != nullptr && icd_downcast(sem)->is_valid();
}
bool is_valid_event_wait_list(cl_uint num_events_in_wait_list,
const cl_event* event_wait_list) {
if (((num_events_in_wait_list > 0) && (event_wait_list == nullptr)) ||
((num_events_in_wait_list == 0) && (event_wait_list != nullptr))) {
return false;
}
for (cl_uint i = 0; i < num_events_in_wait_list; i++) {
if (!is_valid_event(event_wait_list[i])) {
return false;
}
}
return true;
}
bool is_same_context(cl_command_queue queue, cl_mem mem) {
return icd_downcast(queue)->context() == icd_downcast(mem)->context();
}
bool is_same_context(cl_command_queue queue, cl_kernel kernel) {
return icd_downcast(queue)->context() == icd_downcast(kernel)->context();
}
bool is_same_context(cl_command_queue queue, cl_uint num_events,
const cl_event* event_list) {
for (cl_uint i = 0; i < num_events; i++) {
if (icd_downcast(queue)->context() !=
icd_downcast(event_list[i])->context()) {
return false;
}
}
return true;
}
bool is_same_context(cl_command_queue queue, cl_uint num_semas,
const cl_semaphore_khr* semas) {
for (cl_uint i = 0; i < num_semas; i++) {
if (icd_downcast(queue)->context() !=
icd_downcast(semas[i])->context()) {
return false;
}
}
return true;
}
bool is_valid_device_type(cl_device_type type) {
return (type < (CL_DEVICE_TYPE_CUSTOM << 1) && type != 0) ||
(type == CL_DEVICE_TYPE_ALL);
}
bool map_flags_are_valid(cl_map_flags flags) {
if ((flags & CL_MAP_WRITE_INVALIDATE_REGION) &&
(flags & (CL_MAP_READ | CL_MAP_WRITE))) {
return false;
}
return true;
}
bool is_compiler_available(cl_uint num_devices, const cl_device_id* devices) {
for (cl_uint i = 0; i < num_devices; i++) {
auto dev = icd_downcast(devices[i]);
if (!dev->compiler_available()) {
return false;
}
}
return true;
}
// Utilities
struct api_query_string : public std::string {
api_query_string() : std::string() {}
api_query_string(const char* init) : std::string(init) {}
api_query_string(const std::string& other) : std::string(other) {}
size_t size_with_null() const { return size() + 1; }
};
} // namespace
// Platform API
cl_int cvk_get_platform_ids(const clvk_global_state* state, cl_uint num_entries,
cl_platform_id* platforms, cl_uint* num_platforms) {
if ((num_platforms == nullptr) && (platforms == nullptr)) {
return CL_INVALID_VALUE;
}
if ((num_entries == 0) && (platforms != nullptr)) {
return CL_INVALID_VALUE;
}
if (platforms != nullptr) {
platforms[0] = state->platform();
}
if (num_platforms != nullptr) {
*num_platforms = 1;
}
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetPlatformIDs(cl_uint num_entries,
cl_platform_id* platforms,
cl_uint* num_platforms) {
auto state = get_or_init_global_state();
TRACE_FUNCTION("num_entries", num_entries);
LOG_API_CALL("num_entries = %u, platforms = %p, num_platforms = %p",
num_entries, platforms, num_platforms);
return cvk_get_platform_ids(state, num_entries, platforms, num_platforms);
}
cl_int CLVK_API_CALL clGetPlatformInfo(cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
auto state = get_or_init_global_state();
TRACE_FUNCTION("platform", (uintptr_t)platform, "param_name", param_name);
LOG_API_CALL("platform = %p, param_name = %u, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
platform, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t size_ret = 0;
const void* copy_ptr = nullptr;
cl_version val_version;
api_query_string val_string;
cl_ulong val_ulong;
if (!is_valid_platform(platform)) {
return CL_INVALID_PLATFORM;
}
const cvk_platform* plat = state->platform();
if (platform != nullptr) {
plat = icd_downcast(platform);
}
switch (param_name) {
case CL_PLATFORM_NAME:
val_string = plat->name();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_VERSION:
val_string = plat->version_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_VENDOR:
val_string = plat->vendor();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_PROFILE:
val_string = plat->profile();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_EXTENSIONS:
val_string = plat->extension_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_ICD_SUFFIX_KHR:
val_string = plat->icd_suffix();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_PLATFORM_NUMERIC_VERSION:
val_version = plat->version();
copy_ptr = &val_version;
size_ret = sizeof(val_version);
break;
case CL_PLATFORM_EXTENSIONS_WITH_VERSION:
copy_ptr = plat->extensions().data();
size_ret = plat->extensions().size() * sizeof(cl_name_version);
break;
case CL_PLATFORM_HOST_TIMER_RESOLUTION:
val_ulong = plat->host_timer_resolution();
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (param_value_size < size_ret)) {
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
memcpy(param_value, copy_ptr, std::min(param_value_size, size_ret));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = size_ret;
}
return ret;
}
static const std::unordered_map<std::string, void*> gExtensionEntrypoints = {
#define FUNC_PTR(X) reinterpret_cast<void*>(X)
#define EXTENSION_ENTRYPOINT(X) \
{ #X, FUNC_PTR(X) }
EXTENSION_ENTRYPOINT(clCreateProgramWithILKHR),
EXTENSION_ENTRYPOINT(clIcdGetPlatformIDsKHR),
EXTENSION_ENTRYPOINT(clCreateCommandQueueWithPropertiesKHR),
EXTENSION_ENTRYPOINT(clGetKernelSuggestedLocalWorkSizeKHR),
{"clGetKernelSubGroupInfoKHR", FUNC_PTR(clGetKernelSubGroupInfo)},
EXTENSION_ENTRYPOINT(clCreateSemaphoreWithPropertiesKHR),
EXTENSION_ENTRYPOINT(clEnqueueWaitSemaphoresKHR),
EXTENSION_ENTRYPOINT(clEnqueueSignalSemaphoresKHR),
EXTENSION_ENTRYPOINT(clGetSemaphoreInfoKHR),
EXTENSION_ENTRYPOINT(clRetainSemaphoreKHR),
EXTENSION_ENTRYPOINT(clReleaseSemaphoreKHR),
#undef EXTENSION_ENTRYPOINT
#undef FUNC_PTR
};
void* cvk_get_extension_function_pointer(const char* funcname) {
if (gExtensionEntrypoints.find(funcname) != gExtensionEntrypoints.end()) {
return gExtensionEntrypoints.at(funcname);
} else {
return nullptr;
}
}
void* CLVK_API_CALL clGetExtensionFunctionAddressForPlatform(
cl_platform_id platform, const char* funcname) {
TRACE_FUNCTION("platform", (uintptr_t)platform);
LOG_API_CALL("platform = %p, funcname = '%s'", platform, funcname);
if (platform == nullptr) {
return nullptr;
}
return cvk_get_extension_function_pointer(funcname);
}
void* CLVK_API_CALL clGetExtensionFunctionAddress(const char* funcname) {
TRACE_FUNCTION();
LOG_API_CALL("funcname = '%s'", funcname);
return cvk_get_extension_function_pointer(funcname);
}
// Device APIs
cl_int CLVK_API_CALL clGetDeviceIDs(cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries, cl_device_id* devices,
cl_uint* num_devices) {
auto state = get_or_init_global_state();
TRACE_FUNCTION("platform", (uintptr_t)platform, "device_type",
TRACE_STRING(cl_device_type_to_string(device_type)),
"num_entries", num_entries);
LOG_API_CALL(
"platform = %p, device_type = %lu (%s), num_entries = %u, devices "
"= %p, num_devices = %p",
platform, device_type, cl_device_type_to_string(device_type),
num_entries, devices, num_devices);
if (platform == nullptr) {
platform = state->platform();
} else if (platform != state->platform()) {
return CL_INVALID_PLATFORM;
}
if ((num_entries == 0) && (devices != nullptr)) {
return CL_INVALID_VALUE;
}
if ((num_devices == nullptr) && (devices == nullptr)) {
return CL_INVALID_VALUE;
}
if (!is_valid_device_type(device_type)) {
return CL_INVALID_DEVICE_TYPE;
}
cl_uint num = 0;
for (auto dev : icd_downcast(platform)->devices()) {
if ((dev->type() & device_type) ||
(device_type == CL_DEVICE_TYPE_DEFAULT && num == 0) ||
(device_type == CL_DEVICE_TYPE_ALL)) {
if ((devices != nullptr) && (num < num_entries)) {
devices[num] = dev;
}
num++;
if (device_type == CL_DEVICE_TYPE_DEFAULT) {
break;
}
}
}
if (num_devices != nullptr) {
*num_devices = num;
}
if (num == 0) {
return CL_DEVICE_NOT_FOUND;
}
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetDeviceInfo(cl_device_id dev,
cl_device_info param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("device", (uintptr_t)dev, "param_name", param_name);
LOG_API_CALL(
"device = %p, param_name = %d, size = %zu, value = %p, size_ret = %p",
dev, param_name, param_value_size, param_value, param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t size_ret = 0;
const void* copy_ptr = nullptr;
size_t val_sizet;
cl_uint val_uint;
api_query_string val_string;
cl_device_type val_devicetype;
cl_bool val_bool;
cl_device_fp_config val_fpconfig;
size_t work_item_sizes[3];
cl_device_mem_cache_type val_cache_type;
cl_ulong val_ulong;
cl_device_local_mem_type val_local_mem_type;
cl_device_partition_property val_partition_property;
cl_device_affinity_domain val_affinity_domain;
cl_device_exec_capabilities val_exec_capabilities;
cl_command_queue_properties val_queue_properties;
cl_platform_id val_platform;
cl_device_id val_deviceid;
cl_version val_version;
cl_device_svm_capabilities val_svmcaps;
cl_device_device_enqueue_capabilities val_dev_enqueue_caps;
cl_device_pci_bus_info_khr val_pci_bus_info;
cl_device_atomic_capabilities val_atomic_capabilities;
cl_device_integer_dot_product_capabilities_khr val_int_dot_product;
cl_device_integer_dot_product_acceleration_properties_khr
val_int_dot_product_props;
std::vector<size_t> val_subgroup_sizes;
auto device = icd_downcast(dev);
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
switch (param_name) {
case CL_DEVICE_PLATFORM:
val_platform = device->platform();
copy_ptr = &val_platform;
size_ret = sizeof(val_platform);
break;
case CL_DEVICE_TYPE:
val_devicetype = device->type();
copy_ptr = &val_devicetype;
size_ret = sizeof(val_devicetype);
break;
case CL_DEVICE_NAME:
val_string = device->name();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_VENDOR:
val_string = device->vendor();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_VENDOR_ID:
val_uint = device->vendor_id();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DRIVER_VERSION:
val_string = device->driver_version();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_VERSION:
val_string = device->version_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_OPENCL_C_VERSION:
val_string = device->c_version_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_PROFILE:
val_string = device->profile();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_BUILT_IN_KERNELS:
val_string = "";
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_EXTENSIONS:
val_string = device->extension_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_AVAILABLE:
val_bool = CL_TRUE;
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_COMPILER_AVAILABLE:
case CL_DEVICE_LINKER_AVAILABLE:
val_bool = device->compiler_available();
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_IMAGE_SUPPORT:
val_bool = device->supports_images();
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_IMAGE2D_MAX_WIDTH:
case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
val_sizet = device->vulkan_limits().maxImageDimension2D;
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_IMAGE3D_MAX_WIDTH:
case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
case CL_DEVICE_IMAGE3D_MAX_DEPTH:
val_sizet = device->vulkan_limits().maxImageDimension3D;
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_HALF_FP_CONFIG:
case CL_DEVICE_SINGLE_FP_CONFIG:
case CL_DEVICE_DOUBLE_FP_CONFIG:
val_fpconfig = device->fp_config(param_name);
copy_ptr = &val_fpconfig;
size_ret = sizeof(val_fpconfig);
break;
case CL_DEVICE_ADDRESS_BITS:
val_uint = device->address_bits();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
val_uint = device->mem_base_addr_align();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
val_uint = 128; // Alignment in bytes of long16
copy_ptr = &val_uint, size_ret = sizeof(val_uint);
break;
case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
val_cache_type = CL_NONE; // FIXME
copy_ptr = &val_cache_type;
size_ret = sizeof(val_cache_type);
break;
case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
val_ulong = device->global_mem_cache_size();
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
val_uint = device->vulkan_limits().nonCoherentAtomSize;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
val_bool = CL_FALSE; // FIXME
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_ENDIAN_LITTLE:
val_bool = CL_TRUE; // FIXME
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_HOST_UNIFIED_MEMORY:
val_bool = device->has_host_unified_memory();
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_MAX_WORK_GROUP_SIZE:
val_sizet = device->max_work_group_size();
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_MAX_COMPUTE_UNITS:
val_uint = device->num_compute_units();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
val_uint = device->max_work_item_dimensions();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MAX_WORK_ITEM_SIZES:
work_item_sizes[0] = device->vulkan_limits().maxComputeWorkGroupSize[0];
work_item_sizes[1] = device->vulkan_limits().maxComputeWorkGroupSize[1];
work_item_sizes[2] = device->vulkan_limits().maxComputeWorkGroupSize[2];
copy_ptr = work_item_sizes;
size_ret = sizeof(work_item_sizes);
break;
case CL_DEVICE_MAX_PARAMETER_SIZE:
val_sizet = 1024; // FIXME this is the minimum, revisit when looking
// into push constants
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_MAX_CONSTANT_ARGS:
val_uint = 8; // TODO be smarter
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: // TODO be smarter
val_ulong = 64 * 1024;
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
// FIXME can we do better for vector width queries?
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
val_uint = 1;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
val_uint = device->supports_fp16() ? 1 : 0;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
val_uint = device->supports_fp64() ? 1 : 0;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
val_sizet = 1;
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_GLOBAL_MEM_SIZE:
val_ulong = device->global_mem_size();
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
val_ulong = device->max_mem_alloc_size();
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
case CL_DEVICE_LOCAL_MEM_SIZE:
val_ulong = device->vulkan_limits().maxComputeSharedMemorySize;
copy_ptr = &val_ulong;
size_ret = sizeof(val_ulong);
break;
case CL_DEVICE_LOCAL_MEM_TYPE:
val_local_mem_type = CL_LOCAL; // FIXME try to be a bit smarter
copy_ptr = &val_local_mem_type;
size_ret = sizeof(val_local_mem_type);
break;
case CL_DEVICE_MAX_CLOCK_FREQUENCY:
val_uint = 0; // FIXME can we do better?
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PARTITION_MAX_SUB_DEVICES:
val_uint = 0; // TODO support
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PARTITION_PROPERTIES:
case CL_DEVICE_PARTITION_TYPE:
val_partition_property = 0; // TODO
copy_ptr = &val_partition_property;
size_ret = sizeof(val_partition_property);
break;
case CL_DEVICE_PARTITION_AFFINITY_DOMAIN:
val_affinity_domain = 0; // TODO
copy_ptr = &val_affinity_domain;
size_ret = sizeof(val_affinity_domain);
break;
case CL_DEVICE_EXECUTION_CAPABILITIES:
val_exec_capabilities = CL_EXEC_KERNEL;
copy_ptr = &val_exec_capabilities;
size_ret = sizeof(val_exec_capabilities);
break;
case CL_DEVICE_QUEUE_PROPERTIES:
val_queue_properties = CL_QUEUE_PROFILING_ENABLE;
copy_ptr = &val_queue_properties;
size_ret = sizeof(val_queue_properties);
break;
case CL_DEVICE_REFERENCE_COUNT:
val_uint = 1; // FIXME partitioning
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PARENT_DEVICE:
val_deviceid = nullptr; // TODO partitioning
copy_ptr = &val_deviceid;
size_ret = sizeof(val_deviceid);
break;
case CL_DEVICE_PREFERRED_INTEROP_USER_SYNC:
val_bool = CL_TRUE;
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_PRINTF_BUFFER_SIZE:
val_sizet = 1024 * 1024; // FIXME
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_MAX_SAMPLERS:
val_uint = device->max_samplers();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE:
val_sizet = device->image_max_buffer_size();
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE:
val_sizet = device->vulkan_limits().maxImageArrayLayers;
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_MAX_READ_IMAGE_ARGS:
val_uint = device->vulkan_limits().maxPerStageDescriptorSampledImages;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS:
if (!device->supports_read_write_images()) {
val_uint = 0;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
}
[[fallthrough]];
case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
val_uint = device->vulkan_limits().maxPerStageDescriptorStorageImages;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_IL_VERSION:
val_string = device->ils_string();
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_NUMERIC_VERSION:
val_version = device->version();
copy_ptr = &val_version;
size_ret = sizeof(val_version);
break;
case CL_DEVICE_OPENCL_C_NUMERIC_VERSION_KHR:
val_version = device->c_version();
copy_ptr = &val_version;
size_ret = sizeof(val_version);
break;
case CL_DEVICE_EXTENSIONS_WITH_VERSION:
copy_ptr = device->extensions().data();
size_ret = device->extensions().size() * sizeof(cl_name_version);
break;
case CL_DEVICE_ILS_WITH_VERSION:
copy_ptr = device->ils().data();
size_ret = device->ils().size() * sizeof(cl_name_version);
break;
case CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION:
copy_ptr = nullptr;
size_ret = 0;
break;
case CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT:
val_bool = device->supports_non_uniform_workgroup();
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE:
case CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE:
val_sizet = 0;
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_IMAGE_PITCH_ALIGNMENT:
case CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT:
val_uint = 0;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_SVM_CAPABILITIES:
val_svmcaps = 0;
copy_ptr = &val_svmcaps;
size_ret = sizeof(val_svmcaps);
break;
case CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS:
case CL_DEVICE_PIPE_SUPPORT:
case CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT:
val_bool = CL_FALSE;
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT:
// TODO(#216) re-enable when clspv ready
val_bool = CL_FALSE;
copy_ptr = &val_bool;
size_ret = sizeof(val_bool);
break;
case CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES:
val_dev_enqueue_caps = 0;
copy_ptr = &val_dev_enqueue_caps;
size_ret = sizeof(val_dev_enqueue_caps);
break;
case CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES:
val_queue_properties = 0;
copy_ptr = &val_queue_properties;
size_ret = sizeof(val_queue_properties);
break;
case CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE:
case CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE:
case CL_DEVICE_MAX_ON_DEVICE_QUEUES:
case CL_DEVICE_MAX_ON_DEVICE_EVENTS:
case CL_DEVICE_MAX_PIPE_ARGS:
case CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS:
case CL_DEVICE_PIPE_MAX_PACKET_SIZE:
val_uint = 0;
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_MAX_NUM_SUB_GROUPS:
val_uint = device->max_num_sub_groups();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_OPENCL_C_ALL_VERSIONS:
copy_ptr = device->opencl_c_versions().data();
size_ret = device->opencl_c_versions().size() * sizeof(cl_name_version);
break;
case CL_DEVICE_OPENCL_C_FEATURES:
copy_ptr = device->opencl_c_features().data();
size_ret = device->opencl_c_features().size() * sizeof(cl_name_version);
break;
case CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
val_sizet = device->preferred_work_group_size_multiple();
copy_ptr = &val_sizet;
size_ret = sizeof(val_sizet);
break;
case CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES:
val_atomic_capabilities =
CL_DEVICE_ATOMIC_ORDER_RELAXED | CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP;
if (device->supports_atomic_order_acq_rel()) {
val_atomic_capabilities |= CL_DEVICE_ATOMIC_ORDER_ACQ_REL;
}
if (device->supports_atomic_scope_device()) {
val_atomic_capabilities |= CL_DEVICE_ATOMIC_SCOPE_DEVICE;
}
copy_ptr = &val_atomic_capabilities;
size_ret = sizeof(val_atomic_capabilities);
break;
case CL_DEVICE_ATOMIC_FENCE_CAPABILITIES:
val_atomic_capabilities = CL_DEVICE_ATOMIC_ORDER_RELAXED |
CL_DEVICE_ATOMIC_SCOPE_WORK_ITEM |
CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP;
if (device->vulkan_memory_model_features().vulkanMemoryModel) {
val_atomic_capabilities |= CL_DEVICE_ATOMIC_ORDER_ACQ_REL;
}
if (device->vulkan_memory_model_features()
.vulkanMemoryModelDeviceScope) {
val_atomic_capabilities |= CL_DEVICE_ATOMIC_SCOPE_DEVICE;
}
copy_ptr = &val_atomic_capabilities;
size_ret = sizeof(val_atomic_capabilities);
break;
case CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED:
val_string = "v2023-12-12-00";
copy_ptr = val_string.c_str();
size_ret = val_string.size_with_null();
break;
case CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT:
case CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT:
case CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT:
val_uint = 0; // Natural size of the types
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_UUID_KHR:
copy_ptr = device->uuid();
size_ret = CL_UUID_SIZE_KHR;
break;
case CL_DRIVER_UUID_KHR:
copy_ptr = device->driver_uuid();
size_ret = CL_UUID_SIZE_KHR;
break;
case CL_DEVICE_LUID_VALID_KHR:
val_bool = device->luid_valid();
copy_ptr = &val_bool;
size_ret = sizeof(cl_bool);
break;
case CL_DEVICE_LUID_KHR:
copy_ptr = device->luid();
size_ret = CL_LUID_SIZE_KHR;
break;
case CL_DEVICE_NODE_MASK_KHR:
val_uint = device->node_mask();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_DEVICE_PCI_BUS_INFO_KHR:
val_pci_bus_info = device->pci_bus_info();
copy_ptr = &val_pci_bus_info;
size_ret = sizeof(val_pci_bus_info);
break;
case CL_DEVICE_INTEGER_DOT_PRODUCT_CAPABILITIES_KHR:
val_int_dot_product = device->dot_product_capabilities();
copy_ptr = &val_int_dot_product;
size_ret = sizeof(val_int_dot_product);
break;
case CL_DEVICE_INTEGER_DOT_PRODUCT_ACCELERATION_PROPERTIES_8BIT_KHR:
val_int_dot_product_props = device->dot_product_8bit_properties();
copy_ptr = &val_int_dot_product_props;
size_ret = sizeof(val_int_dot_product_props);
break;
case CL_DEVICE_INTEGER_DOT_PRODUCT_ACCELERATION_PROPERTIES_4x8BIT_PACKED_KHR:
val_int_dot_product_props =
device->dot_product_4x8bit_packed_properties();
copy_ptr = &val_int_dot_product_props;
size_ret = sizeof(val_int_dot_product_props);
break;
case CL_DEVICE_SUB_GROUP_SIZES_INTEL:
if (device->supports_subgroup_size_selection()) {
uint32_t size = device->min_sub_group_size();
while (size <= device->max_sub_group_size()) {
val_subgroup_sizes.push_back((size_t)size);
size *= 2;
}
copy_ptr = val_subgroup_sizes.data();
size_ret = sizeof(size_t) * val_subgroup_sizes.size();
} else {
ret = CL_INVALID_VALUE;
}
break;
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < size_ret) {
return CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, size_ret);
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = size_ret;
}
return ret;
}
cl_int CLVK_API_CALL clCreateSubDevices(
cl_device_id in_device, const cl_device_partition_property* properties,
cl_uint num_devices, cl_device_id* out_devices, cl_uint* num_devices_ret) {
TRACE_FUNCTION("in_device", (uintptr_t)in_device, "num_devices",
num_devices);
LOG_API_CALL("in_device = %p, properties = %p, num_devices = %u, "
"out_devices = %p, num_devices_ret = %p",
in_device, properties, num_devices, out_devices,
num_devices_ret);
// TODO CL_INVALID_DEVICE if in_device is not valid.
// TODO CL_INVALID_VALUE if values specified in properties are not valid or
// if values specified in properties are valid but not supported by the
// device.
// TODO CL_INVALID_VALUE if out_devices is not NULL and num_devices is less
// than the number of sub-devices created by the partition scheme.
// TODO CL_DEVICE_PARTITION_FAILED if the partition name is supported by the
// implementation but in_device could not be further partitioned.
// TODO CL_INVALID_DEVICE_PARTITION_COUNT if the partition name specified in
// properties is CL_DEVICE_PARTITION_BY_COUNTS and the number of sub-devices
// requested exceeds CL_DEVICE_PARTITION_MAX_SUB_DEVICES or the total number
// of compute units requested exceeds CL_DEVICE_PARTITION_MAX_COMPUTE_UNITS
// for in_device, or the number of compute units requested for one or more
// sub-devices is less than zero or the number of sub-devices requested
// exceeds CL_DEVICE_PARTITION_MAX_COMPUTE_UNITS for in_device.
// TODO CL_OUT_OF_RESOURCES if there is a failure to allocate resources
// required by the OpenCL implementation on the device.
// TODO CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources
// required by the OpenCL implementation on the host.
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clRetainDevice(cl_device_id device) {
TRACE_FUNCTION("device", (uintptr_t)device);
LOG_API_CALL("device = %p", device);
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseDevice(cl_device_id device) {
TRACE_FUNCTION("device", (uintptr_t)device);
LOG_API_CALL("device = %p", device);
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
return CL_SUCCESS;
}
cvk_context* cvk_create_context(
const cl_context_properties* properties, cl_uint num_devices,
const cl_device_id* devices,
void(CL_CALLBACK* pfn_notify)(const char*, const void*, size_t, void*),
void* user_data, cl_int* errcode_ret) {
if ((devices == nullptr) || (num_devices == 0) ||
((pfn_notify == nullptr) && (user_data != nullptr))) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
if (num_devices > 1) {
cvk_error("Only one device per context is supported.");
return nullptr;
} else if (!is_valid_device(devices[0])) {
*errcode_ret = CL_INVALID_DEVICE;
return nullptr;
}
auto context =
new cvk_context(icd_downcast(devices[0]), properties, user_data);
*errcode_ret = context->init();
if (*errcode_ret != CL_SUCCESS) {
return nullptr;
}
return context;
}
// Context APIs
cl_context CLVK_API_CALL clCreateContext(
const cl_context_properties* properties, cl_uint num_devices,
const cl_device_id* devices,
void(CL_CALLBACK* pfn_notify)(const char*, const void*, size_t, void*),
void* user_data, cl_int* errcode_ret) {
TRACE_FUNCTION("num_devices", num_devices);
LOG_API_CALL("properties = %p, num_devices = %u, devices = %p, pfn_notify "
"= %p, user_data = %p, errcode_ret = %p",
properties, num_devices, devices, pfn_notify, user_data,
errcode_ret);
cl_int err;
auto context = cvk_create_context(properties, num_devices, devices,
pfn_notify, user_data, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return context;
}
cl_context CLVK_API_CALL clCreateContextFromType(
const cl_context_properties* properties, cl_device_type device_type,
void(CL_CALLBACK* pfn_notify)(const char*, const void*, size_t, void*),
void* user_data, cl_int* errcode_ret) {
TRACE_FUNCTION("device_type",
TRACE_STRING(cl_device_type_to_string(device_type)));
LOG_API_CALL("properties = %p, device_type = %lu (%s), pfn_notify = %p, "
"user_data = %p, errcode_ret = %p",
properties, device_type, cl_device_type_to_string(device_type),
pfn_notify, user_data, errcode_ret);
cl_device_id device;
// TODO introduce cvk_ functions to get correct logging
cl_int err = clGetDeviceIDs(nullptr, device_type, 1, &device, nullptr);
cvk_context* context = nullptr;
if (err == CL_SUCCESS) {
context = cvk_create_context(properties, 1, &device, pfn_notify,
user_data, &err);
}
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return context;
}
cl_int CLVK_API_CALL clRetainContext(cl_context context) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p", context);
if (!is_valid_context(context)) {
return CL_INVALID_CONTEXT;
}
icd_downcast(context)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseContext(cl_context context) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p", context);
if (!is_valid_context(context)) {
return CL_INVALID_CONTEXT;
}
icd_downcast(context)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clSetContextDestructorCallback(
cl_context context,
void(CL_CALLBACK* pfn_notify)(cl_context context, void* user_data),
void* user_data) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, pfn_notify = %p, user_data = %p", context,
pfn_notify, user_data);
if (!is_valid_context(context)) {
return CL_INVALID_CONTEXT;
}
if (pfn_notify == nullptr) {
return CL_INVALID_VALUE;
}
icd_downcast(context)->add_destructor_callback(pfn_notify, user_data);
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetContextInfo(cl_context ctx,
cl_context_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("context", (uintptr_t)ctx, "param_name", param_name);
LOG_API_CALL(
"context = %p, param_name = %u, size = %zu, value = %p, size_ret = %p",
ctx, param_name, param_value_size, param_value, param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t size_ret = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_device_id val_device;
if (!is_valid_context(ctx)) {
return CL_INVALID_CONTEXT;
}
auto context = icd_downcast(ctx);
switch (param_name) {
case CL_CONTEXT_REFERENCE_COUNT:
val_uint = context->refcount();
copy_ptr = &val_uint;
size_ret = sizeof(val_uint);
break;
case CL_CONTEXT_DEVICES:
val_device = context->device();
copy_ptr = &val_device;
size_ret = sizeof(val_device);
break;
case CL_CONTEXT_NUM_DEVICES:
val_uint = context->num_devices();
copy_ptr = &val_uint;
size_ret = sizeof(cl_uint);
break;
case CL_CONTEXT_PROPERTIES:
if (context->properties().size() == 0) {
size_ret = 0;
copy_ptr = nullptr;
} else {
copy_ptr = context->properties().data();
size_ret =
context->properties().size() * sizeof(cl_context_properties);
}
break;
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < size_ret) {
return CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, size_ret);
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = size_ret;
}
return ret;
}
// Event APIs
cl_int CLVK_API_CALL clWaitForEvents(cl_uint num_events,
const cl_event* event_list) {
TRACE_FUNCTION("num_events", num_events);
LOG_API_CALL("num_events = %u, event_list = %p", num_events, event_list);
if ((num_events == 0) || (event_list == nullptr)) {
return CL_INVALID_VALUE;
}
// TODO validate that all events belong to the same context
for (cl_uint i = 0; i < num_events; i++) {
if (!is_valid_event(event_list[i])) {
return CL_INVALID_EVENT;
}
}
return cvk_command_queue::wait_for_events(num_events, event_list);
}
cl_int CLVK_API_CALL clEnqueueWaitForEvents(cl_command_queue command_queue,
cl_uint num_events,
const cl_event* event_list) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue, "num_events",
num_events);
LOG_API_CALL("command_queue = %p, num_events = %u, event_list = %p",
command_queue, num_events, event_list);
return CL_INVALID_OPERATION; // TODO implement
}
cl_int CLVK_API_CALL clReleaseEvent(cl_event event) {
TRACE_FUNCTION("event", (uintptr_t)event);
LOG_API_CALL("event = %p", event);
if (!is_valid_event(event)) {
return CL_INVALID_EVENT;
}
icd_downcast(event)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clRetainEvent(cl_event event) {
TRACE_FUNCTION("event", (uintptr_t)event);
LOG_API_CALL("event = %p", event);
if (!is_valid_event(event)) {
return CL_INVALID_EVENT;
}
icd_downcast(event)->retain();
return CL_SUCCESS;
}
cl_event CLVK_API_CALL clCreateUserEvent(cl_context context,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, errcode_ret = %p", context, errcode_ret);
if (!is_valid_context(context)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
}
auto event = new cvk_event_command(icd_downcast(context), nullptr, nullptr);
if (errcode_ret != nullptr) {
*errcode_ret = CL_SUCCESS;
}
return event;
}
cl_int CLVK_API_CALL clSetUserEventStatus(cl_event event,
cl_int execution_status) {
TRACE_FUNCTION(
"event", (uintptr_t)event, "execution_status",
TRACE_STRING(cl_command_execution_status_to_string(execution_status)));
LOG_API_CALL("event = %p, execution_status = %d (%s)", event,
execution_status,
cl_command_execution_status_to_string(execution_status));
if (!is_valid_event(event) || !icd_downcast(event)->is_user_event()) {
return CL_INVALID_EVENT;
}
if (execution_status > CL_COMPLETE) {
return CL_INVALID_VALUE;
}
icd_downcast(event)->set_status(execution_status);
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clSetEventCallback(
cl_event event, cl_int command_exec_callback_type,
void(CL_CALLBACK* pfn_event_notify)(cl_event event,
cl_int event_command_exec_status,
void* user_data),
void* user_data) {
TRACE_FUNCTION("event", (uintptr_t)event, "execution_status",
TRACE_STRING(cl_command_execution_status_to_string(
command_exec_callback_type)));
LOG_API_CALL(
"event = %p, callback_type = %d (%s), pfn_event_notify = %p, user_data "
"= %p",
event, command_exec_callback_type,
cl_command_execution_status_to_string(command_exec_callback_type),
pfn_event_notify, user_data);
if (!is_valid_event(event)) {
return CL_INVALID_EVENT;
}
if (pfn_event_notify == nullptr) {
return CL_INVALID_VALUE;
}
icd_downcast(event)->register_callback(command_exec_callback_type,
pfn_event_notify, user_data);
return CL_SUCCESS;
}
cl_int cvk_enqueue_marker_with_wait_list(cvk_command_queue* command_queue,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
auto cmd = new cvk_command_dep(command_queue, CL_COMMAND_MARKER);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueMarkerWithWaitList(
cl_command_queue command_queue, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
command_queue, num_events_in_wait_list, event_wait_list,
event);
return cvk_enqueue_marker_with_wait_list(icd_downcast(command_queue),
num_events_in_wait_list,
event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueMarker(cl_command_queue command_queue,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p, event = %p", command_queue, event);
return cvk_enqueue_marker_with_wait_list(icd_downcast(command_queue), 0,
nullptr, event);
}
cl_int cvk_enqueue_barrier_with_wait_list(cvk_command_queue* command_queue,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_dep(command_queue, CL_COMMAND_BARRIER);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueBarrierWithWaitList(
cl_command_queue command_queue, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
command_queue, num_events_in_wait_list, event_wait_list,
event);
return cvk_enqueue_barrier_with_wait_list(icd_downcast(command_queue),
num_events_in_wait_list,
event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueBarrier(cl_command_queue command_queue) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
return cvk_enqueue_barrier_with_wait_list(icd_downcast(command_queue), 0,
nullptr, nullptr);
}
cl_int CLVK_API_CALL clGetEventInfo(cl_event evt, cl_event_info param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("event", (uintptr_t)evt, "param_name", param_name);
LOG_API_CALL("event = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
evt, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_int val_int;
cl_context val_context;
cl_command_type val_command_type;
cl_command_queue val_command_queue;
auto event = icd_downcast(evt);
if (!is_valid_event(event)) {
return CL_INVALID_EVENT;
}
switch (param_name) {
case CL_EVENT_REFERENCE_COUNT:
val_uint = event->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_EVENT_COMMAND_EXECUTION_STATUS:
val_int = event->get_status();
copy_ptr = &val_int;
ret_size = sizeof(val_int);
break;
case CL_EVENT_COMMAND_QUEUE:
if (event->is_user_event()) {
val_command_queue = nullptr;
} else {
val_command_queue = event->queue();
}
copy_ptr = &val_command_queue;
ret_size = sizeof(val_command_queue);
break;
case CL_EVENT_COMMAND_TYPE:
if (event->is_user_event()) {
val_command_type = CL_COMMAND_USER;
} else {
val_command_type = event->command_type();
}
copy_ptr = &val_command_type;
ret_size = sizeof(val_command_type);
break;
case CL_EVENT_CONTEXT:
val_context = event->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
// Command Queue APIs
cl_command_queue
cvk_create_command_queue(cl_context context, cl_device_id device,
cl_command_queue_properties properties,
std::vector<cl_queue_properties>&& properties_array,
cl_int* errcode_ret) {
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
if (!is_valid_device(device) ||
icd_downcast(context)->device() != icd_downcast(device)) {
*errcode_ret = CL_INVALID_DEVICE;
return nullptr;
}
const cl_command_queue_properties valid_properties =
CL_QUEUE_PROFILING_ENABLE | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE |
CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT;
if ((properties & ~valid_properties) > 0) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
cl_command_queue_properties supported_properties =
CL_QUEUE_PROFILING_ENABLE;
// We do not support out of order command queues, but allow the property to
// support application using it.
if (config.ignore_out_of_order_execution()) {
supported_properties |= CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
}
if ((properties & ~supported_properties) > 0) {
*errcode_ret = CL_INVALID_QUEUE_PROPERTIES;
return nullptr;
}
auto queue = std::make_unique<cvk_command_queue>(
icd_downcast(context), icd_downcast(device), properties,
std::move(properties_array));
cl_int err = queue->init();
*errcode_ret = err;
if (err != CL_SUCCESS) {
return nullptr;
} else {
return queue.release();
}
}
cl_command_queue CLVK_API_CALL clCreateCommandQueue(
cl_context context, cl_device_id device,
cl_command_queue_properties properties, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "device", (uintptr_t)device,
"properties", properties);
LOG_API_CALL(
"context = %p, device = %p, properties = %lu, errcode_ret = %p",
context, device, properties, errcode_ret);
cl_int err;
std::vector<cl_queue_properties> properties_array;
auto ret = cvk_create_command_queue(context, device, properties,
std::move(properties_array), &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return ret;
}
cl_command_queue cvk_create_command_queue_with_properties(
cl_context context, cl_device_id device,
const cl_queue_properties* properties, cl_int* errcode_ret) {
cl_command_queue_properties props = 0;
std::vector<cl_queue_properties> properties_array;
if (properties) {
while (*properties) {
auto key = *properties;
auto value = *(properties + 1);
properties_array.push_back(key);
properties_array.push_back(value);
if (key == CL_QUEUE_PROPERTIES) {
props = value;
} else {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
properties += 2;
}
properties_array.push_back(0);
}
cl_int err;
auto ret = cvk_create_command_queue(context, device, props,
std::move(properties_array), &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return ret;
}
cl_command_queue CLVK_API_CALL clCreateCommandQueueWithProperties(
cl_context context, cl_device_id device,
const cl_queue_properties* properties, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "device", (uintptr_t)device);
LOG_API_CALL("context = %p, device = %p, properties = %p, errcode_ret = %p",
context, device, properties, errcode_ret);
return cvk_create_command_queue_with_properties(context, device, properties,
errcode_ret);
}
cl_command_queue CLVK_API_CALL clCreateCommandQueueWithPropertiesKHR(
cl_context context, cl_device_id device,
const cl_queue_properties* properties, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "device", (uintptr_t)device);
LOG_API_CALL("context = %p, device = %p, properties = %p, errcode_ret = %p",
context, device, properties, errcode_ret);
return cvk_create_command_queue_with_properties(context, device, properties,
errcode_ret);
}
cl_int CLVK_API_CALL clReleaseCommandQueue(cl_command_queue command_queue) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
cl_int err = icd_downcast(command_queue)->flush();
icd_downcast(command_queue)->release();
return err;
}
cl_int CLVK_API_CALL clRetainCommandQueue(cl_command_queue command_queue) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
icd_downcast(command_queue)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetCommandQueueInfo(cl_command_queue cq,
cl_command_queue_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "param_name", param_name);
LOG_API_CALL("command_queue = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
cq, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_context val_context;
cl_device_id val_device;
cl_command_queue_properties val_properties;
cl_command_queue val_queue;
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
switch (param_name) {
case CL_QUEUE_REFERENCE_COUNT:
val_uint = command_queue->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_QUEUE_CONTEXT:
val_context = command_queue->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_QUEUE_DEVICE:
val_device = command_queue->device();
copy_ptr = &val_device;
ret_size = sizeof(val_device);
break;
case CL_QUEUE_PROPERTIES:
val_properties = command_queue->properties();
copy_ptr = &val_properties;
ret_size = sizeof(val_properties);
break;
case CL_QUEUE_SIZE:
val_uint = 0;
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
ret = CL_INVALID_COMMAND_QUEUE;
break;
case CL_QUEUE_DEVICE_DEFAULT:
val_queue = nullptr;
copy_ptr = &val_queue;
ret_size = sizeof(val_queue);
break;
case CL_QUEUE_PROPERTIES_ARRAY:
copy_ptr = command_queue->properties_array().data();
ret_size = command_queue->properties_array().size() *
sizeof(cl_queue_properties);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clSetDefaultDeviceCommandQueue(
cl_context context, cl_device_id device, cl_command_queue command_queue) {
TRACE_FUNCTION("context", (uintptr_t)context, "device", (uintptr_t)device,
"command_queue", (uintptr_t)command_queue);
LOG_API_CALL("context = %p, device = %p, command_queue = %p", context,
device, command_queue);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clSetCommandQueueProperty(
cl_command_queue command_queue, cl_command_queue_properties properties,
cl_bool enable, cl_command_queue_properties* old_properties) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue, "properties",
properties);
LOG_API_CALL("command_queue = %p, properties = %lx, enable = %d, "
"old_properties = %p",
command_queue, properties, enable, old_properties);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
// TODO validate properties
// TODO support
return CL_INVALID_OPERATION;
}
// Memory Object APIs
static cl_mem CLVK_API_CALL cvk_create_buffer_with_properties(
cl_context context, const cl_mem_properties* properties, cl_mem_flags flags,
size_t size, void* host_ptr, cl_int* errcode_ret) {
CVK_ASSERT(errcode_ret != nullptr);
// Validate context
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
// Validate properties
std::vector<cl_mem_properties> props;
if (properties != nullptr) {
while (*properties) {
// We dont't currently support any properties so return an error
*errcode_ret = CL_INVALID_PROPERTY;
return nullptr;
props.push_back(*properties);
properties++;
}
props.push_back(0);
}
// Validate flags
if ((flags & CL_MEM_READ_WRITE) && (flags & CL_MEM_WRITE_ONLY)) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
if ((flags & CL_MEM_READ_ONLY) &&
(flags & (CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE))) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
if ((flags & (CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR)) &&
(flags & CL_MEM_USE_HOST_PTR)) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
if ((flags & CL_MEM_HOST_READ_ONLY) && (flags & CL_MEM_HOST_WRITE_ONLY)) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
if ((flags & CL_MEM_HOST_NO_ACCESS) &&
(flags & (CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_WRITE_ONLY))) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
// Validate size
// TODO CL_INVALID_BUFFER_SIZE if CL_MEM_USE_HOST_PTR is set in flags and
// host_ptr is a pointer returned by clSVMAlloc and size is greater than the
// size passed to clSVMAlloc.
if ((size == 0) ||
(!icd_downcast(context)->is_mem_alloc_size_valid(size))) {
*errcode_ret = CL_INVALID_BUFFER_SIZE;
return nullptr;
}
// Validate host_ptr
if ((host_ptr == nullptr) &&
(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) {
*errcode_ret = CL_INVALID_HOST_PTR;
return nullptr;
}
if ((host_ptr != nullptr) &&
!(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) {
*errcode_ret = CL_INVALID_HOST_PTR;
return nullptr;
}
auto buffer = cvk_buffer::create(icd_downcast(context), flags, size,
host_ptr, std::move(props), errcode_ret);
if (*errcode_ret != CL_SUCCESS) {
return nullptr;
} else {
return buffer.release();
}
}
cl_mem CLVK_API_CALL clCreateBuffer(cl_context context, cl_mem_flags flags,
size_t size, void* host_ptr,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL("context = %p, flags = %lu, size = %zu, host_ptr = %p, "
"errcode_ret = %p",
context, flags, size, host_ptr, errcode_ret);
cl_int err;
auto buffer = cvk_create_buffer_with_properties(context, nullptr, flags,
size, host_ptr, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return buffer;
}
cl_mem CLVK_API_CALL clCreateBufferWithProperties(
cl_context context, const cl_mem_properties* properties, cl_mem_flags flags,
size_t size, void* host_ptr, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL("context = %p, properties = %p, flags = %lx, size = %zu, "
"host_ptr = %p, errcode_ret = %p",
context, properties, flags, size, host_ptr, errcode_ret);
cl_int err;
auto buffer = cvk_create_buffer_with_properties(context, properties, flags,
size, host_ptr, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return buffer;
}
cl_mem CLVK_API_CALL clCreateSubBuffer(cl_mem buf, cl_mem_flags flags,
cl_buffer_create_type buffer_create_type,
const void* buffer_create_info,
cl_int* errcode_ret) {
TRACE_FUNCTION("buffer", (uintptr_t)buf, "flags", flags,
"buffer_create_type", buffer_create_type);
LOG_API_CALL("buffer = %p, flags = %lu, buffer_create_type = %u, "
"buffer_create_info = %p, errcode_ret = %p",
buf, flags, buffer_create_type, buffer_create_info,
errcode_ret);
auto buffer = static_cast<cvk_buffer*>(buf);
if (!is_valid_buffer(buffer) || buffer->is_sub_buffer()) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_MEM_OBJECT;
}
return nullptr;
}
// TODO CL_INVALID_VALUE if buffer was created with CL_MEM_WRITE_ONLY and
// flags specifies CL_MEM_READ_WRITE or CL_MEM_READ_ONLY, or if buffer was
// created with CL_MEM_READ_ONLY and flags specifies CL_MEM_READ_WRITE or
// CL_MEM_WRITE_ONLY, or if flags specifies CL_MEM_USE_HOST_PTR or
// CL_MEM_ALLOC_HOST_PTR or CL_MEM_COPY_HOST_PTR.
// TODO CL_INVALID_VALUE if buffer was created with CL_MEM_HOST_WRITE_ONLY
// and flags specifies CL_MEM_HOST_READ_ONLY or if buffer was created with
// CL_MEM_HOST_READ_ONLY and flags specifies CL_MEM_HOST_WRITE_ONLY, or if
// buffer was created with CL_MEM_HOST_NO_ACCESS and flags specifies
// CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_WRITE_ONLY.
// TODO CL_INVALID_VALUE if value specified in buffer_create_type is not
// valid.
// TODO CL_INVALID_VALUE if value(s) specified in buffer_create_info (for a
// given buffer_create_type) is not valid or if buffer_create_info is NULL.
// TODO CL_INVALID_BUFFER_SIZE if size is 0.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for sub-buffer object.
// TODO CL_OUT_OF_RESOURCES if there is a failure to allocate resources
// required by the OpenCL implementation on the device.
// TODO CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources
// required by the OpenCL implementation on the host.
if (buffer_create_type != CL_BUFFER_CREATE_TYPE_REGION) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
auto region = static_cast<const cl_buffer_region*>(buffer_create_info);
LOG_API_CALL("CL_BUFFER_CREATE_TYPE_REGION, origin = %zu, size = %zu",
region->origin, region->size);
cl_int err = CL_SUCCESS;
auto sub = buffer->create_subbuffer(flags, region->origin, region->size);
if (sub == nullptr) {
err = CL_OUT_OF_RESOURCES;
}
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return sub;
}
cl_int CLVK_API_CALL clRetainMemObject(cl_mem memobj) {
TRACE_FUNCTION("memobj", (uintptr_t)memobj);
LOG_API_CALL("memobj = %p", memobj);
if (!is_valid_mem_object(memobj)) {
return CL_INVALID_MEM_OBJECT;
}
icd_downcast(memobj)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseMemObject(cl_mem memobj) {
TRACE_FUNCTION("memobj", (uintptr_t)memobj);
LOG_API_CALL("memobj = %p", memobj);
if (!is_valid_mem_object(memobj)) {
return CL_INVALID_MEM_OBJECT;
}
icd_downcast(memobj)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clSetMemObjectDestructorCallback(
cl_mem memobj,
void(CL_CALLBACK* pfn_notify)(cl_mem memobj, void* user_data),
void* user_data) {
TRACE_FUNCTION("memobj", (uintptr_t)memobj);
LOG_API_CALL("memobj = %p, pfn_notify = %p, user_data = %p", memobj,
pfn_notify, user_data);
if (!is_valid_mem_object(memobj)) {
return CL_INVALID_MEM_OBJECT;
}
if (pfn_notify == nullptr) {
return CL_INVALID_VALUE;
}
icd_downcast(memobj)->add_destructor_callback(pfn_notify, user_data);
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clEnqueueMigrateMemObjects(
cl_command_queue cq, cl_uint num_mem_objects, const cl_mem* mem_objects,
cl_mem_migration_flags flags, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "num_mem_objects",
num_mem_objects, "flags", flags, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, num_mem_objects = %u, mem_objects = %p, "
"flags = %lx, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
cq, num_mem_objects, mem_objects, flags,
num_events_in_wait_list, event_wait_list, event);
if ((num_mem_objects == 0) || (mem_objects == nullptr)) {
return CL_INVALID_VALUE;
}
for (cl_uint i = 0; i < num_mem_objects; i++) {
if (!is_valid_mem_object(mem_objects[i])) {
return CL_INVALID_MEM_OBJECT;
}
}
// TODO CL_INVALID_VALUE if flags is not 0 or is not any of the values
// described in the table above.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for the specified set of memory objects in mem_objects.
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
for (cl_uint i = 0; i < num_mem_objects; i++) {
if (!is_same_context(command_queue, mem_objects[i])) {
return CL_INVALID_CONTEXT;
}
}
auto cmd =
new cvk_command_dep(command_queue, CL_COMMAND_MIGRATE_MEM_OBJECTS);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clGetMemObjectInfo(cl_mem mem, cl_mem_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("memobj", (uintptr_t)mem, "param_name", param_name);
LOG_API_CALL("memobj = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
mem, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_context val_context;
cl_mem_object_type val_object_type;
cl_mem_flags val_flags;
size_t val_sizet;
cl_mem val_memobj;
void* val_ptr;
cl_bool val_bool;
auto memobj = icd_downcast(mem);
if (!is_valid_mem_object(memobj)) {
return CL_INVALID_MEM_OBJECT;
}
switch (param_name) {
case CL_MEM_REFERENCE_COUNT:
val_uint = memobj->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_MEM_CONTEXT:
val_context = memobj->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_MEM_TYPE:
val_object_type = memobj->type();
copy_ptr = &val_object_type;
ret_size = sizeof(val_object_type);
break;
case CL_MEM_FLAGS:
val_flags = memobj->flags();
copy_ptr = &val_flags;
ret_size = sizeof(val_flags);
break;
case CL_MEM_SIZE:
val_sizet = memobj->size();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_MEM_MAP_COUNT:
val_uint = memobj->map_count();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_MEM_ASSOCIATED_MEMOBJECT:
if (memobj->is_image_type()) {
auto img = static_cast<cvk_image*>(memobj);
val_memobj = img->buffer();
} else {
val_memobj = memobj->parent();
}
copy_ptr = &val_memobj;
ret_size = sizeof(val_memobj);
break;
case CL_MEM_OFFSET:
val_sizet = memobj->parent_offset();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_MEM_HOST_PTR:
val_ptr = nullptr;
if (memobj->has_any_flag(CL_MEM_USE_HOST_PTR)) {
val_ptr = memobj->host_ptr();
}
copy_ptr = &val_ptr;
ret_size = sizeof(val_ptr);
break;
case CL_MEM_USES_SVM_POINTER:
val_bool = CL_FALSE;
copy_ptr = &val_bool;
ret_size = sizeof(val_bool);
break;
case CL_MEM_PROPERTIES:
copy_ptr = memobj->properties().data();
ret_size = memobj->properties().size() * sizeof(cl_mem_properties);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
// Program Object APIs
cl_program CLVK_API_CALL clCreateProgramWithSource(cl_context context,
cl_uint count,
const char** strings,
const size_t* lengths,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "count", count);
LOG_API_CALL("context = %p, count = %u, lengths = %p", context, count,
lengths);
if (!is_valid_context(context)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
return nullptr;
}
if (count == 0 || strings == nullptr) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
cvk_program* prog = new cvk_program(icd_downcast(context));
for (cl_uint i = 0; i < count; i++) {
if (strings[i] == nullptr) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
size_t len = (lengths != nullptr) ? lengths[i] : 0;
prog->append_source(strings[i], len);
}
if (errcode_ret != nullptr) {
*errcode_ret = CL_SUCCESS;
}
cvk_debug_fn("\n%s", prog->source().c_str());
return prog;
}
cl_program CLVK_API_CALL clCreateProgramWithBinary(
cl_context ctx, cl_uint num_devices, const cl_device_id* device_list,
const size_t* lengths, const unsigned char** binaries,
cl_int* binary_status, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)ctx, "num_devices", num_devices);
LOG_API_CALL("context = %p, num_devices = %u, device_list = %p, lengths = "
"%p, binaries = %p, binary_status = %p, errcode_ret = %p",
ctx, num_devices, device_list, lengths, binaries,
binary_status, errcode_ret);
auto context = icd_downcast(ctx);
if (!is_valid_context(context)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
return nullptr;
}
if ((num_devices != 1) || (device_list == nullptr)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
if (icd_downcast(device_list[0]) != context->device()) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_DEVICE;
}
return nullptr;
}
if ((lengths == nullptr) || (binaries == nullptr)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
for (cl_uint i = 0; i < num_devices; i++) {
if ((lengths[i] == 0) || (binaries[i] == nullptr)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
}
cvk_program* prog = new cvk_program(icd_downcast(context));
cl_int load_status = CL_SUCCESS;
if (!prog->read(binaries[0], lengths[0])) {
load_status = CL_INVALID_BINARY;
}
if (binary_status != nullptr) {
binary_status[0] = load_status;
}
if (errcode_ret != nullptr) {
*errcode_ret = load_status;
}
return prog;
}
cl_program CLVK_API_CALL clCreateProgramWithBuiltInKernels(
cl_context context, cl_uint num_devices, const cl_device_id* device_list,
const char* kernel_names, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "num_devices", num_devices);
LOG_API_CALL("context = %p, num_devices = %u, device_list = %p, "
"kernel_names = \"%s\", errcode_ret = %p",
context, num_devices, device_list, kernel_names, errcode_ret);
if (!is_valid_context(context)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
return nullptr;
}
if ((device_list == nullptr) || (num_devices == 0)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
if (kernel_names == nullptr) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
// TODO CL_INVALID_DEVICE if any device in device_list is not in the list of
// devices associated with context.
if (errcode_ret != nullptr) {
*errcode_ret =
CL_INVALID_VALUE; // Since no built-in kernels are supported
}
return nullptr;
}
cl_int CLVK_API_CALL
clBuildProgram(cl_program prog, cl_uint num_devices,
const cl_device_id* device_list, const char* options,
void(CL_CALLBACK* pfn_notify)(cl_program /* program */,
void* /* user_data */),
void* user_data) {
TRACE_FUNCTION("program", (uintptr_t)prog, "num_devices", num_devices);
LOG_API_CALL("program = %p, num_device = %d, device_list = %p, options = "
"%s, pfn_notify = %p, user_data = %p",
prog, num_devices, device_list, options, pfn_notify,
user_data);
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
if (((num_devices > 0) && (device_list == nullptr)) ||
((num_devices == 0) && (device_list != nullptr))) {
return CL_INVALID_VALUE;
}
if ((pfn_notify == nullptr) && (user_data != nullptr)) {
return CL_INVALID_VALUE;
}
// TODO CL_INVALID_DEVICE if OpenCL devices listed in device_list are not in
// the list of devices associated with program.
// TODO CL_INVALID_BINARY if program is created with
// clCreateProgramWithBinary and devices listed in device_list do not have a
// valid program binary loaded.
// TODO CL_INVALID_BUILD_OPTIONS if the build options specified by options
// are invalid.
build_operation build_op = build_operation::build;
if (program->loaded_from_binary()) {
build_op = build_operation::build_binary;
} else if (!is_compiler_available(num_devices, device_list)) {
return CL_COMPILER_NOT_AVAILABLE;
}
// TODO CL_BUILD_PROGRAM_FAILURE if there is a failure to build the program
// executable. This error will be returned if clBuildProgram does not return
// until the build has completed.
// TODO CL_INVALID_OPERATION if there are kernel objects attached to
// program.
// TODO CL_INVALID_OPERATION if program was not created with
// clCreateProgramWithSource or clCreateProgramWithBinary or
// clCreateProgramWithILKHR.
return program->build(build_op, num_devices, device_list, options, 0,
nullptr, nullptr, pfn_notify, user_data);
}
cl_int CLVK_API_CALL clCompileProgram(
cl_program prog, cl_uint num_devices, const cl_device_id* device_list,
const char* options, cl_uint num_input_headers,
const cl_program* input_headers, const char** header_include_names,
void(CL_CALLBACK* pfn_notify)(cl_program program, void* user_data),
void* user_data) {
TRACE_FUNCTION("program", (uintptr_t)prog, "num_devices", num_devices,
"num_input_headers", num_input_headers);
LOG_API_CALL("program = %p, num_devices = %u, device_list = %p, options = "
"%p, num_input_headers = %u, input_headers = %p, "
"header_include_names = %p, pfn_notify = %p, user_data = %p",
prog, num_devices, device_list, options, num_input_headers,
input_headers, header_include_names, pfn_notify, user_data);
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
if (((device_list == nullptr) && (num_devices > 0)) ||
((device_list != nullptr) && (num_devices == 0))) {
return CL_INVALID_VALUE;
}
if (((num_input_headers == 0) &&
((header_include_names != nullptr) || (input_headers != nullptr))) ||
((num_input_headers != 0) &&
((header_include_names == nullptr) || (input_headers == nullptr)))) {
return CL_INVALID_VALUE;
}
if ((pfn_notify == nullptr) && (user_data != nullptr)) {
return CL_INVALID_VALUE;
}
// TODO CL_INVALID_DEVICE if OpenCL devices listed in device_list are not in
// the list of devices associated with program.
// TODO CL_INVALID_COMPILER_OPTIONS if the compiler options specified by
// options are invalid.
if (!is_compiler_available(num_devices, device_list)) {
return CL_COMPILER_NOT_AVAILABLE;
}
// TODO CL_COMPILE_PROGRAM_FAILURE if there is a failure to compile the
// program source. This error will be returned if clCompileProgram does not
// return until the compile has completed.
// TODO CL_INVALID_OPERATION if there are kernel objects attached to
// program.
if (program->loaded_from_binary()) {
return CL_INVALID_OPERATION;
}
// TODO Validate program
return program->build(build_operation::compile, num_devices, device_list,
options, num_input_headers, input_headers,
header_include_names, pfn_notify, user_data);
}
cl_program CLVK_API_CALL clLinkProgram(
cl_context context, cl_uint num_devices, const cl_device_id* device_list,
const char* options, cl_uint num_input_programs,
const cl_program* input_programs,
void(CL_CALLBACK* pfn_notify)(cl_program program, void* user_data),
void* user_data, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "num_devices", num_devices,
"num_input_programs", num_input_programs);
LOG_API_CALL("context = %p, num_devices = %d, device_list = %p, options = "
"%p, num_input_programs = %d, input_programs = %p, pfn_notify "
"= %p, user_data = %p, errcode_ret = %p",
context, num_devices, device_list, options, num_input_programs,
input_programs, pfn_notify, user_data, errcode_ret);
if (!is_valid_context(context)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
return nullptr;
}
if (((device_list == nullptr) && (num_devices > 0)) ||
((device_list != nullptr) && (num_devices == 0))) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
if (((input_programs == nullptr) && (num_input_programs == 0)) ||
((num_input_programs == 0) && (input_programs != nullptr)) ||
((num_input_programs != 0) && (input_programs == nullptr))) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
for (cl_uint i = 0; i < num_input_programs; i++) {
if (!is_valid_program(input_programs[i])) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_PROGRAM;
}
return nullptr;
}
}
if ((pfn_notify == nullptr) && (user_data != nullptr)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
// TODO CL_INVALID_DEVICE if OpenCL devices listed in device_list are not in
// the list of devices associated with context.
if (!is_compiler_available(num_devices, device_list)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_LINKER_NOT_AVAILABLE;
}
return nullptr;
}
// TODO CL_INVALID_LINKER_OPTIONS if the linker options specified by options
// are invalid
// TODO CL_INVALID_OPERATION if the rules for devices containing compiled
// binaries or libraries as described in input_programs argument above are
// not followed.
for (cl_uint i = 0; i < num_input_programs; i++) {
if (!icd_downcast(input_programs[i])->can_be_linked()) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_OPERATION;
}
return nullptr;
}
}
cvk_program* prog_ret = new cvk_program(icd_downcast(context));
cl_int ret = prog_ret->build(
build_operation::link, num_devices, device_list, options,
num_input_programs, input_programs, nullptr, pfn_notify, user_data);
if (errcode_ret != nullptr) {
*errcode_ret = ret;
}
return prog_ret;
}
cl_int CLVK_API_CALL clUnloadPlatformCompiler(cl_platform_id platform) {
TRACE_FUNCTION("platform", (uintptr_t)platform);
LOG_API_CALL("platform = %p", platform);
if (!is_valid_platform(platform)) {
return CL_INVALID_PLATFORM;
}
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clUnloadCompiler() {
TRACE_FUNCTION();
LOG_API_CALL("%s", "");
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetProgramInfo(cl_program prog,
cl_program_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("program", (uintptr_t)prog, "param_name", param_name);
LOG_API_CALL("program = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
prog, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_context val_context;
size_t val_sizet;
cl_bool val_bool;
api_query_string val_string;
std::vector<size_t> val_sizet_vec;
std::vector<cl_device_id> val_devices;
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
bool program_built_once = false;
for (auto dev : program->devices()) {
if (program->build_status(dev) == CL_BUILD_SUCCESS) {
program_built_once = true;
}
}
switch (param_name) {
case CL_PROGRAM_NUM_DEVICES:
val_uint = program->num_devices();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_PROGRAM_REFERENCE_COUNT:
val_uint = program->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_PROGRAM_CONTEXT:
val_context = program->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_PROGRAM_DEVICES:
for (auto dev : program->devices()) {
val_devices.push_back(const_cast<cvk_device*>(dev));
}
copy_ptr = val_devices.data();
ret_size = sizeof(cl_device_id) * val_devices.size();
break;
case CL_PROGRAM_NUM_KERNELS:
if (!program_built_once) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
val_sizet = program->num_kernels();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_PROGRAM_SOURCE:
copy_ptr = program->source().c_str();
ret_size = program->source().size() + 1;
break;
case CL_PROGRAM_KERNEL_NAMES: {
if (!program_built_once) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
val_string = "";
std::string sep = "";
for (auto kname : program->kernel_names()) {
val_string += sep + kname;
sep = ";";
}
copy_ptr = val_string.c_str();
ret_size = val_string.size_with_null();
break;
}
case CL_PROGRAM_BINARY_SIZES:
for (uint32_t i = 0; i < program->num_devices(); i++) {
val_sizet_vec.push_back(program->binary_size());
}
copy_ptr = val_sizet_vec.data();
ret_size = val_sizet_vec.size() * sizeof(size_t);
break;
case CL_PROGRAM_BINARIES:
ret_size = program->num_devices() * sizeof(unsigned char*);
if (param_value != nullptr) {
for (uint32_t i = 0; i < program->num_devices(); i++) {
auto dst = static_cast<unsigned char**>(param_value)[i];
if (dst != nullptr) {
auto success = program->write(dst);
if (!success) {
ret = CL_OUT_OF_RESOURCES;
break;
}
}
}
}
break;
case CL_PROGRAM_IL:
copy_ptr = program->il().data();
ret_size = program->il().size();
break;
case CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT:
case CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT:
if (!program_built_once) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
val_bool = CL_FALSE;
copy_ptr = &val_bool;
ret_size = sizeof(val_bool);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr) &&
(param_name != CL_PROGRAM_BINARIES)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clGetProgramBuildInfo(cl_program prog, cl_device_id dev,
cl_program_build_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("program", (uintptr_t)prog, "device", (uintptr_t)dev,
"param_name", param_name);
LOG_API_CALL("program = %p, device = %p, param_name = %x, param_value_size "
"= %zu, param_value = %p, param_value_size_ret = %p",
prog, dev, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_build_status val_status;
api_query_string val_string;
cl_program_binary_type val_binarytype;
size_t val_sizet;
auto device = icd_downcast(dev);
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
if (!is_valid_device(device)) { // TODO check the program knows the device
return CL_INVALID_DEVICE;
}
switch (param_name) {
case CL_PROGRAM_BUILD_STATUS:
val_status = program->build_status(device);
copy_ptr = &val_status;
ret_size = sizeof(val_status);
break;
case CL_PROGRAM_BUILD_LOG:
copy_ptr = program->build_log(device).c_str();
ret_size = program->build_log(device).size() + 1;
break;
case CL_PROGRAM_BUILD_OPTIONS:
copy_ptr = program->build_options().c_str();
ret_size = program->build_options().size() + 1;
break;
case CL_PROGRAM_BINARY_TYPE:
val_binarytype = program->binary_type(device);
copy_ptr = &val_binarytype;
ret_size = sizeof(val_binarytype);
break;
case CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE:
val_sizet = 0;
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (ret_size > param_value_size)) {
ret = CL_INVALID_VALUE;
} else if ((param_value != nullptr) && (copy_ptr != nullptr)) {
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clRetainProgram(cl_program program) {
TRACE_FUNCTION("program", (uintptr_t)program);
LOG_API_CALL("program = %p", program);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
icd_downcast(program)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseProgram(cl_program program) {
TRACE_FUNCTION("program", (uintptr_t)program);
LOG_API_CALL("program = %p", program);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
icd_downcast(program)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clSetProgramReleaseCallback(
cl_program program,
void(CL_CALLBACK* pfn_notify)(cl_program program, void* user_data),
void* user_data) {
TRACE_FUNCTION("program", (uintptr_t)program);
LOG_API_CALL("program = %p, pfn_notify = %p, user_data = %p", program,
pfn_notify, user_data);
return CL_INVALID_OPERATION;
}
// Kernel Object APIs
cl_kernel cvk_create_kernel(cl_program program, const char* kernel_name,
cl_int* errcode_ret) {
if (kernel_name == nullptr) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
auto kernel =
std::make_unique<cvk_kernel>(icd_downcast(program), kernel_name);
*errcode_ret = kernel->init();
if (*errcode_ret != CL_SUCCESS) {
return nullptr;
} else {
return kernel.release();
}
}
cl_kernel CLVK_API_CALL clCreateKernel(cl_program prog, const char* kernel_name,
cl_int* errcode_ret) {
TRACE_FUNCTION("program", (uintptr_t)prog, "kernel_name",
TRACE_STRING(kernel_name));
LOG_API_CALL("program = %p, kernel_name = %s", prog, kernel_name);
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_PROGRAM;
}
return nullptr;
}
if (program->build_status() != CL_BUILD_SUCCESS) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_PROGRAM_EXECUTABLE;
}
return nullptr;
}
cl_int err;
cl_kernel ret = cvk_create_kernel(prog, kernel_name, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return ret;
}
cl_int CLVK_API_CALL clCreateKernelsInProgram(cl_program prog,
cl_uint num_kernels,
cl_kernel* kernels,
cl_uint* num_kernels_ret) {
TRACE_FUNCTION("program", (uintptr_t)prog, "num_kernels", num_kernels);
LOG_API_CALL(
"program = %p, num_kernels = %u, kernels = %p, num_kernels_ret = %p",
prog, num_kernels, kernels, num_kernels_ret);
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
if (program->build_status() != CL_BUILD_SUCCESS) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
cl_uint num_kernels_in_program = program->num_kernels();
if ((kernels != nullptr) && (num_kernels < num_kernels_in_program)) {
return CL_INVALID_VALUE;
}
if (kernels != nullptr) {
cl_uint i = 0;
cl_int err;
for (auto& kname : program->kernel_names()) {
kernels[i] = cvk_create_kernel(program, kname, &err);
if (err != CL_SUCCESS) {
return err;
}
++i;
}
}
if (num_kernels_ret != nullptr) {
*num_kernels_ret = num_kernels_in_program;
}
return CL_SUCCESS;
}
cl_kernel CLVK_API_CALL clCloneKernel(cl_kernel source_kernel,
cl_int* errcode_ret) {
TRACE_FUNCTION("source_kernel", (uintptr_t)source_kernel);
LOG_API_CALL("kernel = %p, errcode_ret = %p", source_kernel, errcode_ret);
if (!is_valid_kernel(source_kernel)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_KERNEL;
}
return nullptr;
}
cl_int err;
auto kernel = icd_downcast(source_kernel)->clone(&err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
if (err != CL_SUCCESS) {
return nullptr;
} else {
return kernel.release();
}
}
cl_int CLVK_API_CALL clSetKernelArg(cl_kernel kern, cl_uint arg_index,
size_t arg_size, const void* arg_value) {
TRACE_FUNCTION("kernel", (uintptr_t)kern, "arg_index", arg_index);
LOG_API_CALL("kernel = %p, arg_index = %u, arg_size = %zu, arg_value = %p",
kern, arg_index, arg_size, arg_value);
auto kernel = icd_downcast(kern);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
// TODO CL_INVALID_ARG_VALUE if arg_value specified is not a valid value.
// TODO CL_INVALID_ARG_SIZE if arg_size does not match the size of the data
// type for an argument that is not a memory object or if the argument is a
// memory object and arg_size != sizeof(cl_mem) or if arg_size is zero and
// the argument is declared with the __local qualifier or if the argument is
// a sampler and arg_size != sizeof(cl_sampler).
// TODO CL_INVALID_ARG_VALUE if the argument is an image declared with the
// read_only qualifier and arg_value refers to an image object created with
// cl_mem_flags of CL_MEM_WRITE or if the image argument is declared with
// the write_only qualifier and arg_value refers to an image object created
// with cl_mem_flags of CL_MEM_READ.
if (arg_index >= kernel->num_args()) {
cvk_error_fn("the program has only %u arguments", kernel->num_args());
return CL_INVALID_ARG_INDEX;
}
// CL_INVALID_MEM_OBJECT and CL_INVALID_SAMPLER are handled in
// cvk_kernel_argument_values::set_arg.
// With opaque pointers, clspv is unable to infer the type of an unused
// kernel argument so allow nullptr for its value. It will not have an
// affect on the kernel's operation.
auto arg_kind = kernel->arg_kind(arg_index);
if ((arg_value == nullptr) &&
!((arg_kind == kernel_argument_kind::local) ||
(arg_kind == kernel_argument_kind::unused))) {
cvk_error_fn("passing a null pointer to clSetKernelArg is only "
"supported for local arguments");
if (arg_kind == kernel_argument_kind::sampler) {
return CL_INVALID_SAMPLER;
} else {
return CL_INVALID_ARG_VALUE;
}
}
if (arg_size == 0 &&
(kernel->arg_kind(arg_index) == kernel_argument_kind::local)) {
return CL_INVALID_ARG_SIZE;
}
return kernel->set_arg(arg_index, arg_size, arg_value);
}
cl_int CLVK_API_CALL clSetKernelExecInfo(cl_kernel kernel,
cl_kernel_exec_info param_name,
size_t param_value_size,
const void* param_value) {
TRACE_FUNCTION("kernel", (uintptr_t)kernel);
LOG_API_CALL("kernel = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p",
kernel, param_name, param_value_size, param_value);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clGetKernelInfo(cl_kernel kern, cl_kernel_info param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("kernel", (uintptr_t)kern, "param_name", param_name);
LOG_API_CALL("kernel = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
kern, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_context val_context;
cl_program val_program;
api_query_string val_string;
auto kernel = icd_downcast(kern);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
switch (param_name) {
case CL_KERNEL_REFERENCE_COUNT:
val_uint = kernel->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_KERNEL_CONTEXT:
val_context = kernel->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_KERNEL_FUNCTION_NAME:
copy_ptr = kernel->name().c_str();
ret_size = kernel->name().size() + 1;
break;
case CL_KERNEL_NUM_ARGS:
val_uint = kernel->num_args();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_KERNEL_PROGRAM:
val_program = kernel->program();
copy_ptr = &val_program;
ret_size = sizeof(val_program);
break;
case CL_KERNEL_ATTRIBUTES: {
val_string = kernel->attributes();
copy_ptr = val_string.c_str();
ret_size = val_string.size_with_null();
break;
}
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clGetKernelArgInfo(cl_kernel kern, cl_uint arg_index,
cl_kernel_arg_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("kernel", (uintptr_t)kern, "arg_index", arg_index,
"param_name", param_name);
LOG_API_CALL("kernel = %p, arg_index = %u, param_name = %x, "
"param_value_size = %zu, param_value = %p, "
"param_value_size_ret = %p",
kern, arg_index, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
const void* copy_ptr = nullptr;
size_t ret_size = 0;
api_query_string val_string;
cl_kernel_arg_address_qualifier val_address_qualifier;
cl_kernel_arg_access_qualifier val_access_qualifier;
cl_kernel_arg_type_qualifier val_type_qualifier;
auto kernel = icd_downcast(kern);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
if (arg_index >= kernel->num_args()) {
return CL_INVALID_ARG_INDEX;
}
if ((param_name != CL_KERNEL_ARG_NAME) &&
!kernel->has_extended_arg_info(arg_index)) {
return CL_KERNEL_ARG_INFO_NOT_AVAILABLE;
}
switch (param_name) {
case CL_KERNEL_ARG_NAME:
val_string = kernel->arg_name(arg_index);
copy_ptr = val_string.data();
ret_size = val_string.size_with_null();
break;
case CL_KERNEL_ARG_ADDRESS_QUALIFIER:
val_address_qualifier = kernel->arg_address_qualifier(arg_index);
copy_ptr = &val_address_qualifier;
ret_size = sizeof(val_address_qualifier);
break;
case CL_KERNEL_ARG_TYPE_NAME:
val_string = kernel->arg_type_name(arg_index);
copy_ptr = val_string.data();
ret_size = val_string.size_with_null();
break;
case CL_KERNEL_ARG_ACCESS_QUALIFIER:
val_access_qualifier = kernel->arg_access_qualifier(arg_index);
copy_ptr = &val_access_qualifier;
ret_size = sizeof(val_access_qualifier);
break;
case CL_KERNEL_ARG_TYPE_QUALIFIER:
val_type_qualifier = kernel->arg_type_qualifier(arg_index);
copy_ptr = &val_type_qualifier;
ret_size = sizeof(val_type_qualifier);
break;
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clGetKernelWorkGroupInfo(
cl_kernel kern, cl_device_id dev, cl_kernel_work_group_info param_name,
size_t param_value_size, void* param_value, size_t* param_value_size_ret) {
TRACE_FUNCTION("kernel", (uintptr_t)kern, "device", (uintptr_t)dev,
"param_name", param_name);
LOG_API_CALL(
"kernel = %p, device = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
kern, dev, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
const void* copy_ptr = nullptr;
size_t val_sizet, ret_size = 0;
cl_ulong val_ulong;
std::array<size_t, 3> val_wgs;
auto device = icd_downcast(dev);
auto kernel = icd_downcast(kern);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
switch (param_name) {
case CL_KERNEL_WORK_GROUP_SIZE:
val_sizet = kernel->max_work_group_size(device);
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
val_sizet = device->preferred_work_group_size_multiple();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_KERNEL_LOCAL_MEM_SIZE:
val_ulong = kernel->local_mem_size();
copy_ptr = &val_ulong;
ret_size = sizeof(val_ulong);
break;
case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: {
auto const& val_wgs_uint = kernel->required_work_group_size();
val_wgs[0] = val_wgs_uint[0];
val_wgs[1] = val_wgs_uint[1];
val_wgs[2] = val_wgs_uint[2];
copy_ptr = val_wgs.data();
ret_size = sizeof(val_wgs);
break;
}
case CL_KERNEL_PRIVATE_MEM_SIZE: // TODO
// Return 0 as it is a lower bound of the private memory size needed by
// a kernel.
val_ulong = 0;
copy_ptr = &val_ulong;
ret_size = sizeof(val_ulong);
break;
case CL_KERNEL_GLOBAL_WORK_SIZE: // TODO
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clGetKernelSubGroupInfo(
cl_kernel kern, cl_device_id dev, cl_kernel_sub_group_info param_name,
size_t input_value_size, const void* input_value, size_t param_value_size,
void* param_value, size_t* param_value_size_ret) {
TRACE_FUNCTION("kernel", (uintptr_t)kern, "device", (uintptr_t)dev,
"param_name", param_name);
LOG_API_CALL("kernel = %p, device = %p, param_name = %x, input_value_size "
"= %zu, input_value = %p, param_value_size = %zu, param_value "
"= %p, param_value_size_ret = %p",
kern, dev, param_name, input_value_size, input_value,
param_value_size, param_value, param_value_size_ret);
cl_int ret = CL_SUCCESS;
const void* copy_ptr = nullptr;
size_t val_sizet, ret_size = 0;
std::array<size_t, 3> val_lws;
auto device = icd_downcast(dev);
auto kernel = icd_downcast(kern);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
switch (param_name) {
case CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE:
val_sizet = kernel->max_sub_group_size_for_ndrange(device);
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE: {
std::array<uint32_t, 3> lws = {1, 1, 1};
unsigned num_dims = input_value_size / sizeof(size_t);
if (input_value_size % sizeof(size_t) != 0) {
ret = CL_INVALID_VALUE;
break;
}
for (unsigned dim = 0; dim < num_dims; dim++) {
lws[dim] = static_cast<const size_t*>(input_value)[dim];
}
val_sizet = kernel->sub_group_count_for_ndrange(device, lws);
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
}
case CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT: {
if (input_value_size % sizeof(size_t) != 0) {
ret = CL_INVALID_VALUE;
break;
}
auto num_sub_groups = *static_cast<const size_t*>(input_value);
val_lws =
kernel->local_size_for_sub_group_count(device, num_sub_groups);
copy_ptr = &val_lws;
ret_size = param_value_size;
break;
}
case CL_KERNEL_MAX_NUM_SUB_GROUPS:
val_sizet = kernel->max_num_sub_groups(device);
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_KERNEL_COMPILE_NUM_SUB_GROUPS: // TODO
default:
ret = CL_INVALID_VALUE;
break;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int CLVK_API_CALL clRetainKernel(cl_kernel kernel) {
TRACE_FUNCTION("kernel", (uintptr_t)kernel);
LOG_API_CALL("kernel = %p", kernel);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
icd_downcast(kernel)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseKernel(cl_kernel kernel) {
TRACE_FUNCTION("kernel", (uintptr_t)kernel);
LOG_API_CALL("kernel = %p", kernel);
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
icd_downcast(kernel)->release();
return CL_SUCCESS;
}
/* Profiling APIs */
cl_int CLVK_API_CALL clGetEventProfilingInfo(cl_event evt,
cl_profiling_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("event", (uintptr_t)evt, "param_name", param_name);
LOG_API_CALL("event = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
evt, param_name, param_value_size, param_value,
param_value_size_ret);
auto event = icd_downcast(evt);
if (!is_valid_event(event)) {
return CL_INVALID_EVENT;
}
switch (param_name) {
case CL_PROFILING_COMMAND_QUEUED:
case CL_PROFILING_COMMAND_SUBMIT:
case CL_PROFILING_COMMAND_START:
case CL_PROFILING_COMMAND_END:
case CL_PROFILING_COMMAND_COMPLETE:
break;
default:
return CL_INVALID_VALUE;
}
if ((param_value_size < sizeof(cl_ulong)) && (param_value != nullptr)) {
return CL_INVALID_VALUE;
}
if (event->is_user_event() || (event->get_status() != CL_COMPLETE)) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
if (!event->queue()->has_property(CL_QUEUE_PROFILING_ENABLE)) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = sizeof(cl_ulong);
}
if (param_value != nullptr) {
if (param_name == CL_PROFILING_COMMAND_COMPLETE) {
param_name = CL_PROFILING_COMMAND_END;
}
cl_ulong value = event->get_profiling_info(param_name);
memcpy(param_value, &value, sizeof(cl_ulong));
}
return CL_SUCCESS;
}
/* Flush and Finish APIs */
cl_int CLVK_API_CALL clFlush(cl_command_queue command_queue) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
return icd_downcast(command_queue)->flush();
}
cl_int CLVK_API_CALL clFinish(cl_command_queue command_queue) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
return icd_downcast(command_queue)->finish();
}
/* Enqueued Commands APIs */
cl_int CLVK_API_CALL clEnqueueReadBuffer(cl_command_queue cq, cl_mem buf,
cl_bool blocking_read, size_t offset,
size_t size, void* ptr,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"blocking_read", blocking_read, "offset", offset, "size",
size, "num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, buffer = %p, blocking = %d, offset = "
"%zu, size = %zu, ptr = %p",
cq, buf, blocking_read, offset, size, ptr);
auto buffer = icd_downcast(buf);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_buffer(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (buffer->has_any_flag(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_buffer_host_copy(
command_queue, CL_COMMAND_READ_BUFFER, static_cast<cvk_buffer*>(buffer),
ptr, offset, size);
auto err = command_queue->enqueue_command_with_deps(
cmd, blocking_read, num_events_in_wait_list, event_wait_list, event);
return err;
}
cl_int CLVK_API_CALL clEnqueueWriteBuffer(cl_command_queue cq, cl_mem buf,
cl_bool blocking_write, size_t offset,
size_t size, const void* ptr,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"blocking_write", blocking_write, "offset", offset, "size",
size, "num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, buffer = %p, blocking = %d, offset = "
"%zu, size = %zu, ptr = %p",
cq, buf, blocking_write, offset, size, ptr);
auto buffer = icd_downcast(buf);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
// TODO validate the contexts
if (!is_valid_buffer(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (buffer->has_any_flag(CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_buffer_host_copy(
command_queue, CL_COMMAND_WRITE_BUFFER,
static_cast<cvk_buffer*>(buffer), ptr, offset, size);
auto err = command_queue->enqueue_command_with_deps(
cmd, blocking_write, num_events_in_wait_list, event_wait_list, event);
return err;
}
cl_int CLVK_API_CALL clEnqueueReadBufferRect(
cl_command_queue cq, cl_mem buf, cl_bool blocking_read,
const size_t* buffer_origin, const size_t* host_origin,
const size_t* region, size_t buffer_row_pitch, size_t buffer_slice_pitch,
size_t host_row_pitch, size_t host_slice_pitch, void* ptr,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"blocking_read", blocking_read, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, buffer = %p, blocking = %d", cq, buf,
blocking_read);
LOG_API_CALL("buffer_origin = {%zu,%zu,%zu}, host_origin = {%zu,%zu,%zu}, "
"region = {%zu,%zu,%zu}",
buffer_origin[0], buffer_origin[1], buffer_origin[2],
host_origin[0], host_origin[1], host_origin[2], region[0],
region[1], region[2]);
LOG_API_CALL("buffer_row_pitch = %zu, buffer_slice_pitch = %zu,"
"host_row_pitch = %zu, host_slice_pitch = %zu",
buffer_row_pitch, buffer_slice_pitch, host_row_pitch,
host_slice_pitch);
LOG_API_CALL("ptr = %p, num_events = %u, event_wait_list = %p, event = %p",
ptr, num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
auto buffer = static_cast<cvk_buffer*>(buf);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_buffer(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (buffer->has_any_flag(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_copy_host_buffer_rect(
command_queue, CL_COMMAND_READ_BUFFER_RECT, buffer, ptr, host_origin,
buffer_origin, region, host_row_pitch, host_slice_pitch,
buffer_row_pitch, buffer_slice_pitch);
auto err = command_queue->enqueue_command_with_deps(
cmd, blocking_read, num_events_in_wait_list, event_wait_list, event);
return err;
}
cl_int CLVK_API_CALL clEnqueueWriteBufferRect(
cl_command_queue cq, cl_mem buf, cl_bool blocking_write,
const size_t* buffer_origin, const size_t* host_origin,
const size_t* region, size_t buffer_row_pitch, size_t buffer_slice_pitch,
size_t host_row_pitch, size_t host_slice_pitch, const void* ptr,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"blocking_write", blocking_write, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, buffer = %p, blocking = %d", cq, buf,
blocking_write);
LOG_API_CALL("buffer_origin = {%zu,%zu,%zu}, host_origin = {%zu,%zu,%zu}, "
"region = {%zu,%zu,%zu}",
buffer_origin[0], buffer_origin[1], buffer_origin[2],
host_origin[0], host_origin[1], host_origin[2], region[0],
region[1], region[2]);
LOG_API_CALL("buffer_row_pitch = %zu, buffer_slice_pitch = %zu, "
"host_row_pitch = %zu, host_slice_pitch = %zu",
buffer_row_pitch, buffer_slice_pitch, host_row_pitch,
host_slice_pitch);
LOG_API_CALL("ptr = %p, num_events = %u, event_wait_list = %p, event = %p",
ptr, num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
auto buffer = static_cast<cvk_buffer*>(buf);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_buffer(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (buffer->has_any_flag(CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_copy_host_buffer_rect(
command_queue, CL_COMMAND_WRITE_BUFFER_RECT, buffer,
const_cast<void*>(ptr), host_origin, buffer_origin, region,
host_row_pitch, host_slice_pitch, buffer_row_pitch, buffer_slice_pitch);
auto err = command_queue->enqueue_command_with_deps(
cmd, blocking_write, num_events_in_wait_list, event_wait_list, event);
return err;
}
cl_int CLVK_API_CALL clEnqueueFillBuffer(
cl_command_queue cq, cl_mem buf, const void* pattern, size_t pattern_size,
size_t offset, size_t size, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"offset", offset, "size", size, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL(
"command_queue = %p, buffer = %p, pattern = %p, pattern_size = %zu,"
"offset = %zu, size = %zu, num_events = %u, event_wait_list = %p, "
"event = %p",
cq, buf, pattern, pattern_size, offset, size, num_events_in_wait_list,
event_wait_list, event);
auto buffer = icd_downcast(buf);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_buffer(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
// TODO check context
// TODO check buffer bounds
if (pattern == nullptr) {
return CL_INVALID_VALUE;
}
// Check the pattern size is valid
size_t valid_pattern_sizes[] = {1, 2, 4, 8, 16, 32, 64, 128};
bool pattern_size_valid = false;
for (auto size : valid_pattern_sizes) {
if (size == pattern_size) {
pattern_size_valid = true;
break;
}
}
if (!pattern_size_valid) {
return CL_INVALID_VALUE;
}
// Check that offset and size are a multiple of pattern_size
if ((offset % pattern_size != 0) || (size % pattern_size != 0)) {
return CL_INVALID_VALUE;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
// TODO check sub-buffer alignment
auto cmd = new cvk_command_fill_buffer(
command_queue, static_cast<cvk_buffer*>(buffer), offset, size, pattern,
pattern_size, CL_COMMAND_FILL_BUFFER);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueCopyBuffer(cl_command_queue cq, cl_mem srcbuf,
cl_mem dstbuf, size_t src_offset,
size_t dst_offset, size_t size,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "src_buffer",
(uintptr_t)srcbuf, "dst_buffer", (uintptr_t)dstbuf,
"src_offset", src_offset, "dst_offset", dst_offset, "size",
size, "num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, src_buffer = %p, dst_buffer = %p, "
"src_offset = %zu,"
"dst_offset = %zu, size = %zu, num_events = %u, "
"event_wait_list = %p, event = %p",
cq, srcbuf, dstbuf, src_offset, dst_offset, size,
num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
auto src_buffer = icd_downcast(srcbuf);
auto dst_buffer = icd_downcast(dstbuf);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_same_context(cq, srcbuf) || !is_same_context(cq, dstbuf)) {
return CL_INVALID_CONTEXT;
}
if (!is_same_context(cq, num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (!is_valid_buffer(src_buffer) || !is_valid_buffer(dst_buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto cmd = new cvk_command_copy_buffer(
command_queue, CL_COMMAND_COPY_BUFFER,
static_cast<cvk_buffer*>(src_buffer),
static_cast<cvk_buffer*>(dst_buffer), src_offset, dst_offset, size);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueCopyBufferRect(
cl_command_queue cq, cl_mem src_buffer, cl_mem dst_buffer,
const size_t* src_origin, const size_t* dst_origin, const size_t* region,
size_t src_row_pitch, size_t src_slice_pitch, size_t dst_row_pitch,
size_t dst_slice_pitch, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "src_buffer",
(uintptr_t)src_buffer, "dst_buffer", (uintptr_t)dst_buffer,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, src_buffer = %p, dst_buffer = %p, "
"src_origin = {%zu,%zu,%zu}, dst_origin = {%zu,%zu,%zu}, "
"region = {%zu,%zu,%zu}, src_row_pitch = %zu, "
"src_slice_pitch = %zu, dst_row_pitch = %zu, "
"dst_slice_pitch = %zu, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
cq, src_buffer, dst_buffer, src_origin[0], src_origin[1],
src_origin[2], dst_origin[0], dst_origin[1], dst_origin[2],
region[0], region[1], region[2], src_row_pitch,
src_slice_pitch, dst_row_pitch, dst_slice_pitch,
num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
// TODO CL_INVALID_COMMAND_QUEUE if command_queue is not a valid
// command-queue.
if (!is_valid_buffer(src_buffer) || !is_valid_buffer(dst_buffer)) {
return CL_INVALID_MEM_OBJECT;
}
// TODO CL_INVALID_VALUE if (src_offset, region) or (dst_offset, region)
// require accessing elements outside the src_buffer and dst_buffer objects
// respectively.
// TODO CL_INVALID_VALUE if any region array element is 0.
// TODO CL_INVALID_VALUE if src_row_pitch is not 0 and is less than
// region[0].
// TODO CL_INVALID_VALUE if dst_row_pitch is not 0 and is less than
// region[0].
// TODO CL_INVALID_VALUE if src_slice_pitch is not 0 and is less than
// region[1] * src_row_pitch or if src_slice_pitch is not 0 and is not a
// multiple of src_row_pitch.
// TODO CL_INVALID_VALUE if dst_slice_pitch is not 0 and is less than
// region[1] * dst_row_pitch or if dst_slice_pitch is not 0 and is not a
// multiple of dst_row_pitch.
// TODO CL_INVALID_VALUE if src_buffer and dst_buffer are the same buffer
// object and src_slice_pitch is not equal to dst_slice_pitch and
// src_row_pitch is not equal to dst_row_pitch.
//
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, src_buffer) ||
!is_same_context(command_queue, dst_buffer) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_MEM_COPY_OVERLAP if src_buffer and dst_buffer are the same buffer
// object and the source and destination regions overlap or if src_buffer
// and dst_buffer are different sub-buffers of the same associated buffer
// object and they overlap. Refer to Appendix E in the OpenCL specification
// for details on how to determine if source and destination regions
// overlap.
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if src_buffer is a sub-buffer object
// and offset specified when the sub-buffer object is created is not aligned
// to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if dst_buffer is a sub-buffer object
// and offset specified when the sub-buffer object is created is not aligned
// to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with src_buffer or dst_buffer.
// TODO CL_OUT_OF_RESOURCES if there is a failure to allocate resources
// required by the OpenCL implementation on the device.
// TODO CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources
// required by the OpenCL implementation on the host.
//
auto srcbuf = static_cast<cvk_buffer*>(src_buffer);
auto dstbuf = static_cast<cvk_buffer*>(dst_buffer);
auto cmd = new cvk_command_copy_buffer_rect(
command_queue, srcbuf, dstbuf, src_origin, dst_origin, region,
src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
void* cvk_enqueue_map_buffer(cvk_command_queue* cq, cvk_buffer* buffer,
cl_bool blocking_map, size_t offset, size_t size,
cl_map_flags map_flags,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event,
cl_int* errcode_ret, cl_command_type type,
cvk_image* image = nullptr) {
auto cmd = new cvk_command_map_buffer(cq, buffer, offset, size, map_flags,
type, image);
void* map_ptr;
cl_int err = cmd->build(&map_ptr);
// FIXME This error cannot occur for objects created with
// CL_MEM_USE_HOST_PTR or CL_MEM_ALLOC_HOST_PTR.
if (err != CL_SUCCESS) {
*errcode_ret = CL_MAP_FAILURE;
return nullptr;
}
err = cq->enqueue_command_with_deps(
cmd, blocking_map, num_events_in_wait_list, event_wait_list, event);
*errcode_ret = err;
return map_ptr;
}
void* CLVK_API_CALL clEnqueueMapBuffer(cl_command_queue cq, cl_mem buf,
cl_bool blocking_map,
cl_map_flags map_flags, size_t offset,
size_t size,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event, cl_int* errcode_ret) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "buffer", (uintptr_t)buf,
"blocking_map", blocking_map, "map_flags", map_flags,
"offset", offset, "size", size, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, buffer = %p, offset = %zu, size = %zu",
cq, buf, offset, size);
auto command_queue = icd_downcast(cq);
auto buffer = static_cast<cvk_buffer*>(buf);
if (!is_valid_command_queue(command_queue)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_COMMAND_QUEUE;
}
return nullptr;
}
if (!is_valid_buffer(buffer)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_MEM_OBJECT;
}
return nullptr;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_EVENT_WAIT_LIST;
}
return nullptr;
}
if (!is_same_context(command_queue, buffer)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_CONTEXT;
}
return nullptr;
}
if ((size == 0) || (offset + size > buffer->size())) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
if (!map_flags_are_valid(map_flags)) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if buffer is a sub-buffer object and
// offset specified when the sub-buffer object is created is not aligned to
// CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
if ((map_flags & CL_MAP_READ) &&
(buffer->has_any_flag(CL_MEM_HOST_WRITE_ONLY |
CL_MEM_HOST_NO_ACCESS))) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_OPERATION;
}
return nullptr;
}
if (((map_flags & CL_MAP_WRITE) ||
(map_flags & CL_MAP_WRITE_INVALIDATE_REGION)) &&
(buffer->has_any_flag(CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS))) {
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_OPERATION;
}
return nullptr;
}
cl_int err;
auto map_ptr = cvk_enqueue_map_buffer(
command_queue, buffer, blocking_map, offset, size, map_flags,
num_events_in_wait_list, event_wait_list, event, &err,
CL_COMMAND_MAP_BUFFER);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return map_ptr;
}
cl_int cvk_enqueue_unmap_image(cvk_command_queue* cq, cvk_image* image,
void* mapped_ptr, bool update_host_ptr,
bool blocking, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
const cvk_image_mapping mapping = image->mapping_for(mapped_ptr);
// We need to hold on the buffer to make sure it has not been release by the
// unmap command buffer creating the copy command which will hold on to it.
cvk_buffer_holder buffer(mapping.buffer);
bool needs_copy =
(mapping.flags & (CL_MAP_WRITE | CL_MAP_WRITE_INVALIDATE_REGION)) != 0;
update_host_ptr = update_host_ptr && image->has_flags(CL_MEM_USE_HOST_PTR);
auto cmd_unmap =
std::make_unique<cvk_command_unmap_image>(cq, image, mapped_ptr);
_cl_event* evt_unmap;
auto err = cq->enqueue_command_with_deps(
cmd_unmap.release(), blocking && !needs_copy, num_events_in_wait_list,
event_wait_list, &evt_unmap);
if (err != CL_SUCCESS) {
return err;
}
if (needs_copy) {
_cl_event* evt_pre_copy = evt_unmap;
_cl_event* evt_host_ptr_update;
if (update_host_ptr) {
size_t zero_origin[3] = {0, 0, 0};
auto cmd_host_ptr_update =
std::make_unique<cvk_command_copy_host_buffer_rect>(
cq, CL_COMMAND_WRITE_BUFFER_RECT, mapping.buffer,
image->host_ptr(), mapping.origin.data(), zero_origin,
mapping.region.data(), image->row_pitch(),
image->slice_pitch(), image->map_buffer_row_pitch(mapping),
image->map_buffer_slice_pitch(mapping),
image->element_size());
err =
cq->enqueue_command_with_deps(cmd_host_ptr_update.release(), 1,
&evt_unmap, &evt_host_ptr_update);
if (err != CL_SUCCESS) {
return err;
}
evt_pre_copy = evt_host_ptr_update;
}
_cl_event* evt_copy;
auto cmd_copy = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_BUFFER_TO_IMAGE, cq, mapping.buffer, image, 0,
mapping.origin, mapping.region);
err = cq->enqueue_command_with_deps(cmd_copy.release(), blocking, 1,
&evt_pre_copy, &evt_copy);
if (err != CL_SUCCESS) {
return err;
}
if (event != nullptr) {
*event = new cvk_event_combine(
cq->context(), CL_COMMAND_UNMAP_MEM_OBJECT, cq,
icd_downcast(evt_unmap), icd_downcast(evt_copy));
}
icd_downcast(evt_unmap)->release();
if (update_host_ptr) {
icd_downcast(evt_host_ptr_update)->release();
}
icd_downcast(evt_copy)->release();
} else {
if (event != nullptr) {
*event = evt_unmap;
} else {
icd_downcast(evt_unmap)->release();
}
}
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clEnqueueUnmapMemObject(cl_command_queue cq, cl_mem mem,
void* mapped_ptr,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "memobj", (uintptr_t)mem,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, memobj = %p, mapped_ptr = %p", cq, mem,
mapped_ptr);
auto command_queue = icd_downcast(cq);
auto memobj = icd_downcast(mem);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_mem_object(memobj)) {
return CL_INVALID_MEM_OBJECT;
}
if (memobj->is_image_type()) {
auto image = static_cast<cvk_image*>(memobj);
if (image->is_backed_by_buffer_view()) {
auto buffer = static_cast<cvk_buffer*>(image->buffer());
auto cmd =
new cvk_command_unmap_buffer(command_queue, buffer, mapped_ptr);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
} else {
return cvk_enqueue_unmap_image(command_queue, image, mapped_ptr,
true, false, num_events_in_wait_list,
event_wait_list, event);
}
} else {
auto buffer = static_cast<cvk_buffer*>(memobj);
auto cmd = std::make_unique<cvk_command_unmap_buffer>(
command_queue, buffer, mapped_ptr);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
}
}
cl_int cvk_enqueue_ndrange_kernel(cvk_command_queue* command_queue,
cvk_kernel* kernel, uint32_t dims,
const cvk_ndrange& ndrange,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
// TODO check that it's a host command queue
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, kernel)) {
return CL_INVALID_CONTEXT;
}
if (!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
auto device = command_queue->device();
if ((dims < 1) || (dims > device->max_work_item_dimensions())) {
return CL_INVALID_WORK_DIMENSION;
}
auto program = kernel->program();
if (program->binary_type(device) != CL_PROGRAM_BINARY_TYPE_EXECUTABLE) {
return CL_INVALID_PROGRAM_EXECUTABLE;
}
if (!kernel->args_valid()) {
return CL_INVALID_KERNEL_ARGS;
}
// TODO CL_INVALID_GLOBAL_WORK_SIZE if any of the values specified in
// global_work_size[0], … global_work_size[work_dim - 1] exceed the maximum
// value representable by size_t on the device on which the kernel-instance
// will be enqueued.
// TODO CL_INVALID_GLOBAL_OFFSET if the value specified in global_work_size
// + the corresponding values in global_work_offset for any dimensions is
// greater than the maximum value representable by size t on the device on
// which the kernel-instance will be enqueued.
// TODO CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and is
// not consistent with the required number of sub-groups for kernel in the
// program source.
// TODO CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and the
// total number of work-items in the work-group computed as
// local_work_size[0] × … local_work_size[work_dim - 1] is greater than the
// value specified by CL_KERNEL_WORK_GROUP_SIZE in the Kernel Object Device
// Queries table.
// TODO CL_INVALID_WORK_GROUP_SIZE if the program was compiled with
// cl-uniform-work-group-size and the number of work-items specified by
// global_work_size is not evenly divisible by size of work-group given by
// local_work_size or by the required work-group size specified in the
// kernel source.
// TODO CL_INVALID_WORK_ITEM_SIZE if the number of work-items specified in
// any of local_work_size[0], … local_work_size[work_dim - 1] is greater
// than the corresponding values specified by
// CL_DEVICE_MAX_WORK_ITEM_SIZES[0], …,
// CL_DEVICE_MAX_WORK_ITEM_SIZES[work_dim - 1].
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if a sub-buffer object is specified
// as the value for an argument that is a buffer object and the offset
// specified when the sub-buffer object is created is not aligned to
// CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
// TODO CL_INVALID_IMAGE_SIZE if an image object is specified as an argument
// value and the image dimensions (image width, height, specified or compute
// row and/or slice pitch) are not supported by device associated with
// queue.
// TODO CL_IMAGE_FORMAT_NOT_SUPPORTED if an image object is specified as an
// argument value and the image format (image channel order and data type)
// is not supported by device associated with queue.
// TODO CL_OUT_OF_RESOURCES if there is a failure to queue the execution
// instance of kernel on the command-queue because of insufficient resources
// needed to execute the kernel. For example, the explicitly specified
// local_work_size causes a failure to execute the kernel because of
// insufficient resources such as registers or local memory. Another example
// would be the number of read-only image args used in kernel exceed the
// CL_DEVICE_MAX_READ_IMAGE_ARGS value for device or the number of
// write-only and read-write image args used in kernel exceed the
// CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS value for device or the number of
// samplers used in kernel exceed CL_DEVICE_MAX_SAMPLERS for device.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with image or buffer objects specified
// as arguments to kernel.
// TODO CL_INVALID_OPERATION if SVM pointers are passed as arguments to a
// kernel and the device does not support SVM or if system pointers are
// passed as arguments to a kernel and/or stored inside SVM allocations
// passed as kernel arguments and the device does not support fine grain
// system SVM allocations.
// Check work-group size matches the required size if specified
auto reqd_work_group_size = kernel->required_work_group_size();
if (reqd_work_group_size[0] != 0) {
if (reqd_work_group_size != ndrange.lws) {
cvk_error_fn("work-group size does not match the required size");
return CL_INVALID_WORK_GROUP_SIZE;
}
}
// Check uniformity of the NDRange if needed
if (!command_queue->device()->supports_non_uniform_workgroup()) {
if (!ndrange.is_uniform()) {
cvk_error_fn("non uniform workgroup not supported");
return CL_INVALID_WORK_GROUP_SIZE;
}
}
auto cmd = new cvk_command_kernel(command_queue, kernel, dims, ndrange);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueTask(cl_command_queue command_queue,
cl_kernel kernel,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue, "kernel",
(uintptr_t)kernel, "num_events_in_wait_list",
(uintptr_t)num_events_in_wait_list);
LOG_API_CALL(
"command_queue = %p, kernel = %p, num_events_in_wait_list = %d,"
" event_wait_list = %p, event = %p",
command_queue, kernel, num_events_in_wait_list, event_wait_list, event);
cvk_ndrange ndrange;
ndrange.gws = {1, 1, 1};
return cvk_enqueue_ndrange_kernel(
icd_downcast(command_queue), icd_downcast(kernel), 1, ndrange,
num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueNDRangeKernel(
cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim,
const size_t* global_work_offset, const size_t* global_work_size,
const size_t* local_work_size, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue, "kernel",
(uintptr_t)kernel, "num_events_in_wait_list",
(uintptr_t)num_events_in_wait_list);
LOG_API_CALL(
"command_queue = %p, kernel = %p, work_dim = %u, "
"num_events_in_wait_list = %u, event_wait_list = %p, event = %p",
command_queue, kernel, work_dim, num_events_in_wait_list,
event_wait_list, event);
static const size_t null_global_work_size[3] = {0};
if (global_work_size == nullptr) {
global_work_size = null_global_work_size;
}
cvk_ndrange ndrange(work_dim, global_work_offset, global_work_size,
local_work_size);
// Try to pick a sensible work-group size if the user didn't specify one.
if (local_work_size == nullptr) {
icd_downcast(command_queue)
->device()
->select_work_group_size(icd_downcast(kernel), ndrange.gws,
ndrange.lws);
cvk_info_fn("selected local work size: {%u,%u,%u}", ndrange.lws[0],
ndrange.lws[1], ndrange.lws[2]);
}
LOG_API_CALL("goff = {%u,%u,%u}", ndrange.offset[0], ndrange.offset[1],
ndrange.offset[2]);
LOG_API_CALL("gws = {%u,%u,%u}", ndrange.gws[0], ndrange.gws[1],
ndrange.gws[2]);
LOG_API_CALL("lws = {%u,%u,%u}", ndrange.lws[0], ndrange.lws[1],
ndrange.lws[2]);
return cvk_enqueue_ndrange_kernel(
icd_downcast(command_queue), icd_downcast(kernel), work_dim, ndrange,
num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueNativeKernel(
cl_command_queue command_queue, void(CL_CALLBACK* user_func)(void*),
void* args, size_t cb_args, cl_uint num_mem_objects, const cl_mem* mem_list,
const void** args_mem_loc, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue,
"num_events_in_wait_list",
(uintptr_t)num_events_in_wait_list);
LOG_API_CALL(
"command_queue = %p, user_func = %p, args = %p, cb_args = %zu, "
"num_mem_objects = %u, mem_list = %p, args_mem_loc = %p, "
"num_events_in_wait_list = %u, event_wait_list = %p, event = %p",
command_queue, user_func, args, cb_args, num_mem_objects, mem_list,
args_mem_loc, num_events_in_wait_list, event_wait_list, event);
return CL_INVALID_OPERATION;
}
cl_sampler cvk_create_sampler(cl_context context, cl_bool normalized_coords,
cl_addressing_mode addressing_mode,
cl_filter_mode filter_mode,
std::vector<cl_sampler_properties>&& properties,
cl_int* errcode_ret) {
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
auto ctx = icd_downcast(context);
if (!ctx->device()->supports_images()) {
*errcode_ret = CL_INVALID_OPERATION;
return nullptr;
}
auto sampler = cvk_sampler::create(ctx, normalized_coords, addressing_mode,
filter_mode, std::move(properties));
if (sampler == nullptr) {
*errcode_ret = CL_OUT_OF_RESOURCES;
} else {
*errcode_ret = CL_SUCCESS;
}
return sampler;
}
cl_sampler CLVK_API_CALL clCreateSampler(cl_context context,
cl_bool normalized_coords,
cl_addressing_mode addressing_mode,
cl_filter_mode filter_mode,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "normalized_coords",
normalized_coords, "addressing_mode", addressing_mode,
"filter_mode", filter_mode);
LOG_API_CALL("context = %p, normalized_coords = %d, addressing_mode = %d, "
"filter_mode = %d, errcode_ret = %p",
context, normalized_coords, addressing_mode, filter_mode,
errcode_ret);
std::vector<cl_sampler_properties> properties;
cl_int err;
auto sampler =
cvk_create_sampler(context, normalized_coords, addressing_mode,
filter_mode, std::move(properties), &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return sampler;
}
cl_sampler CLVK_API_CALL clCreateSamplerWithProperties(
cl_context context, const cl_sampler_properties* sampler_properties,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, sampler_properties = %p, errcode_ret = %p",
context, sampler_properties, errcode_ret);
cl_bool normalized_coords = CL_TRUE;
cl_addressing_mode addressing_mode = CL_ADDRESS_CLAMP;
cl_filter_mode filter_mode = CL_FILTER_NEAREST;
std::vector<cl_sampler_properties> properties;
if (sampler_properties) {
while (*sampler_properties) {
auto key = *sampler_properties;
auto value = *(sampler_properties + 1);
switch (key) {
case CL_SAMPLER_NORMALIZED_COORDS:
normalized_coords = value;
break;
case CL_SAMPLER_ADDRESSING_MODE:
addressing_mode = value;
break;
case CL_SAMPLER_FILTER_MODE:
filter_mode = value;
break;
default:
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_VALUE;
}
return nullptr;
}
properties.push_back(key);
properties.push_back(value);
sampler_properties += 2;
}
properties.push_back(0);
}
cl_int err;
auto sampler =
cvk_create_sampler(context, normalized_coords, addressing_mode,
filter_mode, std::move(properties), &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return sampler;
}
cl_int CLVK_API_CALL clRetainSampler(cl_sampler sampler) {
TRACE_FUNCTION("sampler", (uintptr_t)sampler);
LOG_API_CALL("sampler = %p", sampler);
if (!is_valid_sampler(sampler)) {
return CL_INVALID_SAMPLER;
}
icd_downcast(sampler)->retain();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clReleaseSampler(cl_sampler sampler) {
TRACE_FUNCTION("sampler", (uintptr_t)sampler);
LOG_API_CALL("sampler = %p", sampler);
if (!is_valid_sampler(sampler)) {
return CL_INVALID_SAMPLER;
}
icd_downcast(sampler)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clGetSamplerInfo(cl_sampler samp,
cl_sampler_info param_name,
size_t param_value_size,
void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("sampler", (uintptr_t)samp, "param_name", param_name);
LOG_API_CALL("sampler = %p, param_name = %d, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
samp, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_uint val_uint;
cl_context val_context;
cl_bool val_bool;
cl_addressing_mode val_addressing_mode;
cl_filter_mode val_filter_mode;
auto sampler = icd_downcast(samp);
if (!is_valid_sampler(sampler)) {
return CL_INVALID_SAMPLER;
}
switch (param_name) {
case CL_SAMPLER_REFERENCE_COUNT:
val_uint = sampler->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_SAMPLER_CONTEXT:
val_context = sampler->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_SAMPLER_NORMALIZED_COORDS:
val_bool = sampler->normalized_coords();
copy_ptr = &val_bool;
ret_size = sizeof(val_bool);
break;
case CL_SAMPLER_ADDRESSING_MODE:
val_addressing_mode = sampler->addressing_mode();
copy_ptr = &val_addressing_mode;
ret_size = sizeof(val_addressing_mode);
break;
case CL_SAMPLER_FILTER_MODE:
val_filter_mode = sampler->filter_mode();
copy_ptr = &val_filter_mode;
ret_size = sizeof(val_filter_mode);
break;
case CL_SAMPLER_PROPERTIES:
copy_ptr = sampler->properties().data();
ret_size = sampler->properties().size() * sizeof(cl_sampler_properties);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_mem cvk_create_image(cl_context context, cl_mem_flags flags,
const cl_image_format* image_format,
const cl_image_desc* image_desc, void* host_ptr,
std::vector<cl_mem_properties>&& properties,
cl_int* errcode_ret) {
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
// TODO CL_INVALID_VALUE if values specified in flags are not valid.
// TODO CL_INVALID_IMAGE_FORMAT_DESCRIPTOR if values specified in
// image_format are not valid or if image_format is NULL.
// TODO CL_INVALID_IMAGE_DESCRIPTOR if values specified in image_desc are
// not valid or if image_desc is NULL.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions specified in image_desc
// exceed the minimum maximum image dimensions described in the table of
// allowed values for param_name for clGetDeviceInfo for all devices in
// context.
// TODO CL_INVALID_HOST_PTR if host_ptr in image_desc is NULL and
// CL_MEM_USE_HOST_PTR or CL_MEM_COPY_HOST_PTR are set in flags or if
// host_ptr is not NULL but CL_MEM_COPY_HOST_PTR or CL_MEM_USE_HOST_PTR are
// not set in flags.
// TODO CL_INVALID_VALUE if a 1D image buffer is being created and the
// buffer object was created with CL_MEM_WRITE_ONLY and flags specifies
// CL_MEM_READ_WRITE or CL_MEM_READ_ONLY, or if the buffer object was
// created with CL_MEM_READ_ONLY and flags specifies CL_MEM_READ_WRITE or
// CL_MEM_WRITE_ONLY, or if flags specifies CL_MEM_USE_HOST_PTR or
// CL_MEM_ALLOC_HOST_PTR or CL_MEM_COPY_HOST_PTR.
// TODO CL_INVALID_VALUE if a 1D image buffer is being created and the
// buffer object was created with CL_MEM_HOST_WRITE_ONLY and flags specifies
// CL_MEM_HOST_READ_ONLY, or if the buffer object was created with
// CL_MEM_HOST_READ_ONLY and flags specifies CL_MEM_HOST_WRITE_ONLY, or if
// the buffer object was created with CL_MEM_HOST_NO_ACCESS and flags
// specifies CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_WRITE_ONLY.
// TODO CL_IMAGE_FORMAT_NOT_SUPPORTED if the image_format is not supported.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for image object.
// TODO support creating 2D images from buffers
if ((image_desc->image_type == CL_MEM_OBJECT_IMAGE2D) &&
(image_desc->mem_object != nullptr)) {
*errcode_ret = CL_INVALID_OPERATION;
return nullptr;
}
auto image =
cvk_image::create(icd_downcast(context), flags, image_desc,
image_format, host_ptr, std::move(properties));
*errcode_ret = (image != nullptr)
? CL_SUCCESS
: CL_OUT_OF_RESOURCES; // FIXME do this properly
return image;
}
cl_mem cvk_create_image(cl_context context, cl_mem_flags flags,
const cl_image_format* image_format,
const cl_image_desc* image_desc, void* host_ptr,
cl_int* errcode_ret) {
std::vector<cl_mem_properties> properties;
return cvk_create_image(context, flags, image_format, image_desc, host_ptr,
std::move(properties), errcode_ret);
}
cl_mem CLVK_API_CALL clCreateImage(cl_context context, cl_mem_flags flags,
const cl_image_format* image_format,
const cl_image_desc* image_desc,
void* host_ptr, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL(
"context = %p, flags = %lu, image_format = %p, image_desc = %p,"
" host_ptr = %p, errcode_ret = %p",
context, flags, image_format, image_desc, host_ptr, errcode_ret);
cl_int err;
auto image = cvk_create_image(context, flags, image_format, image_desc,
host_ptr, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return image;
}
cl_mem CLVK_API_CALL clCreateImageWithProperties(
cl_context context, const cl_mem_properties* properties, cl_mem_flags flags,
const cl_image_format* image_format, const cl_image_desc* image_desc,
void* host_ptr, cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL("context = %p, properties = %p, flags = %lx, image_format = "
"%p, image_desc = %p, host_ptr = %p, errcode_ret = %p",
context, properties, flags, image_format, image_desc, host_ptr,
errcode_ret);
cl_int err;
std::vector<cl_mem_properties> props;
if (properties != nullptr) {
while (*properties) {
props.push_back(*properties);
properties++;
}
props.push_back(0);
}
auto image = cvk_create_image(context, flags, image_format, image_desc,
host_ptr, std::move(props), &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return image;
}
cl_mem CLVK_API_CALL clCreateImage2D(cl_context context, cl_mem_flags flags,
const cl_image_format* image_format,
size_t image_width, size_t image_height,
size_t image_row_pitch, void* host_ptr,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL(
"context = %p, flags = %lu, image_format = %p, image_width = %zu, "
"image_height = %zu, image_row_pitch = %zu, host_ptr = %p, "
"errcode_ret = %p",
context, flags, image_format, image_width, image_height,
image_row_pitch, host_ptr, errcode_ret);
cl_image_desc desc = {
CL_MEM_OBJECT_IMAGE2D,
image_width,
image_height,
0, // image_depth
0, // image_array_size
image_row_pitch,
0, // image_slice_pitch
0, // num_mip_levels
0, // num_samples
{nullptr} // buffer
};
cl_int err;
auto image =
cvk_create_image(context, flags, image_format, &desc, host_ptr, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return image;
}
cl_mem CLVK_API_CALL clCreateImage3D(cl_context context, cl_mem_flags flags,
const cl_image_format* image_format,
size_t image_width, size_t image_height,
size_t image_depth, size_t image_row_pitch,
size_t image_slice_pitch, void* host_ptr,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags);
LOG_API_CALL(
"context = %p, flags = %lu, image_format = %p, image_width = %zu, "
"image_height = %zu, image_depth = %zu, image_row_pitch = %zu, "
"image_slice_pitch = %zu, host_ptr = %p, errcode_ret = %p",
context, flags, image_format, image_width, image_height, image_depth,
image_row_pitch, image_slice_pitch, host_ptr, errcode_ret);
cl_image_desc desc = {
CL_MEM_OBJECT_IMAGE3D,
image_width,
image_height,
image_depth,
0, // image_array_size
image_row_pitch,
image_slice_pitch,
0, // num_mip_levels
0, // num_samples
{nullptr} // buffer
};
cl_int err;
auto image =
cvk_create_image(context, flags, image_format, &desc, host_ptr, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return image;
}
cl_int CLVK_API_CALL clGetImageInfo(cl_mem image, cl_image_info param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("image", (uintptr_t)image, "param_name", param_name);
LOG_API_CALL("image = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
image, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_image_format val_image_format;
size_t val_sizet;
cl_mem val_mem;
cl_uint val_uint;
if (!is_valid_image(image)) {
return CL_INVALID_MEM_OBJECT;
}
auto img = static_cast<cvk_image*>(image);
switch (param_name) {
case CL_IMAGE_FORMAT:
val_image_format = img->format();
copy_ptr = &val_image_format;
ret_size = sizeof(val_image_format);
break;
case CL_IMAGE_ELEMENT_SIZE:
val_sizet = img->element_size();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_ROW_PITCH:
val_sizet = img->row_pitch();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_SLICE_PITCH:
val_sizet = img->slice_pitch();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_WIDTH:
val_sizet = img->width();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_HEIGHT:
val_sizet = img->height();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_DEPTH:
val_sizet = img->depth();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_ARRAY_SIZE:
val_sizet = img->array_size();
copy_ptr = &val_sizet;
ret_size = sizeof(val_sizet);
break;
case CL_IMAGE_BUFFER:
val_mem = img->buffer();
copy_ptr = &val_mem;
ret_size = sizeof(val_mem);
break;
case CL_IMAGE_NUM_MIP_LEVELS:
val_uint = img->num_mip_levels();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_IMAGE_NUM_SAMPLES:
val_uint = img->num_samples();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
bool operator!=(const VkComponentMapping& lhs, const VkComponentMapping& rhs) {
return lhs.r != rhs.r || lhs.g != rhs.g || lhs.b != rhs.b || lhs.a != rhs.a;
}
static bool is_image_format_supported(
VkPhysicalDevice& pdev, VkFormat format, cl_mem_object_type image_type,
const VkFormatFeatureFlags& required_format_feature_flags,
cl_channel_order image_channel_order) {
VkFormatProperties properties;
vkGetPhysicalDeviceFormatProperties(pdev, format, &properties);
cvk_debug("Vulkan format %d:", format);
cvk_debug(
" linear : %s",
vulkan_format_features_string(properties.linearTilingFeatures).c_str());
cvk_debug(" optimal: %s",
vulkan_format_features_string(properties.optimalTilingFeatures)
.c_str());
cvk_debug(" buffer : %s",
vulkan_format_features_string(properties.bufferFeatures).c_str());
cvk_debug(
"Required format features %s",
vulkan_format_features_string(required_format_feature_flags).c_str());
VkFormatFeatureFlags features;
if (image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER) {
features = properties.bufferFeatures;
} else {
// TODO support linear tiling
features = properties.optimalTilingFeatures;
}
if ((features & required_format_feature_flags) ==
required_format_feature_flags) {
if ((image_channel_order == CL_LUMINANCE ||
image_channel_order == CL_INTENSITY) &&
(image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER)) {
return false;
}
return true;
}
return false;
}
cl_int CLVK_API_CALL clGetSupportedImageFormats(cl_context context,
cl_mem_flags flags,
cl_mem_object_type image_type,
cl_uint num_entries,
cl_image_format* image_formats,
cl_uint* num_image_formats) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags, "image_type",
image_type, "num_entries", num_entries);
LOG_API_CALL(
"context = %p, flags = %lu, image_type = %d, num_entries = %u, "
"image_formats = %p, num_image_formats = %p",
context, flags, image_type, num_entries, image_formats,
num_image_formats);
if (!is_valid_context(context)) {
return CL_INVALID_CONTEXT;
}
if ((num_entries == 0) && (image_formats != nullptr)) {
return CL_INVALID_VALUE;
}
if (!cvk_mem::is_image_type(image_type)) {
return CL_INVALID_VALUE;
}
cl_uint num_formats_found = 0;
auto dev = icd_downcast(context)->device();
auto pdev = dev->vulkan_physical_device();
if (!dev->supports_read_write_images() &&
(flags & CL_MEM_KERNEL_READ_AND_WRITE)) {
if (num_image_formats != nullptr) {
*num_image_formats = 0;
}
return CL_SUCCESS;
}
const VkFormatFeatureFlags required_format_feature_flags =
cvk_image::required_format_feature_flags_for(image_type, flags);
// TODO tiling selection
// No host access => OPTIMAL
// Host ACCESS => LINEAR if supported, OPTIMAL otherwise?
// Iterate over all known CL/VK format associations and report the CL
// formats for which the Vulkan format is supported
for (auto mapping : get_format_maps()) {
VkComponentMapping components_sampled, components_storage;
image_format_support fmt_support;
cl_image_format clfmt = mapping.first;
if (!cl_image_format_to_vulkan_format(clfmt, image_type, dev,
&fmt_support, &components_sampled,
&components_storage)) {
continue;
}
if ((fmt_support.flags & flags) != flags) {
continue;
}
if (!is_image_format_supported(pdev, fmt_support.vkfmt, image_type,
required_format_feature_flags,
clfmt.image_channel_order)) {
continue;
}
// image format is supported
if ((image_formats != nullptr) && (num_formats_found < num_entries)) {
image_formats[num_formats_found] = clfmt;
cvk_debug_fn(
"reporting image format {%s, %s}",
cl_channel_order_to_string(clfmt.image_channel_order).c_str(),
cl_channel_type_to_string(clfmt.image_channel_data_type)
.c_str());
}
num_formats_found++;
}
cvk_debug_fn("reporting %u formats", num_formats_found);
if (num_image_formats != nullptr) {
*num_image_formats = num_formats_found;
}
return CL_SUCCESS;
}
cl_int cvk_enqueue_map_image(cl_command_queue cq, cl_mem img,
bool user_map_image, cl_bool blocking_map,
cl_map_flags map_flags, const size_t* origin,
const size_t* region, size_t* image_row_pitch,
size_t* image_slice_pitch, cvk_buffer** map_buffer,
void** map_ptr, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
auto command_queue = icd_downcast(cq);
auto image = static_cast<cvk_image*>(img);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(image)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (!map_flags_are_valid(map_flags)) {
return CL_INVALID_VALUE;
}
// TODO CL_INVALID_VALUE if region being mapped given by (origin,
// origin+region) is out of bounds
// TODO CL_INVALID_VALUE if values in origin and region do not follow rules
// described in the argument description for origin and region.
if (image_row_pitch == nullptr) {
return CL_INVALID_VALUE;
}
switch (image->type()) {
case CL_MEM_OBJECT_IMAGE3D:
case CL_MEM_OBJECT_IMAGE1D_ARRAY:
case CL_MEM_OBJECT_IMAGE2D_ARRAY:
if (image_slice_pitch == nullptr) {
return CL_INVALID_VALUE;
}
break;
default:
break;
}
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for image are not supported
// by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for image are not supported by device associated with queue.
// TODO CL_MAP_FAILURE if there is a failure to map the requested region
// into the host address space. This error cannot occur for image objects
// created with CL_MEM_USE_HOST_PTR or CL_MEM_ALLOC_HOST_PTR.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with buffer.
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
if ((map_flags & CL_MAP_READ) &&
(image->has_any_flag(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS))) {
return CL_INVALID_OPERATION;
}
if (((map_flags & CL_MAP_WRITE) ||
(map_flags & CL_MAP_WRITE_INVALIDATE_REGION)) &&
(image->has_any_flag(CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS)) &&
user_map_image) {
return CL_INVALID_OPERATION;
}
std::array<size_t, 3> orig = {origin[0], origin[1], origin[2]};
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
bool needs_copy = (map_flags & CL_MAP_WRITE_INVALIDATE_REGION) == 0;
bool update_host_ptr =
user_map_image && image->has_flags(CL_MEM_USE_HOST_PTR);
// Get a mapping
cvk_image_mapping mapping;
if (!image->find_or_create_mapping(mapping, orig, reg, map_flags,
update_host_ptr)) {
cvk_error("cannot find or create a mapping");
return CL_OUT_OF_RESOURCES;
}
if (map_ptr != nullptr) {
*map_ptr = mapping.ptr;
}
if (map_buffer != nullptr) {
*map_buffer = mapping.buffer;
}
if (update_host_ptr) {
*image_row_pitch = image->row_pitch();
if (image_slice_pitch != nullptr) {
*image_slice_pitch = image->slice_pitch();
}
} else {
*image_row_pitch = image->map_buffer_row_pitch(mapping);
if (image_slice_pitch != nullptr) {
*image_slice_pitch = image->map_buffer_slice_pitch(mapping);
}
}
if (needs_copy) {
_cl_event* evt_copy;
auto cmd_copy = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_IMAGE_TO_BUFFER, command_queue, mapping.buffer,
image, 0, orig, reg);
auto err = command_queue->enqueue_command_with_deps(
cmd_copy.release(), blocking_map && !update_host_ptr,
num_events_in_wait_list, event_wait_list, &evt_copy);
if (err != CL_SUCCESS) {
image->cleanup_mapping(mapping);
return err;
}
if (update_host_ptr) {
size_t zero_origin[3] = {0, 0, 0};
auto cmd_host_ptr_update =
std::make_unique<cvk_command_copy_host_buffer_rect>(
command_queue, CL_COMMAND_READ_BUFFER_RECT, mapping.buffer,
image->host_ptr(), orig.data(), zero_origin, reg.data(),
image->row_pitch(), image->slice_pitch(),
image->map_buffer_row_pitch(reg),
image->map_buffer_slice_pitch(reg), image->element_size());
_cl_event* evt_host_ptr_update;
err = command_queue->enqueue_command_with_deps(
cmd_host_ptr_update.release(), blocking_map, 1, &evt_copy,
&evt_host_ptr_update);
if (err != CL_SUCCESS) {
image->cleanup_mapping(mapping);
return err;
}
if (event != nullptr) {
*event = new cvk_event_combine(
command_queue->context(), CL_COMMAND_MAP_IMAGE,
command_queue, icd_downcast(evt_copy),
icd_downcast(evt_host_ptr_update));
}
icd_downcast(evt_copy)->release();
icd_downcast(evt_host_ptr_update)->release();
} else {
if (event != nullptr) {
*event = evt_copy;
} else {
icd_downcast(evt_copy)->release();
}
}
} else {
auto cmd_map =
std::make_unique<cvk_command_map_image>(command_queue, image);
auto err = command_queue->enqueue_command_with_deps(
cmd_map.release(), blocking_map, num_events_in_wait_list,
event_wait_list, event);
if (err != CL_SUCCESS) {
image->cleanup_mapping(mapping);
return err;
}
}
return CL_SUCCESS;
}
cl_int cvk_enqueue_image_copy(
cvk_command_queue* queue, cl_command_type command_type, cvk_mem* image,
bool blocking, const size_t* origin, const size_t* region, size_t row_pitch,
size_t slice_pitch, void* ptr, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
auto img = static_cast<cvk_image*>(image);
if (img->is_backed_by_buffer_view()) {
auto cmd = new cvk_command_buffer_host_copy(
queue, command_type, static_cast<cvk_buffer*>(img->buffer()), ptr,
origin[0] * img->element_size(), region[0] * img->element_size());
auto err = queue->enqueue_command_with_deps(
cmd, blocking, num_events_in_wait_list, event_wait_list, event);
return err;
}
// Create image map command
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
cl_map_flags map_flags;
if (command_type == CL_COMMAND_WRITE_IMAGE) {
map_flags = CL_MAP_WRITE_INVALIDATE_REGION;
} else {
map_flags = CL_MAP_READ;
}
// Create copy command
auto rpitch = row_pitch;
if (rpitch == 0) {
rpitch = region[0] * img->element_size();
}
auto spitch = slice_pitch;
if (spitch == 0) {
spitch = region[1] * rpitch;
}
const size_t zero_origin[3] = {0, 0, 0};
_cl_event* evt_map;
cvk_buffer* map_buffer;
void* map_ptr;
size_t image_row_pitch, image_slice_pitch;
auto err = cvk_enqueue_map_image(
queue, image, false, false, map_flags, origin, region, &image_row_pitch,
&image_slice_pitch, &map_buffer, &map_ptr, num_events_in_wait_list,
event_wait_list, &evt_map);
if (err != CL_SUCCESS) {
return err;
}
auto cmd_copy = std::make_unique<cvk_command_copy_host_buffer_rect>(
queue, command_type, map_buffer, ptr, zero_origin, zero_origin, region,
rpitch, spitch, img->map_buffer_row_pitch(reg),
img->map_buffer_slice_pitch(reg), img->element_size());
_cl_event* evt_copy;
err = queue->enqueue_command_with_deps(cmd_copy.release(), 1, &evt_map,
&evt_copy);
if (err != CL_SUCCESS) {
return err;
}
_cl_event* evt_unmap;
err = cvk_enqueue_unmap_image(queue, img, map_ptr, false, blocking, 1,
&evt_copy, &evt_unmap);
icd_downcast(evt_copy)->release();
if (err != CL_SUCCESS) {
return err;
}
if (event != nullptr) {
*event = new cvk_event_combine(queue->context(), command_type, queue,
icd_downcast(evt_map),
icd_downcast(evt_unmap));
}
icd_downcast(evt_map)->release();
icd_downcast(evt_unmap)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clEnqueueReadImage(
cl_command_queue cq, cl_mem img, cl_bool blocking_read,
const size_t* origin, const size_t* region, size_t row_pitch,
size_t slice_pitch, void* ptr, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "image", (uintptr_t)img,
"blocking_read", blocking_read, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, image = %p, blocking_read = %d, "
"origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"row_pitch = %zu, slice_pitch = %zu, ptr = %p, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p",
cq, img, blocking_read, origin[0], origin[1], origin[2],
region[0], region[1], region[2], row_pitch, slice_pitch, ptr,
num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
auto image = icd_downcast(img);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(image)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_INVALID_VALUE if the region being read specified by origin and
// region is out of bounds or if ptr is a NULL value.
// TODO CL_INVALID_VALUE if values in origin and region do not follow rules
// described in the argument description for origin and region.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for image are not supported
// by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for image are not supported by device associated with queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with image.
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
if (image->has_any_flag(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
return cvk_enqueue_image_copy(command_queue, CL_COMMAND_READ_IMAGE, image,
blocking_read, origin, region, row_pitch,
slice_pitch, ptr, num_events_in_wait_list,
event_wait_list, event);
}
cl_int CLVK_API_CALL clEnqueueWriteImage(
cl_command_queue cq, cl_mem img, cl_bool blocking_write,
const size_t* origin, const size_t* region, size_t input_row_pitch,
size_t input_slice_pitch, const void* ptr, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "image", (uintptr_t)img,
"blocking_write", blocking_write, "num_events_in_wait_list",
num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, image = %p, blocking_write = %d, "
"origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"input_row_pitch = %zu, input_slice_pitch = %zu, ptr = %p, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p",
cq, img, blocking_write, origin[0], origin[1], origin[2],
region[0], region[1], region[2], input_row_pitch,
input_slice_pitch, ptr, num_events_in_wait_list,
event_wait_list, event);
auto command_queue = icd_downcast(cq);
auto image = icd_downcast(img);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(image)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_INVALID_VALUE if the region being written specified by origin and
// region is out of bounds or if ptr is a NULL value.
// TODO CL_INVALID_VALUE if values in origin and region do not follow rules
// described in the argument description for origin and region.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for image are not supported
// by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for image are not supported by device associated with queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with image.
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
if (image->has_any_flag(CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
return cvk_enqueue_image_copy(
command_queue, CL_COMMAND_WRITE_IMAGE, image, blocking_write, origin,
region, input_row_pitch, input_slice_pitch, const_cast<void*>(ptr),
num_events_in_wait_list, event_wait_list, event);
}
cl_int CLVK_API_CALL
clEnqueueCopyImage(cl_command_queue cq, cl_mem src_image, cl_mem dst_image,
const size_t* src_origin, const size_t* dst_origin,
const size_t* region, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "src_image",
(uintptr_t)src_image, "dst_image", (uintptr_t)dst_image,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, src_image = %p, dst_image = %p, "
"src_origin = {%zu,%zu,%zu}, dst_origin = {%zu, %zu, %zu}, "
"region = {%zu, %zu, %zu}, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p",
cq, src_image, dst_image, src_origin[0], src_origin[1],
src_origin[2], dst_origin[0], dst_origin[1], dst_origin[2],
region[0], region[1], region[2], num_events_in_wait_list,
event_wait_list, event);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(src_image) || !is_valid_image(dst_image)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_same_context(command_queue, src_image) ||
!is_same_context(command_queue, dst_image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
auto src_img = static_cast<cvk_image*>(src_image);
auto dst_img = static_cast<cvk_image*>(dst_image);
// TODO CL_INVALID_VALUE if the 2D or 3D rectangular region specified by
// src_origin and src_origin + region refers to a region outside src_image,
// or if the 2D or 3D rectangular region specified by dst_origin and
// dst_origin + region refers to a region outside dst_image.
// TODO CL_INVALID_VALUE if values in src_origin, dst_origin and region do
// not follow rules described in the argument description for src_origin,
// dst_origin and region.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for src_image or dst_image
// are not supported by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for src_image or dst_image are not supported by device
// associated with queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with src_image or dst_image.
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
// TODO CL_MEM_COPY_OVERLAP if src_image and dst_image are the same image
// object and the source and destination regions overlap.
if (!src_img->has_same_format(dst_img)) {
return CL_IMAGE_FORMAT_MISMATCH;
}
std::array<size_t, 3> src_orig = {src_origin[0], src_origin[1],
src_origin[2]};
std::array<size_t, 3> dst_orig = {dst_origin[0], dst_origin[1],
dst_origin[2]};
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
if (src_img->is_backed_by_buffer_view() &&
dst_img->is_backed_by_buffer_view()) {
auto cmd = new cvk_command_copy_buffer(
command_queue, CL_COMMAND_COPY_IMAGE,
static_cast<cvk_buffer*>(src_img->buffer()),
static_cast<cvk_buffer*>(dst_img->buffer()),
src_origin[0] * src_img->element_size(),
dst_origin[0] * dst_img->element_size(),
region[0] * src_img->element_size());
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
} else if (src_img->is_backed_by_buffer_view()) {
auto cmd = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_IMAGE, CL_COMMAND_COPY_BUFFER_TO_IMAGE,
command_queue, static_cast<cvk_buffer*>(src_img->buffer()), dst_img,
src_origin[0] * src_img->element_size(), dst_orig, reg);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
} else if (dst_img->is_backed_by_buffer_view()) {
auto cmd = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_IMAGE, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
command_queue, static_cast<cvk_buffer*>(dst_img->buffer()), src_img,
dst_origin[0] * dst_img->element_size(), src_orig, reg);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
} else {
auto cmd = std::make_unique<cvk_command_image_image_copy>(
command_queue, src_img, dst_img, src_orig, dst_orig, reg);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
}
}
cl_int CLVK_API_CALL clEnqueueFillImage(
cl_command_queue cq, cl_mem image, const void* fill_color,
const size_t* origin, const size_t* region, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "image", (uintptr_t)image,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, image = %p, fill_color = %p, "
"origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p",
cq, image, fill_color, origin[0], origin[1], origin[2],
region[0], region[1], region[2], num_events_in_wait_list,
event_wait_list, event);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(image)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (fill_color == nullptr) {
return CL_INVALID_VALUE;
}
// TODO CL_INVALID_VALUE if the region being written specified by origin and
// region is out of bounds or if ptr is a NULL value.
// TODO CL_INVALID_VALUE if values in origin and region do not follow rules
// described in the argument description for origin and region.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for image are not supported
// by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for image are not supported by device associated with queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with image.
// TODO use Vulkan clear commands when possible
// TODO use a shader when better
auto img = static_cast<cvk_image*>(image);
// Fill
cvk_image::fill_pattern_array pattern;
size_t pattern_size;
img->prepare_fill_pattern(fill_color, pattern, &pattern_size);
if (img->is_backed_by_buffer_view()) {
auto cmd = new cvk_command_fill_buffer(
command_queue, static_cast<cvk_buffer*>(img->buffer()),
origin[0] * img->element_size(), region[0] * img->element_size(),
pattern.data(), pattern_size, CL_COMMAND_FILL_IMAGE);
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
}
// Create image map command
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
void* map_ptr;
_cl_event* evt_map;
size_t image_row_pitch, image_slice_pitch;
auto err = cvk_enqueue_map_image(
command_queue, image, false, false, CL_MAP_WRITE_INVALIDATE_REGION,
origin, region, &image_row_pitch, &image_slice_pitch, nullptr, &map_ptr,
num_events_in_wait_list, event_wait_list, &evt_map);
if (err != CL_SUCCESS) {
return err;
}
auto cmd_fill = std::make_unique<cvk_command_fill_image>(
command_queue, map_ptr, pattern, pattern_size, reg);
_cl_event* evt_fill;
err = command_queue->enqueue_command_with_deps(cmd_fill.release(), 1,
&evt_map, &evt_fill);
if (err != CL_SUCCESS) {
return err;
}
_cl_event* evt_unmap;
err = cvk_enqueue_unmap_image(command_queue, img, map_ptr, false, false, 1,
&evt_fill, &evt_unmap);
icd_downcast(evt_fill)->release();
if (err != CL_SUCCESS) {
return err;
}
if (event != nullptr) {
*event = new cvk_event_combine(
command_queue->context(), CL_COMMAND_FILL_IMAGE, command_queue,
icd_downcast(evt_map), icd_downcast(evt_unmap));
}
icd_downcast(evt_map)->release();
icd_downcast(evt_unmap)->release();
return CL_SUCCESS;
}
cl_int CLVK_API_CALL clEnqueueCopyImageToBuffer(
cl_command_queue cq, cl_mem src_image, cl_mem dst_buffer,
const size_t* src_origin, const size_t* region, size_t dst_offset,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "src_image",
(uintptr_t)src_image, "dst_buffer", (uintptr_t)dst_buffer,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, src_image = %p, dst_buffer = %p, "
"src_origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"dst_offset = %zu, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
cq, src_image, dst_buffer, src_origin[0], src_origin[1],
src_origin[2], region[0], region[1], region[2], dst_offset,
num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(src_image) || !is_valid_buffer(dst_buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, src_image) ||
!is_same_context(command_queue, dst_buffer) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_INVALID_MEM_OBJECT if src_image is a 1D image buffer object
// created from dst_buffer.
// TODO CL_INVALID_VALUE if the 1D, 2D, or 3D rectangular region specified
// by src_origin and src_origin + region refers to a region outside
// src_image, or if the region specified by dst_offset and dst_offset +
// dst_cb refers to a region outside dst_buffer.
// TODO CL_INVALID_VALUE if values in src_origin and region do not follow
// rules described in the argument description for src_origin and region.
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if dst_buffer is a sub-buffer object
// and offset specified when the sub-buffer object is created is not aligned
// to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for src_image are not
// supported by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for src_image are not supported by device associated with
// queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with src_image or dst_buffer.
auto image = static_cast<cvk_image*>(src_image);
auto buffer = static_cast<cvk_buffer*>(dst_buffer);
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
std::array<size_t, 3> origin = {src_origin[0], src_origin[1],
src_origin[2]};
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
if (image->is_backed_by_buffer_view()) {
auto cmd = new cvk_command_copy_buffer(
command_queue, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
static_cast<cvk_buffer*>(image->buffer()), buffer,
src_origin[0] * image->element_size(), dst_offset,
region[0] * image->element_size());
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
} else {
auto cmd = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_IMAGE_TO_BUFFER, command_queue, buffer, image,
dst_offset, origin, reg);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
}
}
cl_int CLVK_API_CALL clEnqueueCopyBufferToImage(
cl_command_queue cq, cl_mem src_buffer, cl_mem dst_image, size_t src_offset,
const size_t* dst_origin, const size_t* region,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "src_buffer",
(uintptr_t)src_buffer, "dst_image", (uintptr_t)dst_image,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, src_buffer = %p, dst_image = %p, "
"src_offset = %zu, "
"dst_origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p",
cq, src_buffer, dst_image, src_offset, dst_origin[0],
dst_origin[1], dst_origin[2], region[0], region[1], region[2],
num_events_in_wait_list, event_wait_list, event);
auto command_queue = icd_downcast(cq);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_image(dst_image) || !is_valid_buffer(src_buffer)) {
return CL_INVALID_MEM_OBJECT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
if (!is_same_context(command_queue, src_buffer) ||
!is_same_context(command_queue, dst_image) ||
!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_INVALID_MEM_OBJECT if dst_image is a 1D image buffer object
// created from src_buffer.
// TODO CL_INVALID_VALUE if the 1D, 2D, or 3D rectangular region specified
// by dst_origin and dst_origin + region refers to a region outside
// dst_origin, or if the region specified by src_offset and src_offset +
// src_cb refers to a region outside src_buffer.
// TODO CL_INVALID_VALUE if values in dst_origin and region do not follow
// rules described in the argument description for dst_origin and region.
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if src_buffer is a sub-buffer object
// and offset specified when the sub-buffer object is created is not aligned
// to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
// TODO CL_INVALID_IMAGE_SIZE if image dimensions (image width, height,
// specified or compute row and/or slice pitch) for dst_image are not
// supported by device associated with queue.
// TODO CL_INVALID_IMAGE_FORMAT if image format (image channel order and
// data type) for dst_image are not supported by device associated with
// queue.
// TODO CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate
// memory for data store associated with src_buffer or dst_image.
auto image = static_cast<cvk_image*>(dst_image);
auto buffer = static_cast<cvk_buffer*>(src_buffer);
if (!command_queue->device()->supports_images()) {
return CL_INVALID_OPERATION;
}
std::array<size_t, 3> origin = {dst_origin[0], dst_origin[1],
dst_origin[2]};
std::array<size_t, 3> reg = {region[0], region[1], region[2]};
if (image->is_backed_by_buffer_view()) {
auto cmd = new cvk_command_copy_buffer(
command_queue, CL_COMMAND_COPY_BUFFER_TO_IMAGE, buffer,
static_cast<cvk_buffer*>(image->buffer()), src_offset,
dst_origin[0] * image->element_size(),
region[0] * image->element_size());
return command_queue->enqueue_command_with_deps(
cmd, num_events_in_wait_list, event_wait_list, event);
} else {
auto cmd = std::make_unique<cvk_command_buffer_image_copy>(
CL_COMMAND_COPY_BUFFER_TO_IMAGE, command_queue, buffer, image,
src_offset, origin, reg);
return command_queue->enqueue_command_with_deps(
cmd.release(), num_events_in_wait_list, event_wait_list, event);
}
}
void* CLVK_API_CALL clEnqueueMapImage(
cl_command_queue cq, cl_mem image, cl_bool blocking_map,
cl_map_flags map_flags, const size_t* origin, const size_t* region,
size_t* image_row_pitch, size_t* image_slice_pitch,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
cl_event* event, cl_int* errcode_ret) {
TRACE_FUNCTION("command_queue", (uintptr_t)cq, "image", (uintptr_t)image,
"blocking_map", blocking_map, "map_flags", map_flags,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p, image = %p, blocking_map = %d, "
"map_flags = %lx, "
"origin = {%zu,%zu,%zu}, region = {%zu, %zu, %zu}, "
"image_row_pitch = %p, image_slice_pitch = %p, "
"num_events_in_wait_list = %u, event_wait_list = %p, "
"event = %p, errcode_ret = %p",
cq, image, blocking_map, map_flags, origin[0], origin[1],
origin[2], region[0], region[1], region[2], image_row_pitch,
image_slice_pitch, num_events_in_wait_list, event_wait_list,
event, errcode_ret);
auto command_queue = icd_downcast(cq);
cl_int err;
void* ret;
auto img = static_cast<cvk_image*>(image);
if (img->is_backed_by_buffer_view()) {
ret = cvk_enqueue_map_buffer(
command_queue, static_cast<cvk_buffer*>(img->buffer()),
blocking_map, origin[0] * img->element_size(),
region[0] * img->element_size(), map_flags, num_events_in_wait_list,
event_wait_list, event, &err, CL_COMMAND_MAP_IMAGE, img);
} else {
err = cvk_enqueue_map_image(
command_queue, image, true, blocking_map, map_flags, origin, region,
image_row_pitch, image_slice_pitch, nullptr, &ret,
num_events_in_wait_list, event_wait_list, event);
}
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
if (err != CL_SUCCESS) {
return nullptr;
}
return ret;
}
cl_program cvk_create_program_with_il(cl_context context, const void* il,
size_t length, cl_int* errcode_ret) {
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
if ((il == nullptr) || (length == 0)) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
// TODO CL_INVALID_VALUE if the length-byte block of memory pointed to by il
// does not contain well-formed intermediate language.
auto program = new cvk_program(icd_downcast(context), il, length);
*errcode_ret = program->parse_user_spec_constants();
return program;
}
cl_program CLVK_API_CALL clCreateProgramWithILKHR(cl_context context,
const void* il, size_t length,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, il = %p, length = %zu, errcode_ret = %p",
context, il, length, errcode_ret);
cl_int errcode;
auto program = cvk_create_program_with_il(context, il, length, &errcode);
if (errcode_ret != nullptr) {
*errcode_ret = errcode;
}
return program;
}
cl_program CLVK_API_CALL clCreateProgramWithIL(cl_context context,
const void* il, size_t length,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, il = %p, length = %zu, errcode_ret = %p",
context, il, length, errcode_ret);
cl_int errcode;
auto program = cvk_create_program_with_il(context, il, length, &errcode);
if (errcode_ret != nullptr) {
*errcode_ret = errcode;
}
return program;
}
cl_int CLVK_API_CALL
clSetProgramSpecializationConstant(cl_program prog, cl_uint spec_id,
size_t spec_size, const void* spec_value) {
TRACE_FUNCTION("program", (uintptr_t)prog, "spec_id", spec_id, "spec_size",
spec_size);
LOG_API_CALL("program = %p, spec_id = %u, spec_size = %zu, spec_value = %p",
prog, spec_id, spec_size, spec_value);
#if !ENABLE_SPIRV_IL
// If the SPIRV Intermediate Language is not enabled, the OpenCL
// specification requires clSetProgramSpecializationConstant to return
// CL_INVALID_OPERATION.
return CL_INVALID_OPERATION;
#endif
auto program = icd_downcast(prog);
if (!is_valid_program(program)) {
return CL_INVALID_PROGRAM;
}
// TODO CL_INVALID_OPERATION if no devices associated with program support
// intermediate language programs.
// TODO CL_COMPILER_NOT_AVAILABLE if program is created with
// clCreateProgramWithIL and a compiler is not available, i.e.
// CL_DEVICE_COMPILER_AVAILABLE specified in the Device Queries table is set
// to CL_FALSE.
if (spec_value == nullptr) {
return CL_INVALID_VALUE;
}
return program->set_user_spec_constant(spec_id, spec_size, spec_value);
}
// Shared Virtual Memory
void* CLVK_API_CALL clSVMAlloc(cl_context context, cl_svm_mem_flags flags,
size_t size, cl_uint alignment) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags, "size", size,
"alignment", alignment);
LOG_API_CALL("context = %p, flags = %lu, size = %zu, alignment = %u",
context, flags, size, alignment);
return nullptr;
}
void CLVK_API_CALL clSVMFree(cl_context context, void* svm_pointer) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, svm_pointer = %p", context, svm_pointer);
}
cl_int CLVK_API_CALL clEnqueueSVMFree(
cl_command_queue command_queue, cl_uint num_svm_pointers,
void* svm_pointers[],
void(CL_CALLBACK* pfn_free_func)(cl_command_queue queue,
cl_uint num_svm_pointers,
void* svm_pointers[], void* user_data),
void* user_data, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue,
"num_svm_pointers", num_svm_pointers,
"num_events_in_wait_list", num_events_in_wait_list);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(num_svm_pointers);
UNUSED(svm_pointers);
UNUSED(pfn_free_func);
UNUSED(user_data);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clEnqueueSVMMap(cl_command_queue command_queue,
cl_bool blocking_map, cl_map_flags flags,
void* svm_ptr, size_t size,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(blocking_map);
UNUSED(flags);
UNUSED(svm_ptr);
UNUSED(size);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clEnqueueSVMMemcpy(cl_command_queue command_queue,
cl_bool blocking_copy, void* dst_ptr,
const void* src_ptr, size_t size,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(blocking_copy);
UNUSED(dst_ptr);
UNUSED(src_ptr);
UNUSED(size);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clEnqueueSVMMemFill(cl_command_queue command_queue,
void* svm_ptr, const void* pattern,
size_t pattern_size, size_t size,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(svm_ptr);
UNUSED(pattern);
UNUSED(pattern_size);
UNUSED(size);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clEnqueueSVMMigrateMem(
cl_command_queue command_queue, cl_uint num_svm_pointers,
const void** svm_pointers, const size_t* sizes,
cl_mem_migration_flags flags, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(num_svm_pointers);
UNUSED(svm_pointers);
UNUSED(sizes);
UNUSED(flags);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clEnqueueSVMUnmap(cl_command_queue command_queue,
void* svm_ptr,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list,
cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p", command_queue);
UNUSED(svm_ptr);
UNUSED(num_events_in_wait_list);
UNUSED(event_wait_list);
UNUSED(event);
return CL_INVALID_OPERATION;
}
cl_int CLVK_API_CALL clSetKernelArgSVMPointer(cl_kernel kernel,
cl_uint arg_index,
const void* arg_value) {
TRACE_FUNCTION("kernel", (uintptr_t)kernel, "arg_index", arg_index);
LOG_API_CALL("kernel = %p, arg_index = %u, arg_value = %p", kernel,
arg_index, arg_value);
UNUSED(kernel);
UNUSED(arg_index);
UNUSED(arg_value);
return CL_INVALID_OPERATION;
}
// Pipes
cl_mem CLVK_API_CALL clCreatePipe(cl_context context, cl_mem_flags flags,
cl_uint pipe_packet_size,
cl_uint pipe_max_packets,
const cl_pipe_properties* properties,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context, "flags", flags,
"pipe_packet_size", pipe_packet_size, "pipe_max_packets",
pipe_max_packets);
LOG_API_CALL("context = %p, flags = %lx, packet_size = %u, max_packets = "
"%u, properties = %p, errcode_ret = %p",
context, flags, pipe_packet_size, pipe_max_packets, properties,
errcode_ret);
if (errcode_ret != nullptr) {
*errcode_ret = CL_INVALID_OPERATION;
}
return nullptr;
}
cl_int CLVK_API_CALL clGetPipeInfo(cl_mem pipe, cl_pipe_info param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("pipe", (uintptr_t)pipe, "param_name", param_name);
LOG_API_CALL("pipe = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
pipe, param_name, param_value_size, param_value,
param_value_size_ret);
return CL_INVALID_MEM_OBJECT;
}
// Timer functions
cl_int CLVK_API_CALL clGetHostTimer(cl_device_id device,
cl_ulong* host_timestamp) {
TRACE_FUNCTION("device", (uintptr_t)device);
LOG_API_CALL("device = %p, host_timestamp = %p", device, host_timestamp);
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
auto dev = icd_downcast(device);
if (!dev->has_timer_support()) {
return CL_INVALID_OPERATION;
}
if (host_timestamp == nullptr) {
return CL_INVALID_VALUE;
}
return dev->get_device_host_timer(nullptr, host_timestamp);
}
cl_int CLVK_API_CALL clGetDeviceAndHostTimer(cl_device_id device,
cl_ulong* device_timestamp,
cl_ulong* host_timestamp) {
TRACE_FUNCTION("device", (uintptr_t)device);
LOG_API_CALL("device = %p, device_timestamp = %p, host_timestamp = %p",
device, device_timestamp, host_timestamp);
if (!is_valid_device(device)) {
return CL_INVALID_DEVICE;
}
auto dev = icd_downcast(device);
if (!dev->has_timer_support()) {
return CL_INVALID_OPERATION;
}
if ((device_timestamp == nullptr) || (host_timestamp == nullptr)) {
return CL_INVALID_VALUE;
}
cl_ulong host;
cl_int err = dev->get_device_host_timer(nullptr, &host);
*device_timestamp = host;
*host_timestamp = host;
return err;
}
// cl_khr_semaphore
cl_semaphore_khr cvk_create_semaphore_with_properties_khr(
cl_context context, const cl_semaphore_properties_khr* sema_props,
cl_int* errcode_ret) {
if (!is_valid_context(context)) {
*errcode_ret = CL_INVALID_CONTEXT;
return nullptr;
}
if (sema_props == nullptr) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
cl_semaphore_type_khr type = 0;
std::vector<cl_semaphore_properties_khr> properties;
std::vector<cl_device_id> devices;
if (sema_props) {
bool has_type = false;
while (*sema_props) {
auto key = *sema_props;
auto value = *(sema_props + 1);
if (key == CL_SEMAPHORE_TYPE_KHR) {
properties.push_back(key);
properties.push_back(value);
type = value;
sema_props += 2;
has_type = true;
} else if (key == CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR) {
properties.push_back(key);
sema_props++;
while (*sema_props != CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR) {
auto devapi = reinterpret_cast<cl_device_id>(*sema_props);
if (!is_valid_device(devapi)) {
*errcode_ret = CL_INVALID_DEVICE;
return nullptr;
}
auto dev = static_cast<cvk_device*>(devapi);
if (!icd_downcast(context)->has_device(dev)) {
*errcode_ret = CL_INVALID_DEVICE;
return nullptr;
}
devices.push_back(devapi);
properties.push_back(*sema_props);
sema_props++;
}
properties.push_back(*sema_props);
sema_props++;
} else {
*errcode_ret = CL_INVALID_PROPERTY;
return nullptr;
}
}
if (!has_type) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
properties.push_back(0);
}
if (type == 0) {
*errcode_ret = CL_INVALID_VALUE;
return nullptr;
}
auto sem = std::make_unique<cvk_semaphore>(
icd_downcast(context), type, std::move(devices), std::move(properties));
auto err = sem->init();
if (err != CL_SUCCESS) {
*errcode_ret = err;
return nullptr;
}
*errcode_ret = CL_SUCCESS;
return sem.release();
}
cl_semaphore_khr clCreateSemaphoreWithPropertiesKHR(
cl_context context, const cl_semaphore_properties_khr* sema_props,
cl_int* errcode_ret) {
TRACE_FUNCTION("context", (uintptr_t)context);
LOG_API_CALL("context = %p, sema_props = %p, errcode_ret = %p", context,
sema_props, errcode_ret);
cl_int err;
auto sem =
cvk_create_semaphore_with_properties_khr(context, sema_props, &err);
if (errcode_ret != nullptr) {
*errcode_ret = err;
}
return sem;
}
cl_int
clEnqueueWaitSemaphoresKHR(cl_command_queue command_queue,
cl_uint num_sema_objects,
const cl_semaphore_khr* sema_objects,
const cl_semaphore_payload_khr* sema_payload_list,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p, num_sema_objects = %u, sema_objects = "
"%p, sema_payload_list = %p, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
command_queue, num_sema_objects, sema_objects,
sema_payload_list, num_events_in_wait_list, event_wait_list,
event);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (num_sema_objects == 0) {
return CL_INVALID_VALUE;
}
for (cl_uint i = 0; i < num_sema_objects; i++) {
if (!is_valid_semaphore(sema_objects[i])) {
return CL_INVALID_SEMAPHORE_KHR;
}
auto sem = icd_downcast(sema_objects[i]);
if (sem->requires_payload() && sema_payload_list == nullptr) {
return CL_INVALID_VALUE;
}
auto queue = icd_downcast(command_queue);
if (!sem->can_be_used_with_device(queue->device())) {
return CL_INVALID_COMMAND_QUEUE;
}
}
if (!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (!is_same_context(command_queue, num_sema_objects, sema_objects)) {
return CL_INVALID_CONTEXT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
for (cl_uint i = 0; i < num_events_in_wait_list; i++) {
if (icd_downcast(event_wait_list[i])->terminated()) {
return CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
}
}
return CL_INVALID_OPERATION;
}
cl_int
clEnqueueSignalSemaphoresKHR(cl_command_queue command_queue,
cl_uint num_sema_objects,
const cl_semaphore_khr* sema_objects,
const cl_semaphore_payload_khr* sema_payload_list,
cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue);
LOG_API_CALL("command_queue = %p, num_sema_objects = %u, sema_objects = "
"%p, sema_payload_list = %p, num_events_in_wait_list = %u, "
"event_wait_list = %p, event = %p",
command_queue, num_sema_objects, sema_objects,
sema_payload_list, num_events_in_wait_list, event_wait_list,
event);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (num_sema_objects == 0) {
return CL_INVALID_VALUE;
}
for (cl_uint i = 0; i < num_sema_objects; i++) {
if (!is_valid_semaphore(sema_objects[i])) {
return CL_INVALID_SEMAPHORE_KHR;
}
auto sem = icd_downcast(sema_objects[i]);
if (sem->requires_payload() && sema_payload_list == nullptr) {
return CL_INVALID_VALUE;
}
auto queue = icd_downcast(command_queue);
if (!sem->can_be_used_with_device(queue->device())) {
return CL_INVALID_COMMAND_QUEUE;
}
}
if (!is_same_context(command_queue, num_events_in_wait_list,
event_wait_list)) {
return CL_INVALID_CONTEXT;
}
if (!is_same_context(command_queue, num_sema_objects, sema_objects)) {
return CL_INVALID_CONTEXT;
}
if (!is_valid_event_wait_list(num_events_in_wait_list, event_wait_list)) {
return CL_INVALID_EVENT_WAIT_LIST;
}
for (cl_uint i = 0; i < num_events_in_wait_list; i++) {
if (icd_downcast(event_wait_list[i])->terminated()) {
return CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
}
}
return CL_INVALID_OPERATION;
}
cl_int clGetSemaphoreInfoKHR(const cl_semaphore_khr sema_object,
cl_semaphore_info_khr param_name,
size_t param_value_size, void* param_value,
size_t* param_value_size_ret) {
TRACE_FUNCTION("sema_object", (uintptr_t)sema_object);
LOG_API_CALL("sema_object = %p, param_name = %x, param_value_size = %zu, "
"param_value = %p, param_value_size_ret = %p",
sema_object, param_name, param_value_size, param_value,
param_value_size_ret);
cl_int ret = CL_SUCCESS;
size_t ret_size = 0;
const void* copy_ptr = nullptr;
cl_context val_context;
cl_uint val_uint;
cl_semaphore_type_khr val_semaphore_type;
cl_semaphore_payload_khr val_semaphore_payload;
auto sem = icd_downcast(sema_object);
if (!is_valid_semaphore(sem)) {
return CL_INVALID_SEMAPHORE_KHR;
}
switch (param_name) {
case CL_SEMAPHORE_CONTEXT_KHR:
val_context = sem->context();
copy_ptr = &val_context;
ret_size = sizeof(val_context);
break;
case CL_SEMAPHORE_REFERENCE_COUNT_KHR:
val_uint = sem->refcount();
copy_ptr = &val_uint;
ret_size = sizeof(val_uint);
break;
case CL_SEMAPHORE_TYPE_KHR:
val_semaphore_type = sem->type();
copy_ptr = &val_semaphore_type;
ret_size = sizeof(val_semaphore_type);
break;
case CL_SEMAPHORE_PAYLOAD_KHR:
val_semaphore_payload = sem->payload();
copy_ptr = &val_semaphore_payload;
ret_size = sizeof(val_semaphore_payload);
break;
case CL_SEMAPHORE_PROPERTIES_KHR:
copy_ptr = sem->properties().data();
ret_size =
sem->properties().size() * sizeof(cl_semaphore_properties_khr);
break;
case CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR:
copy_ptr = sem->devices().data();
ret_size = sem->devices().size() * sizeof(cl_device_id);
break;
default:
ret = CL_INVALID_VALUE;
}
if ((param_value != nullptr) && (copy_ptr != nullptr)) {
if (param_value_size < ret_size) {
ret = CL_INVALID_VALUE;
}
memcpy(param_value, copy_ptr, std::min(param_value_size, ret_size));
}
if (param_value_size_ret != nullptr) {
*param_value_size_ret = ret_size;
}
return ret;
}
cl_int clReleaseSemaphoreKHR(cl_semaphore_khr sema_object) {
TRACE_FUNCTION("sema_object", (uintptr_t)sema_object);
LOG_API_CALL("sema_object = %p", sema_object);
if (!is_valid_semaphore(sema_object)) {
return CL_INVALID_SEMAPHORE_KHR;
}
icd_downcast(sema_object)->release();
return CL_SUCCESS;
}
cl_int clRetainSemaphoreKHR(cl_semaphore_khr sema_object) {
TRACE_FUNCTION("sema_object", (uintptr_t)sema_object);
LOG_API_CALL("sema_object = %p", sema_object);
if (!is_valid_semaphore(sema_object)) {
return CL_INVALID_SEMAPHORE_KHR;
}
icd_downcast(sema_object)->retain();
return CL_SUCCESS;
}
// clang-format off
cl_icd_dispatch gDispatchTable = {
// OpenCL 1.0
{clGetPlatformIDs},
clGetPlatformInfo,
clGetDeviceIDs,
clGetDeviceInfo,
clCreateContext,
clCreateContextFromType,
clRetainContext,
clReleaseContext,
clGetContextInfo,
clCreateCommandQueue,
clRetainCommandQueue,
clReleaseCommandQueue,
clGetCommandQueueInfo,
clSetCommandQueueProperty,
clCreateBuffer,
clCreateImage2D,
clCreateImage3D,
clRetainMemObject,
clReleaseMemObject,
clGetSupportedImageFormats,
clGetMemObjectInfo,
clGetImageInfo,
clCreateSampler,
clRetainSampler,
clReleaseSampler,
clGetSamplerInfo,
clCreateProgramWithSource,
clCreateProgramWithBinary,
clRetainProgram,
clReleaseProgram,
clBuildProgram,
{clUnloadCompiler},
clGetProgramInfo,
clGetProgramBuildInfo,
clCreateKernel,
clCreateKernelsInProgram,
clRetainKernel,
clReleaseKernel,
clSetKernelArg,
clGetKernelInfo,
clGetKernelWorkGroupInfo,
clWaitForEvents,
clGetEventInfo,
clRetainEvent,
clReleaseEvent,
clGetEventProfilingInfo,
clFlush,
clFinish,
clEnqueueReadBuffer,
clEnqueueWriteBuffer,
clEnqueueCopyBuffer,
clEnqueueReadImage,
clEnqueueWriteImage,
clEnqueueCopyImage,
clEnqueueCopyImageToBuffer,
clEnqueueCopyBufferToImage,
clEnqueueMapBuffer,
clEnqueueMapImage,
clEnqueueUnmapMemObject,
clEnqueueNDRangeKernel,
clEnqueueTask,
clEnqueueNativeKernel,
clEnqueueMarker,
clEnqueueWaitForEvents,
clEnqueueBarrier,
clGetExtensionFunctionAddress,
nullptr, // clCreateFromGLBuffer;
nullptr, // clCreateFromGLTexture2D;
nullptr, // clCreateFromGLTexture3D;
nullptr, // clCreateFromGLRenderbuffer;
nullptr, // clGetGLObjectInfo;
nullptr, // clGetGLTextureInfo;
nullptr, // clEnqueueAcquireGLObjects;
nullptr, // clEnqueueReleaseGLObjects;
nullptr, // clGetGLContextInfoKHR;
nullptr, // clGetDeviceIDsFromD3D10KHR;
nullptr, // clCreateFromD3D10BufferKHR;
nullptr, // clCreateFromD3D10Texture2DKHR;
nullptr, // clCreateFromD3D10Texture3DKHR;
nullptr, // clEnqueueAcquireD3D10ObjectsKHR;
nullptr, // clEnqueueReleaseD3D10ObjectsKHR;
// OpenCL 1.1
clSetEventCallback,
clCreateSubBuffer,
clSetMemObjectDestructorCallback,
clCreateUserEvent,
clSetUserEventStatus,
clEnqueueReadBufferRect,
clEnqueueWriteBufferRect,
clEnqueueCopyBufferRect,
/* cl_ext_device_fission */
nullptr, // clCreateSubDevicesEXT;
nullptr, // clRetainDeviceEXT;
nullptr, // clReleaseDeviceEXT;
/* cl_khr_gl_event */
nullptr, // clCreateEventFromGLsyncKHR;
// OpenCL 1.2
clCreateSubDevices,
clRetainDevice,
clReleaseDevice,
clCreateImage,
clCreateProgramWithBuiltInKernels,
clCompileProgram,
clLinkProgram,
clUnloadPlatformCompiler,
clGetKernelArgInfo,
clEnqueueFillBuffer,
clEnqueueFillImage,
clEnqueueMigrateMemObjects,
clEnqueueMarkerWithWaitList,
clEnqueueBarrierWithWaitList,
clGetExtensionFunctionAddressForPlatform,
nullptr, // clCreateFromGLTexture;
/* cl_khr_d3d11_sharing */
nullptr, // clGetDeviceIDsFromD3D11KHR;
nullptr, // clCreateFromD3D11BufferKHR;
nullptr, // clCreateFromD3D11Texture2DKHR;
nullptr, // clCreateFromD3D11Texture3DKHR;
nullptr, // clCreateFromDX9MediaSurfaceKHR;
nullptr, // clEnqueueAcquireD3D11ObjectsKHR;
nullptr, // clEnqueueReleaseD3D11ObjectsKHR;
/* cl_khr_dx9_media_sharing */
nullptr, // clGetDeviceIDsFromDX9MediaAdapterKHR;
nullptr, // clEnqueueAcquireDX9MediaSurfacesKHR;
nullptr, // clEnqueueReleaseDX9MediaSurfacesKHR;
/* cl_khr_egl_image */
nullptr, // clCreateFromEGLImageKHR;
nullptr, // clEnqueueAcquireEGLObjectsKHR;
nullptr, // clEnqueueReleaseEGLObjectsKHR;
/* cl_khr_egl_event */
nullptr, // clCreateEventFromEGLSyncKHR;
/* OpenCL 2.0 */
clCreateCommandQueueWithProperties,
clCreatePipe,
clGetPipeInfo,
clSVMAlloc,
clSVMFree,
clEnqueueSVMFree,
clEnqueueSVMMemcpy,
clEnqueueSVMMemFill,
clEnqueueSVMMap,
clEnqueueSVMUnmap,
clCreateSamplerWithProperties,
clSetKernelArgSVMPointer,
clSetKernelExecInfo,
/* cl_khr_sub_groups */
clGetKernelSubGroupInfo,
/* OpenCL 2.1 */
clCloneKernel,
clCreateProgramWithIL,
clEnqueueSVMMigrateMem,
clGetDeviceAndHostTimer,
clGetHostTimer,
clGetKernelSubGroupInfo,
clSetDefaultDeviceCommandQueue,
/* OpenCL 2.2 */
clSetProgramReleaseCallback,
clSetProgramSpecializationConstant,
/* OpenCL 3.0 */
clCreateBufferWithProperties,
clCreateImageWithProperties,
clSetContextDestructorCallback,
};
// clang-format on
cl_int CLVK_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries,
cl_platform_id* platforms,
cl_uint* num_platforms) {
auto state = get_or_init_global_state();
TRACE_FUNCTION("num_entries", num_entries);
LOG_API_CALL("num_entries = %u, platforms = %p, num_platforms = %p",
num_entries, platforms, num_platforms);
return cvk_get_platform_ids(state, num_entries, platforms, num_platforms);
}
cl_int CLVK_API_CALL clGetKernelSuggestedLocalWorkSizeKHR(
cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim,
const size_t* global_work_offset, const size_t* global_work_size,
size_t* suggested_local_work_size) {
TRACE_FUNCTION("command_queue", (uintptr_t)command_queue, "kernel",
(uintptr_t)kernel);
LOG_API_CALL(
"command_queue = %p, kernel = %p, work_dim = %u, global_work_offset = "
"%p, global_work_size = %p, suggested_local_work_size = %p",
command_queue, kernel, work_dim, global_work_offset, global_work_size,
suggested_local_work_size);
if (!is_valid_command_queue(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid_kernel(kernel)) {
return CL_INVALID_KERNEL;
}
if (!is_same_context(command_queue, kernel)) {
return CL_INVALID_CONTEXT;
}
// TODO CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built
// program executable available for kernel for the device associated with
// command_queue.
// TODO CL_INVALID_KERNEL_ARGS if all argument values for kernel have not
// been set.
// TODO CL_MISALIGNED_SUB_BUFFER_OFFSET if a sub-buffer object is set as an
// argument to kernel and the offset specified when the sub-buffer object
// was created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN for the
// device associated with command_queue.
// TODO CL_INVALID_IMAGE_SIZE if an image object is set as an argument to
// kernel and the image dimensions are not supported by device associated
// with command_queue.
// TODO CL_IMAGE_FORMAT_NOT_SUPPORTED if an image object is set as an
// argument to kernel and the image format is not supported by the device
// associated with command_queue.
// TODO CL_INVALID_OPERATION if an SVM pointer is set as an argument to
// kernel and the device associated with command_queue does not support SVM
// or the required SVM capabilities for the SVM pointer.
if ((work_dim < 1) ||
(work_dim >
icd_downcast(command_queue)->device()->max_work_item_dimensions())) {
return CL_INVALID_WORK_DIMENSION;
}
if (global_work_size == nullptr) {
return CL_INVALID_GLOBAL_WORK_SIZE;
}
for (cl_uint i = 0; i < work_dim; i++) {
if (global_work_size[i] == 0) {
return CL_INVALID_GLOBAL_WORK_SIZE;
}
}
// TODO CL_INVALID_GLOBAL_WORK_SIZE if any of the values specified in
// global_work_size exceed the maximum value representable by size_t on the
// device associated with command_queue.
// TODO CL_INVALID_GLOBAL_OFFSET if the value specified in global_work_size
// plus the corresponding value in global_work_offset for dimension exceeds
// the maximum value representable by size_t on the device associated with
// command_queue.
cvk_ndrange ndrange(work_dim, global_work_offset, global_work_size,
nullptr);
icd_downcast(command_queue)
->device()
->select_work_group_size(icd_downcast(kernel), ndrange.gws,
ndrange.lws);
for (cl_uint i = 0; i < work_dim; i++) {
suggested_local_work_size[i] = ndrange.lws[i];
}
return CL_SUCCESS;
}