diff options
Diffstat (limited to 'src/gallium/frontends/clover/api')
| -rw-r--r-- | src/gallium/frontends/clover/api/context.cpp | 144 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/device.cpp | 421 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/dispatch.cpp | 174 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/dispatch.hpp | 105 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/event.cpp | 309 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/interop.cpp | 69 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/kernel.cpp | 390 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/memory.cpp | 497 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/platform.cpp | 235 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/program.cpp | 479 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/queue.cpp | 135 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/sampler.cpp | 100 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/transfer.cpp | 1059 | ||||
| -rw-r--r-- | src/gallium/frontends/clover/api/util.hpp | 84 |
14 files changed, 4201 insertions, 0 deletions
diff --git a/src/gallium/frontends/clover/api/context.cpp b/src/gallium/frontends/clover/api/context.cpp new file mode 100644 index 00000000000..c0cd2d32b95 --- /dev/null +++ b/src/gallium/frontends/clover/api/context.cpp @@ -0,0 +1,144 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/context.hpp" +#include "core/platform.hpp" + +using namespace clover; + +CLOVER_API cl_context +clCreateContext(const cl_context_properties *d_props, cl_uint num_devs, + const cl_device_id *d_devs, + void (CL_CALLBACK *pfn_notify)(const char *, const void *, + size_t, void *), + void *user_data, cl_int *r_errcode) try { + auto props = obj<property_list_tag>(d_props); + auto devs = objs(d_devs, num_devs); + + if (!pfn_notify && user_data) + throw error(CL_INVALID_VALUE); + + for (auto &prop : props) { + if (prop.first == CL_CONTEXT_PLATFORM) + obj(prop.second.as<cl_platform_id>()); + else + throw error(CL_INVALID_PROPERTY); + } + + const auto notify = (!pfn_notify ? context::notify_action() : + [=](const char *s) { + pfn_notify(s, NULL, 0, user_data); + }); + + ret_error(r_errcode, CL_SUCCESS); + return desc(new context(props, devs, notify)); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_context +clCreateContextFromType(const cl_context_properties *d_props, + cl_device_type type, + void (CL_CALLBACK *pfn_notify)( + const char *, const void *, size_t, void *), + void *user_data, cl_int *r_errcode) try { + cl_platform_id d_platform; + cl_uint num_platforms; + cl_int ret; + std::vector<cl_device_id> devs; + cl_uint num_devices; + + ret = clGetPlatformIDs(1, &d_platform, &num_platforms); + if (ret || !num_platforms) + throw error(CL_INVALID_PLATFORM); + + ret = clGetDeviceIDs(d_platform, type, 0, NULL, &num_devices); + if (ret) + throw error(CL_DEVICE_NOT_FOUND); + devs.resize(num_devices); + ret = clGetDeviceIDs(d_platform, type, num_devices, devs.data(), 0); + if (ret) + throw error(CL_DEVICE_NOT_FOUND); + + return clCreateContext(d_props, num_devices, devs.data(), pfn_notify, + user_data, r_errcode); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clRetainContext(cl_context d_ctx) try { + obj(d_ctx).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseContext(cl_context d_ctx) try { + if (obj(d_ctx).release()) + delete pobj(d_ctx); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetContextInfo(cl_context d_ctx, cl_context_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &ctx = obj(d_ctx); + + switch (param) { + case CL_CONTEXT_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = ctx.ref_count(); + break; + + case CL_CONTEXT_NUM_DEVICES: + buf.as_scalar<cl_uint>() = ctx.devices().size(); + break; + + case CL_CONTEXT_DEVICES: + buf.as_vector<cl_device_id>() = descs(ctx.devices()); + break; + + case CL_CONTEXT_PROPERTIES: + buf.as_vector<cl_context_properties>() = desc(ctx.properties()); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/device.cpp b/src/gallium/frontends/clover/api/device.cpp new file mode 100644 index 00000000000..042f2eda21c --- /dev/null +++ b/src/gallium/frontends/clover/api/device.cpp @@ -0,0 +1,421 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/platform.hpp" +#include "core/device.hpp" +#include "git_sha1.h" + +using namespace clover; + +CLOVER_API cl_int +clGetDeviceIDs(cl_platform_id d_platform, cl_device_type device_type, + cl_uint num_entries, cl_device_id *rd_devices, + cl_uint *rnum_devices) try { + auto &platform = obj(d_platform); + std::vector<cl_device_id> d_devs; + + if ((!num_entries && rd_devices) || + (!rnum_devices && !rd_devices)) + throw error(CL_INVALID_VALUE); + + // Collect matching devices + for (device &dev : platform) { + if (((device_type & CL_DEVICE_TYPE_DEFAULT) && + dev == platform.front()) || + (device_type & dev.type())) + d_devs.push_back(desc(dev)); + } + + if (d_devs.empty()) + throw error(CL_DEVICE_NOT_FOUND); + + // ...and return the requested data. + if (rnum_devices) + *rnum_devices = d_devs.size(); + if (rd_devices) + copy(range(d_devs.begin(), + std::min((unsigned)d_devs.size(), num_entries)), + rd_devices); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clCreateSubDevices(cl_device_id d_dev, + const cl_device_partition_property *props, + cl_uint num_devs, cl_device_id *rd_devs, + cl_uint *rnum_devs) { + // There are no currently supported partitioning schemes. + return CL_INVALID_VALUE; +} + +CLOVER_API cl_int +clRetainDevice(cl_device_id d_dev) try { + obj(d_dev); + + // The reference count doesn't change for root devices. + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseDevice(cl_device_id d_dev) try { + obj(d_dev); + + // The reference count doesn't change for root devices. + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetDeviceInfo(cl_device_id d_dev, cl_device_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &dev = obj(d_dev); + + switch (param) { + case CL_DEVICE_TYPE: + buf.as_scalar<cl_device_type>() = dev.type(); + break; + + case CL_DEVICE_VENDOR_ID: + buf.as_scalar<cl_uint>() = dev.vendor_id(); + break; + + case CL_DEVICE_MAX_COMPUTE_UNITS: + buf.as_scalar<cl_uint>() = dev.max_compute_units(); + break; + + case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: + buf.as_scalar<cl_uint>() = dev.max_block_size().size(); + break; + + case CL_DEVICE_MAX_WORK_ITEM_SIZES: + buf.as_vector<size_t>() = dev.max_block_size(); + break; + + case CL_DEVICE_MAX_WORK_GROUP_SIZE: + buf.as_scalar<size_t>() = dev.max_threads_per_block(); + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR: + buf.as_scalar<cl_uint>() = 16; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT: + buf.as_scalar<cl_uint>() = 8; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT: + buf.as_scalar<cl_uint>() = 4; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG: + buf.as_scalar<cl_uint>() = 2; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT: + buf.as_scalar<cl_uint>() = 4; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE: + buf.as_scalar<cl_uint>() = dev.has_doubles() ? 2 : 0; + break; + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF: + buf.as_scalar<cl_uint>() = dev.has_halves() ? 8 : 0; + break; + + case CL_DEVICE_MAX_CLOCK_FREQUENCY: + buf.as_scalar<cl_uint>() = dev.max_clock_frequency(); + break; + + case CL_DEVICE_ADDRESS_BITS: + buf.as_scalar<cl_uint>() = dev.address_bits(); + break; + + case CL_DEVICE_MAX_READ_IMAGE_ARGS: + buf.as_scalar<cl_uint>() = dev.max_images_read(); + break; + + case CL_DEVICE_MAX_WRITE_IMAGE_ARGS: + buf.as_scalar<cl_uint>() = dev.max_images_write(); + break; + + case CL_DEVICE_MAX_MEM_ALLOC_SIZE: + buf.as_scalar<cl_ulong>() = dev.max_mem_alloc_size(); + break; + + case CL_DEVICE_IMAGE2D_MAX_WIDTH: + case CL_DEVICE_IMAGE2D_MAX_HEIGHT: + buf.as_scalar<size_t>() = 1 << dev.max_image_levels_2d(); + break; + + case CL_DEVICE_IMAGE3D_MAX_WIDTH: + case CL_DEVICE_IMAGE3D_MAX_HEIGHT: + case CL_DEVICE_IMAGE3D_MAX_DEPTH: + buf.as_scalar<size_t>() = 1 << dev.max_image_levels_3d(); + break; + + case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE: + buf.as_scalar<size_t>() = dev.max_image_buffer_size(); + break; + + case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE: + buf.as_scalar<size_t>() = dev.max_image_array_number(); + break; + + case CL_DEVICE_IMAGE_SUPPORT: + buf.as_scalar<cl_bool>() = dev.image_support(); + break; + + case CL_DEVICE_MAX_PARAMETER_SIZE: + buf.as_scalar<size_t>() = dev.max_mem_input(); + break; + + case CL_DEVICE_MAX_SAMPLERS: + buf.as_scalar<cl_uint>() = dev.max_samplers(); + break; + + case CL_DEVICE_MEM_BASE_ADDR_ALIGN: + buf.as_scalar<cl_uint>() = 8 * + std::max(dev.mem_base_addr_align(), (cl_uint) sizeof(cl_long) * 16); + break; + + case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE: + buf.as_scalar<cl_uint>() = 128; + break; + + case CL_DEVICE_HALF_FP_CONFIG: + // This is the "mandated minimum half precision floating-point + // capability" for OpenCL 1.x. + buf.as_scalar<cl_device_fp_config>() = + CL_FP_INF_NAN | CL_FP_ROUND_TO_NEAREST; + break; + + case CL_DEVICE_SINGLE_FP_CONFIG: + // This is the "mandated minimum single precision floating-point + // capability" for OpenCL 1.1. In OpenCL 1.2, nothing is required for + // custom devices. + buf.as_scalar<cl_device_fp_config>() = + CL_FP_INF_NAN | CL_FP_ROUND_TO_NEAREST; + break; + + case CL_DEVICE_DOUBLE_FP_CONFIG: + if (dev.has_doubles()) + // This is the "mandated minimum double precision floating-point + // capability" + buf.as_scalar<cl_device_fp_config>() = + CL_FP_FMA + | CL_FP_ROUND_TO_NEAREST + | CL_FP_ROUND_TO_ZERO + | CL_FP_ROUND_TO_INF + | CL_FP_INF_NAN + | CL_FP_DENORM; + else + buf.as_scalar<cl_device_fp_config>() = 0; + break; + + case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE: + buf.as_scalar<cl_device_mem_cache_type>() = CL_NONE; + break; + + case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE: + buf.as_scalar<cl_uint>() = 0; + break; + + case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE: + buf.as_scalar<cl_ulong>() = 0; + break; + + case CL_DEVICE_GLOBAL_MEM_SIZE: + buf.as_scalar<cl_ulong>() = dev.max_mem_global(); + break; + + case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: + buf.as_scalar<cl_ulong>() = dev.max_const_buffer_size(); + break; + + case CL_DEVICE_MAX_CONSTANT_ARGS: + buf.as_scalar<cl_uint>() = dev.max_const_buffers(); + break; + + case CL_DEVICE_LOCAL_MEM_TYPE: + buf.as_scalar<cl_device_local_mem_type>() = CL_LOCAL; + break; + + case CL_DEVICE_LOCAL_MEM_SIZE: + buf.as_scalar<cl_ulong>() = dev.max_mem_local(); + break; + + case CL_DEVICE_ERROR_CORRECTION_SUPPORT: + buf.as_scalar<cl_bool>() = CL_FALSE; + break; + + case CL_DEVICE_PROFILING_TIMER_RESOLUTION: + buf.as_scalar<size_t>() = 0; + break; + + case CL_DEVICE_ENDIAN_LITTLE: + buf.as_scalar<cl_bool>() = (dev.endianness() == PIPE_ENDIAN_LITTLE); + break; + + case CL_DEVICE_AVAILABLE: + case CL_DEVICE_COMPILER_AVAILABLE: + case CL_DEVICE_LINKER_AVAILABLE: + buf.as_scalar<cl_bool>() = CL_TRUE; + break; + + case CL_DEVICE_EXECUTION_CAPABILITIES: + buf.as_scalar<cl_device_exec_capabilities>() = CL_EXEC_KERNEL; + break; + + case CL_DEVICE_QUEUE_PROPERTIES: + buf.as_scalar<cl_command_queue_properties>() = CL_QUEUE_PROFILING_ENABLE; + break; + + case CL_DEVICE_BUILT_IN_KERNELS: + buf.as_string() = ""; + break; + + case CL_DEVICE_NAME: + buf.as_string() = dev.device_name(); + break; + + case CL_DEVICE_VENDOR: + buf.as_string() = dev.vendor_name(); + break; + + case CL_DRIVER_VERSION: + buf.as_string() = PACKAGE_VERSION; + break; + + case CL_DEVICE_PROFILE: + buf.as_string() = "FULL_PROFILE"; + break; + + case CL_DEVICE_VERSION: + buf.as_string() = "OpenCL " + dev.device_version() + " Mesa " PACKAGE_VERSION MESA_GIT_SHA1; + break; + + case CL_DEVICE_EXTENSIONS: + buf.as_string() = dev.supported_extensions(); + break; + + case CL_DEVICE_PLATFORM: + buf.as_scalar<cl_platform_id>() = desc(dev.platform); + break; + + case CL_DEVICE_HOST_UNIFIED_MEMORY: + buf.as_scalar<cl_bool>() = dev.has_unified_memory(); + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR: + buf.as_scalar<cl_uint>() = 16; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT: + buf.as_scalar<cl_uint>() = 8; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT: + buf.as_scalar<cl_uint>() = 4; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG: + buf.as_scalar<cl_uint>() = 2; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT: + buf.as_scalar<cl_uint>() = 4; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE: + buf.as_scalar<cl_uint>() = dev.has_doubles() ? 2 : 0; + break; + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF: + buf.as_scalar<cl_uint>() = dev.has_halves() ? 8 : 0; + break; + + case CL_DEVICE_OPENCL_C_VERSION: + buf.as_string() = "OpenCL C " + dev.device_clc_version() + " "; + break; + + case CL_DEVICE_PRINTF_BUFFER_SIZE: + // Per the spec, the minimum value for the FULL profile is 1 MB. + // However, clover is not ready yet to support it + buf.as_scalar<size_t>() = 0 /* 1024 */; + break; + + case CL_DEVICE_PREFERRED_INTEROP_USER_SYNC: + buf.as_scalar<cl_bool>() = CL_TRUE; + break; + + case CL_DEVICE_PARENT_DEVICE: + buf.as_scalar<cl_device_id>() = NULL; + break; + + case CL_DEVICE_PARTITION_MAX_SUB_DEVICES: + buf.as_scalar<cl_uint>() = 0; + break; + + case CL_DEVICE_PARTITION_PROPERTIES: + buf.as_vector<cl_device_partition_property>() = + desc(property_list<cl_device_partition_property>()); + break; + + case CL_DEVICE_PARTITION_AFFINITY_DOMAIN: + buf.as_scalar<cl_device_affinity_domain>() = 0; + break; + + case CL_DEVICE_PARTITION_TYPE: + buf.as_vector<cl_device_partition_property>() = + desc(property_list<cl_device_partition_property>()); + break; + + case CL_DEVICE_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = 1; + break; + + case CL_DEVICE_SVM_CAPABILITIES: + case CL_DEVICE_SVM_CAPABILITIES_ARM: + buf.as_scalar<cl_device_svm_capabilities>() = dev.svm_support(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/dispatch.cpp b/src/gallium/frontends/clover/api/dispatch.cpp new file mode 100644 index 00000000000..6e1b0351afa --- /dev/null +++ b/src/gallium/frontends/clover/api/dispatch.cpp @@ -0,0 +1,174 @@ +// +// Copyright 2013 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/dispatch.hpp" + +namespace clover { + const cl_icd_dispatch _dispatch = { + clGetPlatformIDs, + GetPlatformInfo, + clGetDeviceIDs, + clGetDeviceInfo, + clCreateContext, + clCreateContextFromType, + clRetainContext, + clReleaseContext, + clGetContextInfo, + clCreateCommandQueue, + clRetainCommandQueue, + clReleaseCommandQueue, + clGetCommandQueueInfo, + NULL, // 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, + GetExtensionFunctionAddress, + NULL, // clCreateFromGLBuffer + NULL, // clCreateFromGLTexture2D + NULL, // clCreateFromGLTexture3D + NULL, // clCreateFromGLRenderbuffer + NULL, // clGetGLObjectInfo + NULL, // clGetGLTextureInfo + NULL, // clEnqueueAcquireGLObjects + NULL, // clEnqueueReleaseGLObjects + NULL, // clGetGLContextInfoKHR + NULL, // clGetDeviceIDsFromD3D10KHR + NULL, // clCreateFromD3D10BufferKHR + NULL, // clCreateFromD3D10Texture2DKHR + NULL, // clCreateFromD3D10Texture3DKHR + NULL, // clEnqueueAcquireD3D10ObjectsKHR + NULL, // clEnqueueReleaseD3D10ObjectsKHR + clSetEventCallback, + clCreateSubBuffer, + clSetMemObjectDestructorCallback, + clCreateUserEvent, + clSetUserEventStatus, + clEnqueueReadBufferRect, + clEnqueueWriteBufferRect, + clEnqueueCopyBufferRect, + NULL, // clCreateSubDevicesEXT + NULL, // clRetainDeviceEXT + NULL, // clReleaseDeviceEXT + NULL, // clCreateEventFromGLsyncKHR + clCreateSubDevices, + clRetainDevice, + clReleaseDevice, + clCreateImage, + clCreateProgramWithBuiltInKernels, + clCompileProgram, + clLinkProgram, + clUnloadPlatformCompiler, + clGetKernelArgInfo, + clEnqueueFillBuffer, + clEnqueueFillImage, + clEnqueueMigrateMemObjects, + clEnqueueMarkerWithWaitList, + clEnqueueBarrierWithWaitList, + GetExtensionFunctionAddressForPlatform, + NULL, // clCreateFromGLTexture + NULL, // clGetDeviceIDsFromD3D11KHR + NULL, // clCreateFromD3D11BufferKHR + NULL, // clCreateFromD3D11Texture2DKHR + NULL, // clCreateFromD3D11Texture3DKHR + NULL, // clCreateFromDX9MediaSurfaceKHR + NULL, // clEnqueueAcquireD3D11ObjectsKHR + NULL, // clEnqueueReleaseD3D11ObjectsKHR + NULL, // clGetDeviceIDsFromDX9MediaAdapterKHR + NULL, // clEnqueueAcquireDX9MediaSurfacesKHR + NULL, // clEnqueueReleaseDX9MediaSurfacesKHR + NULL, // clCreateFromEGLImageKHR + NULL, // clEnqueueAcquireEGLObjectsKHR + NULL, // clEnqueueReleaseEGLObjectsKHR + NULL, // clCreateEventFromEGLSyncKHR + clCreateCommandQueueWithProperties, + NULL, // clCreatePipe + NULL, // clGetPipeInfo + clSVMAlloc, + clSVMFree, + clEnqueueSVMFree, + clEnqueueSVMMemcpy, + clEnqueueSVMMemFill, + clEnqueueSVMMap, + clEnqueueSVMUnmap, + NULL, // clCreateSamplerWithProperties + clSetKernelArgSVMPointer, + clSetKernelExecInfo, + NULL, // clGetKernelSubGroupInfoKHR + NULL, // clCloneKernel + NULL, // clCreateProgramWithIL + clEnqueueSVMMigrateMem, + NULL, // clGetDeviceAndHostTimer + NULL, // clGetHostTimer + NULL, // clGetKernelSubGroupInfo + NULL, // clSetDefaultDeviceCommandQueue + NULL, // clSetProgramReleaseCallback + NULL, // clSetProgramSpecializationConstant + }; +} diff --git a/src/gallium/frontends/clover/api/dispatch.hpp b/src/gallium/frontends/clover/api/dispatch.hpp new file mode 100644 index 00000000000..ea835ed6da4 --- /dev/null +++ b/src/gallium/frontends/clover/api/dispatch.hpp @@ -0,0 +1,105 @@ +// +// Copyright 2013 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#ifndef API_DISPATCH_HPP +#define API_DISPATCH_HPP + +#include "CL/cl.h" +#include "CL/cl_ext.h" +#include "CL/cl_egl.h" +#include "CL/cl_gl.h" +#include "CL/cl_icd.h" + +namespace clover { + extern const cl_icd_dispatch _dispatch; + + cl_int + GetPlatformInfo(cl_platform_id d_platform, cl_platform_info param, + size_t size, void *r_buf, size_t *r_size); + + void * + GetExtensionFunctionAddress(const char *p_name); + + void * + GetExtensionFunctionAddressForPlatform(cl_platform_id d_platform, + const char *p_name); + + cl_int + IcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms, + cl_uint *rnum_platforms); + + cl_int + EnqueueSVMFree(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, + cl_int cmd); + + cl_int + EnqueueSVMMemcpy(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, + cl_int cmd); + + cl_int + EnqueueSVMMap(cl_command_queue command_queue, + cl_bool blocking_map, + cl_map_flags map_flags, + void *svm_ptr, + size_t size, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event, + cl_int cmd); + + cl_int + EnqueueSVMMemFill(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, + cl_int cmd); + + cl_int + EnqueueSVMUnmap(cl_command_queue command_queue, + void *svm_ptr, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event, + cl_int cmd); +} + +#endif diff --git a/src/gallium/frontends/clover/api/event.cpp b/src/gallium/frontends/clover/api/event.cpp new file mode 100644 index 00000000000..3f89644d0a4 --- /dev/null +++ b/src/gallium/frontends/clover/api/event.cpp @@ -0,0 +1,309 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/event.hpp" + +using namespace clover; + +CLOVER_API cl_event +clCreateUserEvent(cl_context d_ctx, cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + + ret_error(r_errcode, CL_SUCCESS); + return desc(new soft_event(ctx, {}, false)); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clSetUserEventStatus(cl_event d_ev, cl_int status) try { + auto &sev = obj<soft_event>(d_ev); + + if (status > 0) + return CL_INVALID_VALUE; + + if (sev.status() <= 0) + return CL_INVALID_OPERATION; + + if (status) + sev.abort(status); + else + sev.trigger(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clWaitForEvents(cl_uint num_evs, const cl_event *d_evs) try { + auto evs = objs(d_evs, num_evs); + + for (auto &ev : evs) { + if (ev.context() != evs.front().context()) + throw error(CL_INVALID_CONTEXT); + + if (ev.status() < 0) + throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST); + } + + // Create a temporary soft event that depends on all the events in + // the wait list + auto sev = create<soft_event>(evs.front().context(), evs, true); + + // ...and wait on it. + sev().wait(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetEventInfo(cl_event d_ev, cl_event_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &ev = obj(d_ev); + + switch (param) { + case CL_EVENT_COMMAND_QUEUE: + buf.as_scalar<cl_command_queue>() = desc(ev.queue()); + break; + + case CL_EVENT_CONTEXT: + buf.as_scalar<cl_context>() = desc(ev.context()); + break; + + case CL_EVENT_COMMAND_TYPE: + buf.as_scalar<cl_command_type>() = ev.command(); + break; + + case CL_EVENT_COMMAND_EXECUTION_STATUS: + buf.as_scalar<cl_int>() = ev.status(); + break; + + case CL_EVENT_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = ev.ref_count(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clSetEventCallback(cl_event d_ev, cl_int type, + void (CL_CALLBACK *pfn_notify)(cl_event, cl_int, void *), + void *user_data) try { + auto &ev = obj(d_ev); + + if (!pfn_notify || + (type != CL_COMPLETE && type != CL_SUBMITTED && type != CL_RUNNING)) + throw error(CL_INVALID_VALUE); + + // Create a temporary soft event that depends on ev, with + // pfn_notify as completion action. + create<soft_event>(ev.context(), ref_vector<event> { ev }, true, + [=, &ev](event &) { + ev.wait(); + pfn_notify(desc(ev), ev.status(), user_data); + }); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clRetainEvent(cl_event d_ev) try { + obj(d_ev).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseEvent(cl_event d_ev) try { + if (obj(d_ev).release()) + delete pobj(d_ev); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueMarker(cl_command_queue d_q, cl_event *rd_ev) try { + auto &q = obj(d_q); + + if (!rd_ev) + throw error(CL_INVALID_VALUE); + + *rd_ev = desc(new hard_event(q, CL_COMMAND_MARKER, {})); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueMarkerWithWaitList(cl_command_queue d_q, cl_uint num_deps, + const cl_event *d_deps, cl_event *rd_ev) try { + auto &q = obj(d_q); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + + for (auto &ev : deps) { + if (ev.context() != q.context()) + throw error(CL_INVALID_CONTEXT); + } + + // Create a hard event that depends on the events in the wait list: + // previous commands in the same queue are implicitly serialized + // with respect to it -- hard events always are. + auto hev = create<hard_event>(q, CL_COMMAND_MARKER, deps); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueBarrier(cl_command_queue d_q) try { + obj(d_q); + + // No need to do anything, q preserves data ordering strictly. + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueBarrierWithWaitList(cl_command_queue d_q, cl_uint num_deps, + const cl_event *d_deps, cl_event *rd_ev) try { + auto &q = obj(d_q); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + + for (auto &ev : deps) { + if (ev.context() != q.context()) + throw error(CL_INVALID_CONTEXT); + } + + // Create a hard event that depends on the events in the wait list: + // subsequent commands in the same queue will be implicitly + // serialized with respect to it -- hard events always are. + auto hev = create<hard_event>(q, CL_COMMAND_BARRIER, deps); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueWaitForEvents(cl_command_queue d_q, cl_uint num_evs, + const cl_event *d_evs) try { + // The wait list is mandatory for clEnqueueWaitForEvents(). + objs(d_evs, num_evs); + + return clEnqueueBarrierWithWaitList(d_q, num_evs, d_evs, NULL); + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetEventProfilingInfo(cl_event d_ev, cl_profiling_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + hard_event &hev = dynamic_cast<hard_event &>(obj(d_ev)); + + if (hev.status() != CL_COMPLETE) + throw error(CL_PROFILING_INFO_NOT_AVAILABLE); + + switch (param) { + case CL_PROFILING_COMMAND_QUEUED: + buf.as_scalar<cl_ulong>() = hev.time_queued(); + break; + + case CL_PROFILING_COMMAND_SUBMIT: + buf.as_scalar<cl_ulong>() = hev.time_submit(); + break; + + case CL_PROFILING_COMMAND_START: + buf.as_scalar<cl_ulong>() = hev.time_start(); + break; + + case CL_PROFILING_COMMAND_END: + buf.as_scalar<cl_ulong>() = hev.time_end(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (std::bad_cast &e) { + return CL_PROFILING_INFO_NOT_AVAILABLE; + +} catch (lazy<cl_ulong>::undefined_error &e) { + return CL_PROFILING_INFO_NOT_AVAILABLE; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clFinish(cl_command_queue d_q) try { + auto &q = obj(d_q); + + // Create a temporary hard event -- it implicitly depends on all + // the previously queued hard events. + auto hev = create<hard_event>(q, 0, ref_vector<event> {}); + + // And wait on it. + hev().wait(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/interop.cpp b/src/gallium/frontends/clover/api/interop.cpp new file mode 100644 index 00000000000..b96069f5167 --- /dev/null +++ b/src/gallium/frontends/clover/api/interop.cpp @@ -0,0 +1,69 @@ +// +// Copyright 2015 Advanced Micro Devices, Inc. +// All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "core/event.hpp" +#include "api/util.hpp" + +using namespace clover; + +extern "C" { + +PUBLIC bool +opencl_dri_event_add_ref(cl_event event) +{ + /* This should fail if the event hasn't been created by + * clEnqueueReleaseGLObjects or clEnqueueReleaseEGLObjects. + * + * TODO: implement the CL functions + */ + return false; /*return clRetainEvent(event) == CL_SUCCESS;*/ +} + +PUBLIC bool +opencl_dri_event_release(cl_event event) +{ + return clReleaseEvent(event) == CL_SUCCESS; +} + +PUBLIC bool +opencl_dri_event_wait(cl_event event, uint64_t timeout) try { + if (!timeout) { + return obj(event).status() == CL_COMPLETE; + } + + obj(event).wait(); + return true; + +} catch (error &) { + return false; +} + +PUBLIC struct pipe_fence_handle * +opencl_dri_event_get_fence(cl_event event) try { + return obj(event).fence(); + +} catch (error &) { + return NULL; +} + +} diff --git a/src/gallium/frontends/clover/api/kernel.cpp b/src/gallium/frontends/clover/api/kernel.cpp new file mode 100644 index 00000000000..31a87b63868 --- /dev/null +++ b/src/gallium/frontends/clover/api/kernel.cpp @@ -0,0 +1,390 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/kernel.hpp" +#include "core/event.hpp" + +using namespace clover; + +CLOVER_API cl_kernel +clCreateKernel(cl_program d_prog, const char *name, cl_int *r_errcode) try { + auto &prog = obj(d_prog); + + if (!name) + throw error(CL_INVALID_VALUE); + + auto &sym = find(name_equals(name), prog.symbols()); + + ret_error(r_errcode, CL_SUCCESS); + return new kernel(prog, name, range(sym.args)); + +} catch (std::out_of_range &e) { + ret_error(r_errcode, CL_INVALID_KERNEL_NAME); + return NULL; + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clCreateKernelsInProgram(cl_program d_prog, cl_uint count, + cl_kernel *rd_kerns, cl_uint *r_count) try { + auto &prog = obj(d_prog); + auto &syms = prog.symbols(); + + if (rd_kerns && count < syms.size()) + throw error(CL_INVALID_VALUE); + + if (rd_kerns) + copy(map([&](const module::symbol &sym) { + return desc(new kernel(prog, + std::string(sym.name.begin(), + sym.name.end()), + range(sym.args))); + }, syms), + rd_kerns); + + if (r_count) + *r_count = syms.size(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clRetainKernel(cl_kernel d_kern) try { + obj(d_kern).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseKernel(cl_kernel d_kern) try { + if (obj(d_kern).release()) + delete pobj(d_kern); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clSetKernelArg(cl_kernel d_kern, cl_uint idx, size_t size, + const void *value) try { + obj(d_kern).args().at(idx).set(size, value); + return CL_SUCCESS; + +} catch (std::out_of_range &e) { + return CL_INVALID_ARG_INDEX; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetKernelInfo(cl_kernel d_kern, cl_kernel_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &kern = obj(d_kern); + + switch (param) { + case CL_KERNEL_FUNCTION_NAME: + buf.as_string() = kern.name(); + break; + + case CL_KERNEL_NUM_ARGS: + buf.as_scalar<cl_uint>() = kern.args().size(); + break; + + case CL_KERNEL_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = kern.ref_count(); + break; + + case CL_KERNEL_CONTEXT: + buf.as_scalar<cl_context>() = desc(kern.program().context()); + break; + + case CL_KERNEL_PROGRAM: + buf.as_scalar<cl_program>() = desc(kern.program()); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetKernelWorkGroupInfo(cl_kernel d_kern, cl_device_id d_dev, + cl_kernel_work_group_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &kern = obj(d_kern); + auto &dev = (d_dev ? *pobj(d_dev) : unique(kern.program().devices())); + + if (!count(dev, kern.program().devices())) + throw error(CL_INVALID_DEVICE); + + switch (param) { + case CL_KERNEL_WORK_GROUP_SIZE: + buf.as_scalar<size_t>() = dev.max_threads_per_block(); + break; + + case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: + buf.as_vector<size_t>() = kern.required_block_size(); + break; + + case CL_KERNEL_LOCAL_MEM_SIZE: + buf.as_scalar<cl_ulong>() = kern.mem_local(); + break; + + case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: + buf.as_scalar<size_t>() = dev.subgroup_size(); + break; + + case CL_KERNEL_PRIVATE_MEM_SIZE: + buf.as_scalar<cl_ulong>() = kern.mem_private(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); + +} catch (std::out_of_range &e) { + return CL_INVALID_DEVICE; +} + +CLOVER_API cl_int +clGetKernelArgInfo(cl_kernel d_kern, + cl_uint idx, cl_kernel_arg_info param, + size_t size, void *r_buf, size_t *r_size) { + CLOVER_NOT_SUPPORTED_UNTIL("1.2"); + return CL_KERNEL_ARG_INFO_NOT_AVAILABLE; +} + +namespace { + /// + /// Common argument checking shared by kernel invocation commands. + /// + void + validate_common(const command_queue &q, kernel &kern, + const ref_vector<event> &deps) { + if (kern.program().context() != q.context() || + any_of([&](const event &ev) { + return ev.context() != q.context(); + }, deps)) + throw error(CL_INVALID_CONTEXT); + + if (any_of([](kernel::argument &arg) { + return !arg.set(); + }, kern.args())) + throw error(CL_INVALID_KERNEL_ARGS); + + // If the command queue's device is not associated to the program, we get + // a module, with no sections, which will also fail the following test. + auto &m = kern.program().build(q.device()).binary; + if (!any_of(type_equals(module::section::text_executable), m.secs)) + throw error(CL_INVALID_PROGRAM_EXECUTABLE); + } + + std::vector<size_t> + validate_grid_size(const command_queue &q, cl_uint dims, + const size_t *d_grid_size) { + auto grid_size = range(d_grid_size, dims); + + if (dims < 1 || dims > q.device().max_block_size().size()) + throw error(CL_INVALID_WORK_DIMENSION); + + if (!d_grid_size || any_of(is_zero(), grid_size)) + throw error(CL_INVALID_GLOBAL_WORK_SIZE); + + return grid_size; + } + + std::vector<size_t> + validate_grid_offset(const command_queue &q, cl_uint dims, + const size_t *d_grid_offset) { + if (d_grid_offset) + return range(d_grid_offset, dims); + else + return std::vector<size_t>(dims, 0); + } + + std::vector<size_t> + validate_block_size(const command_queue &q, const kernel &kern, + cl_uint dims, const size_t *d_grid_size, + const size_t *d_block_size) { + auto grid_size = range(d_grid_size, dims); + + if (d_block_size) { + auto block_size = range(d_block_size, dims); + + if (any_of(is_zero(), block_size) || + any_of(greater(), block_size, q.device().max_block_size())) + throw error(CL_INVALID_WORK_ITEM_SIZE); + + if (any_of(modulus(), grid_size, block_size)) + throw error(CL_INVALID_WORK_GROUP_SIZE); + + if (fold(multiplies(), 1u, block_size) > + q.device().max_threads_per_block()) + throw error(CL_INVALID_WORK_GROUP_SIZE); + + return block_size; + + } else { + return kern.optimal_block_size(q, grid_size); + } + } +} + +CLOVER_API cl_int +clEnqueueNDRangeKernel(cl_command_queue d_q, cl_kernel d_kern, + cl_uint dims, const size_t *d_grid_offset, + const size_t *d_grid_size, const size_t *d_block_size, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &kern = obj(d_kern); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto grid_size = validate_grid_size(q, dims, d_grid_size); + auto grid_offset = validate_grid_offset(q, dims, d_grid_offset); + auto block_size = validate_block_size(q, kern, dims, + d_grid_size, d_block_size); + + validate_common(q, kern, deps); + + auto hev = create<hard_event>( + q, CL_COMMAND_NDRANGE_KERNEL, deps, + [=, &kern, &q](event &) { + kern.launch(q, grid_offset, grid_size, block_size); + }); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueTask(cl_command_queue d_q, cl_kernel d_kern, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &kern = obj(d_kern); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + + validate_common(q, kern, deps); + + auto hev = create<hard_event>( + q, CL_COMMAND_TASK, deps, + [=, &kern, &q](event &) { + kern.launch(q, { 0 }, { 1 }, { 1 }); + }); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueNativeKernel(cl_command_queue d_q, void (*func)(void *), + void *args, size_t args_size, + cl_uint num_mems, const cl_mem *d_mems, + const void **mem_handles, cl_uint num_deps, + const cl_event *d_deps, cl_event *rd_ev) { + return CL_INVALID_OPERATION; +} + +CLOVER_API cl_int +clSetKernelArgSVMPointer(cl_kernel d_kern, + cl_uint arg_index, + const void *arg_value) try { + obj(d_kern).args().at(arg_index).set_svm(arg_value); + return CL_SUCCESS; + +} catch (std::out_of_range &e) { + return CL_INVALID_ARG_INDEX; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clSetKernelExecInfo(cl_kernel d_kern, + cl_kernel_exec_info param_name, + size_t param_value_size, + const void *param_value) try { + auto &kern = obj(d_kern); + const bool has_system_svm = all_of(std::mem_fn(&device::has_system_svm), + kern.program().context().devices()); + + if (!param_value) + return CL_INVALID_VALUE; + + switch (param_name) { + case CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM: + case CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM_ARM: { + if (param_value_size != sizeof(cl_bool)) + return CL_INVALID_VALUE; + + cl_bool val = *static_cast<const cl_bool*>(param_value); + if (val == CL_TRUE && !has_system_svm) + return CL_INVALID_OPERATION; + else + return CL_SUCCESS; + } + + case CL_KERNEL_EXEC_INFO_SVM_PTRS: + case CL_KERNEL_EXEC_INFO_SVM_PTRS_ARM: + if (has_system_svm) + return CL_SUCCESS; + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + + default: + return CL_INVALID_VALUE; + } + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/memory.cpp b/src/gallium/frontends/clover/api/memory.cpp new file mode 100644 index 00000000000..e03793339c1 --- /dev/null +++ b/src/gallium/frontends/clover/api/memory.cpp @@ -0,0 +1,497 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "util/u_math.h" +#include "api/util.hpp" +#include "core/memory.hpp" +#include "core/format.hpp" + +using namespace clover; + +namespace { + cl_mem_flags + validate_flags(cl_mem d_parent, cl_mem_flags d_flags, bool svm) { + const cl_mem_flags dev_access_flags = + CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY; + const cl_mem_flags host_ptr_flags = + CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR; + const cl_mem_flags host_access_flags = + CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS; + const cl_mem_flags svm_flags = + CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS; + + const cl_mem_flags valid_flags = + dev_access_flags + | (svm || d_parent ? 0 : host_ptr_flags) + | (svm ? svm_flags : host_access_flags); + + if ((d_flags & ~valid_flags) || + util_bitcount(d_flags & dev_access_flags) > 1 || + util_bitcount(d_flags & host_access_flags) > 1) + throw error(CL_INVALID_VALUE); + + if ((d_flags & CL_MEM_USE_HOST_PTR) && + (d_flags & (CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR))) + throw error(CL_INVALID_VALUE); + + if ((d_flags & CL_MEM_SVM_ATOMICS) && + !(d_flags & CL_MEM_SVM_FINE_GRAIN_BUFFER)) + throw error(CL_INVALID_VALUE); + + if (d_parent) { + const auto &parent = obj(d_parent); + const cl_mem_flags flags = (d_flags | + (d_flags & dev_access_flags ? 0 : + parent.flags() & dev_access_flags) | + (d_flags & host_access_flags ? 0 : + parent.flags() & host_access_flags) | + (parent.flags() & host_ptr_flags)); + + if (~flags & parent.flags() & (dev_access_flags & ~CL_MEM_READ_WRITE)) + throw error(CL_INVALID_VALUE); + + // Check if new host access flags cause a mismatch between + // host-read/write-only. + if (!(flags & CL_MEM_HOST_NO_ACCESS) && + (~flags & parent.flags() & host_access_flags)) + throw error(CL_INVALID_VALUE); + + return flags; + + } else { + return d_flags | (d_flags & dev_access_flags ? 0 : CL_MEM_READ_WRITE); + } + } +} + +CLOVER_API cl_mem +clCreateBuffer(cl_context d_ctx, cl_mem_flags d_flags, size_t size, + void *host_ptr, cl_int *r_errcode) try { + const cl_mem_flags flags = validate_flags(NULL, d_flags, false); + auto &ctx = obj(d_ctx); + + if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_COPY_HOST_PTR))) + throw error(CL_INVALID_HOST_PTR); + + if (!size || + size > fold(maximum(), cl_ulong(0), + map(std::mem_fn(&device::max_mem_alloc_size), ctx.devices()) + )) + throw error(CL_INVALID_BUFFER_SIZE); + + ret_error(r_errcode, CL_SUCCESS); + return new root_buffer(ctx, flags, size, host_ptr); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_mem +clCreateSubBuffer(cl_mem d_mem, cl_mem_flags d_flags, + cl_buffer_create_type op, + const void *op_info, cl_int *r_errcode) try { + auto &parent = obj<root_buffer>(d_mem); + const cl_mem_flags flags = validate_flags(d_mem, d_flags, false); + + if (op == CL_BUFFER_CREATE_TYPE_REGION) { + auto reg = reinterpret_cast<const cl_buffer_region *>(op_info); + + if (!reg || + reg->origin > parent.size() || + reg->origin + reg->size > parent.size()) + throw error(CL_INVALID_VALUE); + + if (!reg->size) + throw error(CL_INVALID_BUFFER_SIZE); + + ret_error(r_errcode, CL_SUCCESS); + return new sub_buffer(parent, flags, reg->origin, reg->size); + + } else { + throw error(CL_INVALID_VALUE); + } + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_mem +clCreateImage(cl_context d_ctx, cl_mem_flags d_flags, + const cl_image_format *format, + const cl_image_desc *desc, + void *host_ptr, cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + + if (!any_of(std::mem_fn(&device::image_support), ctx.devices())) + throw error(CL_INVALID_OPERATION); + + if (!format) + throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR); + + if (!desc) + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + + if (desc->image_array_size == 0 && + (desc->image_type == CL_MEM_OBJECT_IMAGE1D_ARRAY || + desc->image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY)) + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + + if (!host_ptr && + (desc->image_row_pitch || desc->image_slice_pitch)) + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + + if (desc->num_mip_levels || desc->num_samples) + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + + if (bool(desc->buffer) != (desc->image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER)) + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + + if (bool(host_ptr) != bool(d_flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_COPY_HOST_PTR))) + throw error(CL_INVALID_HOST_PTR); + + const cl_mem_flags flags = validate_flags(desc->buffer, d_flags, false); + + if (!supported_formats(ctx, desc->image_type).count(*format)) + throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED); + + ret_error(r_errcode, CL_SUCCESS); + + switch (desc->image_type) { + case CL_MEM_OBJECT_IMAGE2D: + if (!desc->image_width || !desc->image_height) + throw error(CL_INVALID_IMAGE_SIZE); + + if (all_of([=](const device &dev) { + const size_t max = 1 << dev.max_image_levels_2d(); + return (desc->image_width > max || + desc->image_height > max); + }, ctx.devices())) + throw error(CL_INVALID_IMAGE_SIZE); + + return new image2d(ctx, flags, format, + desc->image_width, desc->image_height, + desc->image_row_pitch, host_ptr); + + case CL_MEM_OBJECT_IMAGE3D: + if (!desc->image_width || !desc->image_height || !desc->image_depth) + throw error(CL_INVALID_IMAGE_SIZE); + + if (all_of([=](const device &dev) { + const size_t max = 1 << dev.max_image_levels_3d(); + return (desc->image_width > max || + desc->image_height > max || + desc->image_depth > max); + }, ctx.devices())) + throw error(CL_INVALID_IMAGE_SIZE); + + return new image3d(ctx, flags, format, + desc->image_width, desc->image_height, + desc->image_depth, desc->image_row_pitch, + desc->image_slice_pitch, host_ptr); + + case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE1D_ARRAY: + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + case CL_MEM_OBJECT_IMAGE2D_ARRAY: + // XXX - Not implemented. + throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED); + + default: + throw error(CL_INVALID_IMAGE_DESCRIPTOR); + } + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_mem +clCreateImage2D(cl_context d_ctx, cl_mem_flags d_flags, + const cl_image_format *format, + size_t width, size_t height, size_t row_pitch, + void *host_ptr, cl_int *r_errcode) { + const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, + row_pitch, 0, 0, 0, NULL }; + + return clCreateImage(d_ctx, d_flags, format, &desc, host_ptr, r_errcode); +} + +CLOVER_API cl_mem +clCreateImage3D(cl_context d_ctx, cl_mem_flags d_flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, + void *host_ptr, cl_int *r_errcode) { + const cl_image_desc desc = { CL_MEM_OBJECT_IMAGE3D, width, height, depth, 0, + row_pitch, slice_pitch, 0, 0, NULL }; + + return clCreateImage(d_ctx, d_flags, format, &desc, host_ptr, r_errcode); +} + +CLOVER_API cl_int +clGetSupportedImageFormats(cl_context d_ctx, cl_mem_flags flags, + cl_mem_object_type type, cl_uint count, + cl_image_format *r_buf, cl_uint *r_count) try { + auto &ctx = obj(d_ctx); + auto formats = supported_formats(ctx, type); + + validate_flags(NULL, flags, false); + + if (r_buf && !r_count) + throw error(CL_INVALID_VALUE); + + if (r_buf) + std::copy_n(formats.begin(), + std::min((cl_uint)formats.size(), count), + r_buf); + + if (r_count) + *r_count = formats.size(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetMemObjectInfo(cl_mem d_mem, cl_mem_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &mem = obj(d_mem); + + switch (param) { + case CL_MEM_TYPE: + buf.as_scalar<cl_mem_object_type>() = mem.type(); + break; + + case CL_MEM_FLAGS: + buf.as_scalar<cl_mem_flags>() = mem.flags(); + break; + + case CL_MEM_SIZE: + buf.as_scalar<size_t>() = mem.size(); + break; + + case CL_MEM_HOST_PTR: + buf.as_scalar<void *>() = mem.host_ptr(); + break; + + case CL_MEM_MAP_COUNT: + buf.as_scalar<cl_uint>() = 0; + break; + + case CL_MEM_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = mem.ref_count(); + break; + + case CL_MEM_CONTEXT: + buf.as_scalar<cl_context>() = desc(mem.context()); + break; + + case CL_MEM_ASSOCIATED_MEMOBJECT: { + sub_buffer *sub = dynamic_cast<sub_buffer *>(&mem); + buf.as_scalar<cl_mem>() = (sub ? desc(sub->parent()) : NULL); + break; + } + case CL_MEM_OFFSET: { + sub_buffer *sub = dynamic_cast<sub_buffer *>(&mem); + buf.as_scalar<size_t>() = (sub ? sub->offset() : 0); + break; + } + case CL_MEM_USES_SVM_POINTER: + case CL_MEM_USES_SVM_POINTER_ARM: { + // with system SVM all host ptrs are SVM pointers + // TODO: once we support devices with lower levels of SVM, we have to + // check the ptr in more detail + const bool system_svm = all_of(std::mem_fn(&device::has_system_svm), + mem.context().devices()); + buf.as_scalar<cl_bool>() = mem.host_ptr() && system_svm; + break; + } + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetImageInfo(cl_mem d_mem, cl_image_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &img = obj<image>(d_mem); + + switch (param) { + case CL_IMAGE_FORMAT: + buf.as_scalar<cl_image_format>() = img.format(); + break; + + case CL_IMAGE_ELEMENT_SIZE: + buf.as_scalar<size_t>() = 0; + break; + + case CL_IMAGE_ROW_PITCH: + buf.as_scalar<size_t>() = img.row_pitch(); + break; + + case CL_IMAGE_SLICE_PITCH: + buf.as_scalar<size_t>() = img.slice_pitch(); + break; + + case CL_IMAGE_WIDTH: + buf.as_scalar<size_t>() = img.width(); + break; + + case CL_IMAGE_HEIGHT: + buf.as_scalar<size_t>() = img.height(); + break; + + case CL_IMAGE_DEPTH: + buf.as_scalar<size_t>() = img.depth(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clRetainMemObject(cl_mem d_mem) try { + obj(d_mem).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseMemObject(cl_mem d_mem) try { + if (obj(d_mem).release()) + delete pobj(d_mem); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clSetMemObjectDestructorCallback(cl_mem d_mem, + void (CL_CALLBACK *pfn_notify)(cl_mem, void *), + void *user_data) try { + auto &mem = obj(d_mem); + + if (!pfn_notify) + return CL_INVALID_VALUE; + + mem.destroy_notify([=]{ pfn_notify(d_mem, user_data); }); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueFillBuffer(cl_command_queue command_queue, cl_mem buffer, + 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) { + CLOVER_NOT_SUPPORTED_UNTIL("1.2"); + return CL_INVALID_VALUE; +} + +CLOVER_API cl_int +clEnqueueFillImage(cl_command_queue command_queue, 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) { + CLOVER_NOT_SUPPORTED_UNTIL("1.2"); + return CL_INVALID_VALUE; +} + +CLOVER_API void * +clSVMAlloc(cl_context d_ctx, + cl_svm_mem_flags flags, + size_t size, + unsigned int alignment) try { + auto &ctx = obj(d_ctx); + validate_flags(NULL, flags, true); + + if (!size || + size > fold(minimum(), cl_ulong(ULONG_MAX), + map(std::mem_fn(&device::max_mem_alloc_size), ctx.devices()))) + return nullptr; + + if (!util_is_power_of_two_or_zero(alignment)) + return nullptr; + + if (!alignment) + alignment = 0x80; // sizeof(long16) + + bool can_emulate = all_of(std::mem_fn(&device::has_system_svm), ctx.devices()); + if (can_emulate) { + // we can ignore all the flags as it's not required to honor them. + void *ptr = nullptr; + if (alignment < sizeof(void*)) + alignment = sizeof(void*); + posix_memalign(&ptr, alignment, size); + return ptr; + } + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return nullptr; + +} catch (error &e) { + return nullptr; +} + +CLOVER_API void +clSVMFree(cl_context d_ctx, + void *svm_pointer) try { + auto &ctx = obj(d_ctx); + bool can_emulate = all_of(std::mem_fn(&device::has_system_svm), ctx.devices()); + + if (can_emulate) + return free(svm_pointer); + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + +} catch (error &e) { +} diff --git a/src/gallium/frontends/clover/api/platform.cpp b/src/gallium/frontends/clover/api/platform.cpp new file mode 100644 index 00000000000..7360461e62f --- /dev/null +++ b/src/gallium/frontends/clover/api/platform.cpp @@ -0,0 +1,235 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include <unordered_map> + +#include "api/dispatch.hpp" +#include "api/util.hpp" +#include "core/platform.hpp" +#include "git_sha1.h" +#include "util/u_debug.h" + +using namespace clover; + +namespace { + platform _clover_platform; +} + +CLOVER_API cl_int +clGetPlatformIDs(cl_uint num_entries, cl_platform_id *rd_platforms, + cl_uint *rnum_platforms) { + if ((!num_entries && rd_platforms) || + (!rnum_platforms && !rd_platforms)) + return CL_INVALID_VALUE; + + if (rnum_platforms) + *rnum_platforms = 1; + if (rd_platforms) + *rd_platforms = desc(_clover_platform); + + return CL_SUCCESS; +} + +cl_int +clover::GetPlatformInfo(cl_platform_id d_platform, cl_platform_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + + auto &platform = obj(d_platform); + + switch (param) { + case CL_PLATFORM_PROFILE: + buf.as_string() = "FULL_PROFILE"; + break; + + case CL_PLATFORM_VERSION: { + static const std::string version_string = + debug_get_option("CLOVER_PLATFORM_VERSION_OVERRIDE", "1.1"); + + buf.as_string() = "OpenCL " + version_string + " Mesa " PACKAGE_VERSION MESA_GIT_SHA1; + break; + } + case CL_PLATFORM_NAME: + buf.as_string() = "Clover"; + break; + + case CL_PLATFORM_VENDOR: + buf.as_string() = "Mesa"; + break; + + case CL_PLATFORM_EXTENSIONS: + buf.as_string() = platform.supported_extensions(); + break; + + case CL_PLATFORM_ICD_SUFFIX_KHR: + buf.as_string() = "MESA"; + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +void * +clover::GetExtensionFunctionAddressForPlatform(cl_platform_id d_platform, + const char *p_name) try { + obj(d_platform); + return GetExtensionFunctionAddress(p_name); + +} catch (error &e) { + return NULL; +} + +namespace { + +cl_int +enqueueSVMFreeARM(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) { + + return EnqueueSVMFree(command_queue, num_svm_pointers, svm_pointers, + pfn_free_func, user_data, num_events_in_wait_list, + event_wait_list, event, CL_COMMAND_SVM_FREE_ARM); +} + +cl_int +enqueueSVMMapARM(cl_command_queue command_queue, + cl_bool blocking_map, + cl_map_flags map_flags, + void *svm_ptr, + size_t size, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + + return EnqueueSVMMap(command_queue, blocking_map, map_flags, svm_ptr, size, + num_events_in_wait_list, event_wait_list, event, + CL_COMMAND_SVM_MAP_ARM); +} + +cl_int +enqueueSVMMemcpyARM(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) { + + return EnqueueSVMMemcpy(command_queue, blocking_copy, dst_ptr, src_ptr, + size, num_events_in_wait_list, event_wait_list, + event, CL_COMMAND_SVM_MEMCPY_ARM); +} + +cl_int +enqueueSVMMemFillARM(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) { + + return EnqueueSVMMemFill(command_queue, svm_ptr, pattern, pattern_size, + size, num_events_in_wait_list, event_wait_list, + event, CL_COMMAND_SVM_MEMFILL_ARM); +} + +cl_int +enqueueSVMUnmapARM(cl_command_queue command_queue, + void *svm_ptr, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + + return EnqueueSVMUnmap(command_queue, svm_ptr, num_events_in_wait_list, + event_wait_list, event, CL_COMMAND_SVM_UNMAP_ARM); +} + +const std::unordered_map<std::string, void *> +ext_funcs = { + // cl_arm_shared_virtual_memory + { "clEnqueueSVMFreeARM", reinterpret_cast<void *>(enqueueSVMFreeARM) }, + { "clEnqueueSVMMapARM", reinterpret_cast<void *>(enqueueSVMMapARM) }, + { "clEnqueueSVMMemcpyARM", reinterpret_cast<void *>(enqueueSVMMemcpyARM) }, + { "clEnqueueSVMMemFillARM", reinterpret_cast<void *>(enqueueSVMMemFillARM) }, + { "clEnqueueSVMUnmapARM", reinterpret_cast<void *>(enqueueSVMUnmapARM) }, + { "clSetKernelArgSVMPointerARM", reinterpret_cast<void *>(clSetKernelArgSVMPointer) }, + { "clSetKernelExecInfoARM", reinterpret_cast<void *>(clSetKernelExecInfo) }, + { "clSVMAllocARM", reinterpret_cast<void *>(clSVMAlloc) }, + { "clSVMFreeARM", reinterpret_cast<void *>(clSVMFree) }, + + // cl_khr_icd + { "clIcdGetPlatformIDsKHR", reinterpret_cast<void *>(IcdGetPlatformIDsKHR) }, +}; + +} // anonymous namespace + +void * +clover::GetExtensionFunctionAddress(const char *p_name) try { + return ext_funcs.at(p_name); +} catch (...) { + return nullptr; +} + +cl_int +clover::IcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms, + cl_uint *rnum_platforms) { + return clGetPlatformIDs(num_entries, rd_platforms, rnum_platforms); +} + +CLOVER_ICD_API cl_int +clGetPlatformInfo(cl_platform_id d_platform, cl_platform_info param, + size_t size, void *r_buf, size_t *r_size) { + return GetPlatformInfo(d_platform, param, size, r_buf, r_size); +} + +CLOVER_ICD_API void * +clGetExtensionFunctionAddress(const char *p_name) { + return GetExtensionFunctionAddress(p_name); +} + +CLOVER_ICD_API void * +clGetExtensionFunctionAddressForPlatform(cl_platform_id d_platform, + const char *p_name) { + return GetExtensionFunctionAddressForPlatform(d_platform, p_name); +} + +CLOVER_ICD_API cl_int +clIcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms, + cl_uint *rnum_platforms) { + return IcdGetPlatformIDsKHR(num_entries, rd_platforms, rnum_platforms); +} diff --git a/src/gallium/frontends/clover/api/program.cpp b/src/gallium/frontends/clover/api/program.cpp new file mode 100644 index 00000000000..33f843e9c87 --- /dev/null +++ b/src/gallium/frontends/clover/api/program.cpp @@ -0,0 +1,479 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/program.hpp" +#include "util/u_debug.h" + +#include <sstream> + +using namespace clover; + +namespace { + void + validate_build_common(const program &prog, cl_uint num_devs, + const cl_device_id *d_devs, + void (*pfn_notify)(cl_program, void *), + void *user_data) { + if (!pfn_notify && user_data) + throw error(CL_INVALID_VALUE); + + if (prog.kernel_ref_count()) + throw error(CL_INVALID_OPERATION); + + if (any_of([&](const device &dev) { + return !count(dev, prog.devices()); + }, objs<allow_empty_tag>(d_devs, num_devs))) + throw error(CL_INVALID_DEVICE); + } +} + +CLOVER_API cl_program +clCreateProgramWithSource(cl_context d_ctx, cl_uint count, + const char **strings, const size_t *lengths, + cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + std::string source; + + if (!count || !strings || + any_of(is_zero(), range(strings, count))) + throw error(CL_INVALID_VALUE); + + // Concatenate all the provided fragments together + for (unsigned i = 0; i < count; ++i) + source += (lengths && lengths[i] ? + std::string(strings[i], strings[i] + lengths[i]) : + std::string(strings[i])); + + // ...and create a program object for them. + ret_error(r_errcode, CL_SUCCESS); + return new program(ctx, source); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_program +clCreateProgramWithBinary(cl_context d_ctx, cl_uint n, + const cl_device_id *d_devs, + const size_t *lengths, + const unsigned char **binaries, + cl_int *r_status, cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + auto devs = objs(d_devs, n); + + if (!lengths || !binaries) + throw error(CL_INVALID_VALUE); + + if (any_of([&](const device &dev) { + return !count(dev, ctx.devices()); + }, devs)) + throw error(CL_INVALID_DEVICE); + + // Deserialize the provided binaries, + std::vector<std::pair<cl_int, module>> result = map( + [](const unsigned char *p, size_t l) -> std::pair<cl_int, module> { + if (!p || !l) + return { CL_INVALID_VALUE, {} }; + + try { + std::stringbuf bin( { (char*)p, l } ); + std::istream s(&bin); + + return { CL_SUCCESS, module::deserialize(s) }; + + } catch (std::istream::failure &e) { + return { CL_INVALID_BINARY, {} }; + } + }, + range(binaries, n), + range(lengths, n)); + + // update the status array, + if (r_status) + copy(map(keys(), result), r_status); + + if (any_of(key_equals(CL_INVALID_VALUE), result)) + throw error(CL_INVALID_VALUE); + + if (any_of(key_equals(CL_INVALID_BINARY), result)) + throw error(CL_INVALID_BINARY); + + // initialize a program object with them. + ret_error(r_errcode, CL_SUCCESS); + return new program(ctx, devs, map(values(), result)); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_program +clCreateProgramWithBuiltInKernels(cl_context d_ctx, cl_uint n, + const cl_device_id *d_devs, + const char *kernel_names, + cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + auto devs = objs(d_devs, n); + + if (any_of([&](const device &dev) { + return !count(dev, ctx.devices()); + }, devs)) + throw error(CL_INVALID_DEVICE); + + // No currently supported built-in kernels. + throw error(CL_INVALID_VALUE); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + + +CLOVER_API cl_int +clRetainProgram(cl_program d_prog) try { + obj(d_prog).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseProgram(cl_program d_prog) try { + if (obj(d_prog).release()) + delete pobj(d_prog); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clBuildProgram(cl_program d_prog, cl_uint num_devs, + const cl_device_id *d_devs, const char *p_opts, + void (*pfn_notify)(cl_program, void *), + void *user_data) try { + auto &prog = obj(d_prog); + auto devs = + (d_devs ? objs(d_devs, num_devs) : ref_vector<device>(prog.devices())); + const auto opts = std::string(p_opts ? p_opts : "") + " " + + debug_get_option("CLOVER_EXTRA_BUILD_OPTIONS", ""); + + validate_build_common(prog, num_devs, d_devs, pfn_notify, user_data); + + if (prog.has_source) { + prog.compile(devs, opts); + prog.link(devs, opts, { prog }); + } else if (any_of([&](const device &dev){ + return prog.build(dev).binary_type() != CL_PROGRAM_BINARY_TYPE_EXECUTABLE; + }, devs)) { + // According to the OpenCL 1.2 specification, “if program is created + // with clCreateProgramWithBinary, then the program binary must be an + // executable binary (not a compiled binary or library).” + throw error(CL_INVALID_BINARY); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clCompileProgram(cl_program d_prog, cl_uint num_devs, + const cl_device_id *d_devs, const char *p_opts, + cl_uint num_headers, const cl_program *d_header_progs, + const char **header_names, + void (*pfn_notify)(cl_program, void *), + void *user_data) try { + auto &prog = obj(d_prog); + auto devs = + (d_devs ? objs(d_devs, num_devs) : ref_vector<device>(prog.devices())); + const auto opts = std::string(p_opts ? p_opts : "") + " " + + debug_get_option("CLOVER_EXTRA_COMPILE_OPTIONS", ""); + header_map headers; + + validate_build_common(prog, num_devs, d_devs, pfn_notify, user_data); + + if (bool(num_headers) != bool(header_names)) + throw error(CL_INVALID_VALUE); + + if (!prog.has_source) + throw error(CL_INVALID_OPERATION); + + for_each([&](const char *name, const program &header) { + if (!header.has_source) + throw error(CL_INVALID_OPERATION); + + if (!any_of(key_equals(name), headers)) + headers.push_back(std::pair<std::string, std::string>( + name, header.source())); + }, + range(header_names, num_headers), + objs<allow_empty_tag>(d_header_progs, num_headers)); + + prog.compile(devs, opts, headers); + return CL_SUCCESS; + +} catch (invalid_build_options_error &e) { + return CL_INVALID_COMPILER_OPTIONS; + +} catch (build_error &e) { + return CL_COMPILE_PROGRAM_FAILURE; + +} catch (error &e) { + return e.get(); +} + +namespace { + ref_vector<device> + validate_link_devices(const ref_vector<program> &progs, + const ref_vector<device> &all_devs, + const std::string &opts) { + std::vector<device *> devs; + const bool create_library = + opts.find("-create-library") != std::string::npos; + const bool enable_link_options = + opts.find("-enable-link-options") != std::string::npos; + const bool has_link_options = + opts.find("-cl-denorms-are-zero") != std::string::npos || + opts.find("-cl-no-signed-zeroes") != std::string::npos || + opts.find("-cl-unsafe-math-optimizations") != std::string::npos || + opts.find("-cl-finite-math-only") != std::string::npos || + opts.find("-cl-fast-relaxed-math") != std::string::npos || + opts.find("-cl-no-subgroup-ifp") != std::string::npos; + + // According to the OpenCL 1.2 specification, "[the + // -enable-link-options] option must be specified with the + // create-library option". + if (enable_link_options && !create_library) + throw error(CL_INVALID_LINKER_OPTIONS); + + // According to the OpenCL 1.2 specification, "the + // [program linking options] can be specified when linking a program + // executable". + if (has_link_options && create_library) + throw error(CL_INVALID_LINKER_OPTIONS); + + for (auto &dev : all_devs) { + const auto has_binary = [&](const program &prog) { + const auto t = prog.build(dev).binary_type(); + return t == CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT || + t == CL_PROGRAM_BINARY_TYPE_LIBRARY; + }; + + // According to the OpenCL 1.2 specification, a library is made of + // “compiled binaries specified in input_programs argument to + // clLinkProgram“; compiled binaries does not refer to libraries: + // “input_programs is an array of program objects that are compiled + // binaries or libraries that are to be linked to create the program + // executable”. + if (create_library && any_of([&](const program &prog) { + const auto t = prog.build(dev).binary_type(); + return t != CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT; + }, progs)) + throw error(CL_INVALID_OPERATION); + + // According to the CL 1.2 spec, when "all programs specified [..] + // contain a compiled binary or library for the device [..] a link is + // performed", + else if (all_of(has_binary, progs)) + devs.push_back(&dev); + + // otherwise if "none of the programs contain a compiled binary or + // library for that device [..] no link is performed. All other + // cases will return a CL_INVALID_OPERATION error." + else if (any_of(has_binary, progs)) + throw error(CL_INVALID_OPERATION); + + // According to the OpenCL 1.2 specification, "[t]he linker may apply + // [program linking options] to all compiled program objects + // specified to clLinkProgram. The linker may apply these options + // only to libraries which were created with the + // -enable-link-option." + else if (has_link_options && any_of([&](const program &prog) { + const auto t = prog.build(dev).binary_type(); + return !(t == CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT || + (t == CL_PROGRAM_BINARY_TYPE_LIBRARY && + prog.build(dev).opts.find("-enable-link-options") != + std::string::npos)); + }, progs)) + throw error(CL_INVALID_LINKER_OPTIONS); + } + + return map(derefs(), devs); + } +} + +CLOVER_API cl_program +clLinkProgram(cl_context d_ctx, cl_uint num_devs, const cl_device_id *d_devs, + const char *p_opts, cl_uint num_progs, const cl_program *d_progs, + void (*pfn_notify) (cl_program, void *), void *user_data, + cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + const auto opts = std::string(p_opts ? p_opts : "") + " " + + debug_get_option("CLOVER_EXTRA_LINK_OPTIONS", ""); + auto progs = objs(d_progs, num_progs); + auto all_devs = + (d_devs ? objs(d_devs, num_devs) : ref_vector<device>(ctx.devices())); + auto prog = create<program>(ctx, all_devs); + auto devs = validate_link_devices(progs, all_devs, opts); + + validate_build_common(prog, num_devs, d_devs, pfn_notify, user_data); + + try { + prog().link(devs, opts, progs); + ret_error(r_errcode, CL_SUCCESS); + + } catch (build_error &e) { + ret_error(r_errcode, CL_LINK_PROGRAM_FAILURE); + } + + return ret_object(prog); + +} catch (invalid_build_options_error &e) { + ret_error(r_errcode, CL_INVALID_LINKER_OPTIONS); + return NULL; + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clUnloadCompiler() { + return CL_SUCCESS; +} + +CLOVER_API cl_int +clUnloadPlatformCompiler(cl_platform_id d_platform) { + return CL_SUCCESS; +} + +CLOVER_API cl_int +clGetProgramInfo(cl_program d_prog, cl_program_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &prog = obj(d_prog); + + switch (param) { + case CL_PROGRAM_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = prog.ref_count(); + break; + + case CL_PROGRAM_CONTEXT: + buf.as_scalar<cl_context>() = desc(prog.context()); + break; + + case CL_PROGRAM_NUM_DEVICES: + buf.as_scalar<cl_uint>() = (prog.devices().size() ? + prog.devices().size() : + prog.context().devices().size()); + break; + + case CL_PROGRAM_DEVICES: + buf.as_vector<cl_device_id>() = (prog.devices().size() ? + descs(prog.devices()) : + descs(prog.context().devices())); + break; + + case CL_PROGRAM_SOURCE: + buf.as_string() = prog.source(); + break; + + case CL_PROGRAM_BINARY_SIZES: + buf.as_vector<size_t>() = map([&](const device &dev) { + return prog.build(dev).binary.size(); + }, + prog.devices()); + break; + + case CL_PROGRAM_BINARIES: + buf.as_matrix<unsigned char>() = map([&](const device &dev) { + std::stringbuf bin; + std::ostream s(&bin); + prog.build(dev).binary.serialize(s); + return bin.str(); + }, + prog.devices()); + break; + + case CL_PROGRAM_NUM_KERNELS: + buf.as_scalar<cl_uint>() = prog.symbols().size(); + break; + + case CL_PROGRAM_KERNEL_NAMES: + buf.as_string() = fold([](const std::string &a, const module::symbol &s) { + return ((a.empty() ? "" : a + ";") + s.name); + }, std::string(), prog.symbols()); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetProgramBuildInfo(cl_program d_prog, cl_device_id d_dev, + cl_program_build_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &prog = obj(d_prog); + auto &dev = obj(d_dev); + + if (!count(dev, prog.context().devices())) + return CL_INVALID_DEVICE; + + switch (param) { + case CL_PROGRAM_BUILD_STATUS: + buf.as_scalar<cl_build_status>() = prog.build(dev).status(); + break; + + case CL_PROGRAM_BUILD_OPTIONS: + buf.as_string() = prog.build(dev).opts; + break; + + case CL_PROGRAM_BUILD_LOG: + buf.as_string() = prog.build(dev).log; + break; + + case CL_PROGRAM_BINARY_TYPE: + buf.as_scalar<cl_program_binary_type>() = prog.build(dev).binary_type(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/queue.cpp b/src/gallium/frontends/clover/api/queue.cpp new file mode 100644 index 00000000000..65b271b216f --- /dev/null +++ b/src/gallium/frontends/clover/api/queue.cpp @@ -0,0 +1,135 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/queue.hpp" + +using namespace clover; + +CLOVER_API cl_command_queue +clCreateCommandQueue(cl_context d_ctx, cl_device_id d_dev, + cl_command_queue_properties props, + cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + auto &dev = obj(d_dev); + + if (!count(dev, ctx.devices())) + throw error(CL_INVALID_DEVICE); + + if (props & ~(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | + CL_QUEUE_PROFILING_ENABLE)) + throw error(CL_INVALID_VALUE); + + ret_error(r_errcode, CL_SUCCESS); + return new command_queue(ctx, dev, props); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clRetainCommandQueue(cl_command_queue d_q) try { + obj(d_q).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseCommandQueue(cl_command_queue d_q) try { + auto &q = obj(d_q); + + q.flush(); + + if (q.release()) + delete pobj(d_q); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetCommandQueueInfo(cl_command_queue d_q, cl_command_queue_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &q = obj(d_q); + + switch (param) { + case CL_QUEUE_CONTEXT: + buf.as_scalar<cl_context>() = desc(q.context()); + break; + + case CL_QUEUE_DEVICE: + buf.as_scalar<cl_device_id>() = desc(q.device()); + break; + + case CL_QUEUE_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = q.ref_count(); + break; + + case CL_QUEUE_PROPERTIES: + buf.as_scalar<cl_command_queue_properties>() = q.properties(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clFlush(cl_command_queue d_q) try { + obj(d_q).flush(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_command_queue +clCreateCommandQueueWithProperties(cl_context context, cl_device_id device, + const cl_queue_properties *properties, + cl_int *errcode_ret) try { + cl_command_queue_properties props = 0; + if (properties) { + for (auto idx = 0; properties[idx]; idx += 2) { + if (properties[idx] == CL_QUEUE_PROPERTIES) + props |= properties[idx + 1]; + else + throw error(CL_INVALID_VALUE); + } + } + + return clCreateCommandQueue(context, device, props, errcode_ret); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} diff --git a/src/gallium/frontends/clover/api/sampler.cpp b/src/gallium/frontends/clover/api/sampler.cpp new file mode 100644 index 00000000000..482e55a9ce9 --- /dev/null +++ b/src/gallium/frontends/clover/api/sampler.cpp @@ -0,0 +1,100 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include "api/util.hpp" +#include "core/sampler.hpp" + +using namespace clover; + +CLOVER_API cl_sampler +clCreateSampler(cl_context d_ctx, cl_bool norm_mode, + cl_addressing_mode addr_mode, cl_filter_mode filter_mode, + cl_int *r_errcode) try { + auto &ctx = obj(d_ctx); + + if (!any_of(std::mem_fn(&device::image_support), ctx.devices())) + throw error(CL_INVALID_OPERATION); + + ret_error(r_errcode, CL_SUCCESS); + return new sampler(ctx, norm_mode, addr_mode, filter_mode); + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clRetainSampler(cl_sampler d_s) try { + obj(d_s).retain(); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clReleaseSampler(cl_sampler d_s) try { + if (obj(d_s).release()) + delete pobj(d_s); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clGetSamplerInfo(cl_sampler d_s, cl_sampler_info param, + size_t size, void *r_buf, size_t *r_size) try { + property_buffer buf { r_buf, size, r_size }; + auto &s = obj(d_s); + + switch (param) { + case CL_SAMPLER_REFERENCE_COUNT: + buf.as_scalar<cl_uint>() = s.ref_count(); + break; + + case CL_SAMPLER_CONTEXT: + buf.as_scalar<cl_context>() = desc(s.context()); + break; + + case CL_SAMPLER_NORMALIZED_COORDS: + buf.as_scalar<cl_bool>() = s.norm_mode(); + break; + + case CL_SAMPLER_ADDRESSING_MODE: + buf.as_scalar<cl_addressing_mode>() = s.addr_mode(); + break; + + case CL_SAMPLER_FILTER_MODE: + buf.as_scalar<cl_filter_mode>() = s.filter_mode(); + break; + + default: + throw error(CL_INVALID_VALUE); + } + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/frontends/clover/api/transfer.cpp b/src/gallium/frontends/clover/api/transfer.cpp new file mode 100644 index 00000000000..fa8741e02b4 --- /dev/null +++ b/src/gallium/frontends/clover/api/transfer.cpp @@ -0,0 +1,1059 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#include <cstring> + +#include "util/bitscan.h" + +#include "api/dispatch.hpp" +#include "api/util.hpp" +#include "core/event.hpp" +#include "core/memory.hpp" + +using namespace clover; + +namespace { + typedef resource::vector vector_t; + + vector_t + vector(const size_t *p) { + return range(p, 3); + } + + vector_t + pitch(const vector_t ®ion, vector_t pitch) { + for (auto x : zip(tail(pitch), + map(multiplies(), region, pitch))) { + // The spec defines a value of zero as the natural pitch, + // i.e. the unaligned size of the previous dimension. + if (std::get<0>(x) == 0) + std::get<0>(x) = std::get<1>(x); + } + + return pitch; + } + + /// + /// Size of a region in bytes. + /// + size_t + size(const vector_t &pitch, const vector_t ®ion) { + if (any_of(is_zero(), region)) + return 0; + else + return dot(pitch, region - vector_t{ 0, 1, 1 }); + } + + /// + /// Common argument checking shared by memory transfer commands. + /// + void + validate_common(command_queue &q, + const ref_vector<event> &deps) { + if (any_of([&](const event &ev) { + return ev.context() != q.context(); + }, deps)) + throw error(CL_INVALID_CONTEXT); + } + + /// + /// Common error checking for a buffer object argument. + /// + void + validate_object(command_queue &q, buffer &mem, const vector_t &origin, + const vector_t &pitch, const vector_t ®ion) { + if (mem.context() != q.context()) + throw error(CL_INVALID_CONTEXT); + + // The region must fit within the specified pitch, + if (any_of(greater(), map(multiplies(), pitch, region), tail(pitch))) + throw error(CL_INVALID_VALUE); + + // ...and within the specified object. + if (dot(pitch, origin) + size(pitch, region) > mem.size()) + throw error(CL_INVALID_VALUE); + + if (any_of(is_zero(), region)) + throw error(CL_INVALID_VALUE); + } + + /// + /// Common error checking for an image argument. + /// + void + validate_object(command_queue &q, image &img, + const vector_t &orig, const vector_t ®ion) { + vector_t size = { img.width(), img.height(), img.depth() }; + + if (!q.device().image_support()) + throw error(CL_INVALID_OPERATION); + + if (img.context() != q.context()) + throw error(CL_INVALID_CONTEXT); + + if (any_of(greater(), orig + region, size)) + throw error(CL_INVALID_VALUE); + + if (any_of(is_zero(), region)) + throw error(CL_INVALID_VALUE); + } + + /// + /// Common error checking for a host pointer argument. + /// + void + validate_object(command_queue &q, const void *ptr, const vector_t &orig, + const vector_t &pitch, const vector_t ®ion) { + if (!ptr) + throw error(CL_INVALID_VALUE); + + // The region must fit within the specified pitch. + if (any_of(greater(), map(multiplies(), pitch, region), tail(pitch))) + throw error(CL_INVALID_VALUE); + } + + /// + /// Common argument checking for a copy between two buffer objects. + /// + void + validate_copy(command_queue &q, buffer &dst_mem, + const vector_t &dst_orig, const vector_t &dst_pitch, + buffer &src_mem, + const vector_t &src_orig, const vector_t &src_pitch, + const vector_t ®ion) { + if (dst_mem == src_mem) { + auto dst_offset = dot(dst_pitch, dst_orig); + auto src_offset = dot(src_pitch, src_orig); + + if (interval_overlaps()( + dst_offset, dst_offset + size(dst_pitch, region), + src_offset, src_offset + size(src_pitch, region))) + throw error(CL_MEM_COPY_OVERLAP); + } + } + + /// + /// Common argument checking for a copy between two image objects. + /// + void + validate_copy(command_queue &q, + image &dst_img, const vector_t &dst_orig, + image &src_img, const vector_t &src_orig, + const vector_t ®ion) { + if (dst_img.format() != src_img.format()) + throw error(CL_IMAGE_FORMAT_MISMATCH); + + if (dst_img == src_img) { + if (all_of(interval_overlaps(), + dst_orig, dst_orig + region, + src_orig, src_orig + region)) + throw error(CL_MEM_COPY_OVERLAP); + } + } + + /// + /// Checks that the host access flags of the memory object are + /// within the allowed set \a flags. + /// + void + validate_object_access(const memory_obj &mem, const cl_mem_flags flags) { + if (mem.flags() & ~flags & + (CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_WRITE_ONLY | + CL_MEM_HOST_NO_ACCESS)) + throw error(CL_INVALID_OPERATION); + } + + /// + /// Checks that the mapping flags are correct. + /// + void + validate_map_flags(const memory_obj &mem, const cl_map_flags flags) { + if ((flags & (CL_MAP_WRITE | CL_MAP_READ)) && + (flags & CL_MAP_WRITE_INVALIDATE_REGION)) + throw error(CL_INVALID_VALUE); + + if (flags & CL_MAP_READ) + validate_object_access(mem, CL_MEM_HOST_READ_ONLY); + + if (flags & (CL_MAP_WRITE | CL_MAP_WRITE_INVALIDATE_REGION)) + validate_object_access(mem, CL_MEM_HOST_WRITE_ONLY); + } + + /// + /// Class that encapsulates the task of mapping an object of type + /// \a T. The return value of get() should be implicitly + /// convertible to \a void *. + /// + template<typename T> + struct _map { + static mapping + get(command_queue &q, T obj, cl_map_flags flags, + size_t offset, size_t size) { + return { q, obj->resource(q), flags, true, + {{ offset }}, {{ size, 1, 1 }} }; + } + }; + + template<> + struct _map<void *> { + static void * + get(command_queue &q, void *obj, cl_map_flags flags, + size_t offset, size_t size) { + return (char *)obj + offset; + } + }; + + template<> + struct _map<const void *> { + static const void * + get(command_queue &q, const void *obj, cl_map_flags flags, + size_t offset, size_t size) { + return (const char *)obj + offset; + } + }; + + /// + /// Software copy from \a src_obj to \a dst_obj. They can be + /// either pointers or memory objects. + /// + template<typename T, typename S> + std::function<void (event &)> + soft_copy_op(command_queue &q, + T dst_obj, const vector_t &dst_orig, const vector_t &dst_pitch, + S src_obj, const vector_t &src_orig, const vector_t &src_pitch, + const vector_t ®ion) { + return [=, &q](event &) { + auto dst = _map<T>::get(q, dst_obj, CL_MAP_WRITE, + dot(dst_pitch, dst_orig), + size(dst_pitch, region)); + auto src = _map<S>::get(q, src_obj, CL_MAP_READ, + dot(src_pitch, src_orig), + size(src_pitch, region)); + vector_t v = {}; + + for (v[2] = 0; v[2] < region[2]; ++v[2]) { + for (v[1] = 0; v[1] < region[1]; ++v[1]) { + std::memcpy( + static_cast<char *>(dst) + dot(dst_pitch, v), + static_cast<const char *>(src) + dot(src_pitch, v), + src_pitch[0] * region[0]); + } + } + }; + } + + /// + /// Hardware copy from \a src_obj to \a dst_obj. + /// + template<typename T, typename S> + std::function<void (event &)> + hard_copy_op(command_queue &q, T dst_obj, const vector_t &dst_orig, + S src_obj, const vector_t &src_orig, const vector_t ®ion) { + return [=, &q](event &) { + dst_obj->resource(q).copy(q, dst_orig, region, + src_obj->resource(q), src_orig); + }; + } +} + +CLOVER_API cl_int +clEnqueueReadBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + size_t offset, size_t size, void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &mem = obj<buffer>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + vector_t region = { size, 1, 1 }; + vector_t obj_origin = { offset }; + auto obj_pitch = pitch(region, {{ 1 }}); + + validate_common(q, deps); + validate_object(q, ptr, {}, obj_pitch, region); + validate_object(q, mem, obj_origin, obj_pitch, region); + validate_object_access(mem, CL_MEM_HOST_READ_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_READ_BUFFER, deps, + soft_copy_op(q, ptr, {}, obj_pitch, + &mem, obj_origin, obj_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueWriteBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + size_t offset, size_t size, const void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &mem = obj<buffer>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + vector_t region = { size, 1, 1 }; + vector_t obj_origin = { offset }; + auto obj_pitch = pitch(region, {{ 1 }}); + + validate_common(q, deps); + validate_object(q, mem, obj_origin, obj_pitch, region); + validate_object(q, ptr, {}, obj_pitch, region); + validate_object_access(mem, CL_MEM_HOST_WRITE_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_WRITE_BUFFER, deps, + soft_copy_op(q, &mem, obj_origin, obj_pitch, + ptr, {}, obj_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueReadBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + const size_t *p_obj_origin, + const size_t *p_host_origin, + const size_t *p_region, + size_t obj_row_pitch, size_t obj_slice_pitch, + size_t host_row_pitch, size_t host_slice_pitch, + void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &mem = obj<buffer>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto obj_origin = vector(p_obj_origin); + auto obj_pitch = pitch(region, {{ 1, obj_row_pitch, obj_slice_pitch }}); + auto host_origin = vector(p_host_origin); + auto host_pitch = pitch(region, {{ 1, host_row_pitch, host_slice_pitch }}); + + validate_common(q, deps); + validate_object(q, ptr, host_origin, host_pitch, region); + validate_object(q, mem, obj_origin, obj_pitch, region); + validate_object_access(mem, CL_MEM_HOST_READ_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_READ_BUFFER_RECT, deps, + soft_copy_op(q, ptr, host_origin, host_pitch, + &mem, obj_origin, obj_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueWriteBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + const size_t *p_obj_origin, + const size_t *p_host_origin, + const size_t *p_region, + size_t obj_row_pitch, size_t obj_slice_pitch, + size_t host_row_pitch, size_t host_slice_pitch, + const void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &mem = obj<buffer>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto obj_origin = vector(p_obj_origin); + auto obj_pitch = pitch(region, {{ 1, obj_row_pitch, obj_slice_pitch }}); + auto host_origin = vector(p_host_origin); + auto host_pitch = pitch(region, {{ 1, host_row_pitch, host_slice_pitch }}); + + validate_common(q, deps); + validate_object(q, mem, obj_origin, obj_pitch, region); + validate_object(q, ptr, host_origin, host_pitch, region); + validate_object_access(mem, CL_MEM_HOST_WRITE_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_WRITE_BUFFER_RECT, deps, + soft_copy_op(q, &mem, obj_origin, obj_pitch, + ptr, host_origin, host_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueCopyBuffer(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem, + size_t src_offset, size_t dst_offset, size_t size, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &src_mem = obj<buffer>(d_src_mem); + auto &dst_mem = obj<buffer>(d_dst_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + vector_t region = { size, 1, 1 }; + vector_t dst_origin = { dst_offset }; + auto dst_pitch = pitch(region, {{ 1 }}); + vector_t src_origin = { src_offset }; + auto src_pitch = pitch(region, {{ 1 }}); + + validate_common(q, deps); + validate_object(q, dst_mem, dst_origin, dst_pitch, region); + validate_object(q, src_mem, src_origin, src_pitch, region); + validate_copy(q, dst_mem, dst_origin, dst_pitch, + src_mem, src_origin, src_pitch, region); + + auto hev = create<hard_event>( + q, CL_COMMAND_COPY_BUFFER, deps, + hard_copy_op(q, &dst_mem, dst_origin, + &src_mem, src_origin, region)); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueCopyBufferRect(cl_command_queue d_q, cl_mem d_src_mem, + cl_mem d_dst_mem, + const size_t *p_src_origin, const size_t *p_dst_origin, + const size_t *p_region, + size_t src_row_pitch, size_t src_slice_pitch, + size_t dst_row_pitch, size_t dst_slice_pitch, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &src_mem = obj<buffer>(d_src_mem); + auto &dst_mem = obj<buffer>(d_dst_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto dst_origin = vector(p_dst_origin); + auto dst_pitch = pitch(region, {{ 1, dst_row_pitch, dst_slice_pitch }}); + auto src_origin = vector(p_src_origin); + auto src_pitch = pitch(region, {{ 1, src_row_pitch, src_slice_pitch }}); + + validate_common(q, deps); + validate_object(q, dst_mem, dst_origin, dst_pitch, region); + validate_object(q, src_mem, src_origin, src_pitch, region); + validate_copy(q, dst_mem, dst_origin, dst_pitch, + src_mem, src_origin, src_pitch, region); + + auto hev = create<hard_event>( + q, CL_COMMAND_COPY_BUFFER_RECT, deps, + soft_copy_op(q, &dst_mem, dst_origin, dst_pitch, + &src_mem, src_origin, src_pitch, + region)); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueReadImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + const size_t *p_origin, const size_t *p_region, + size_t row_pitch, size_t slice_pitch, void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &img = obj<image>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto dst_pitch = pitch(region, {{ img.pixel_size(), + row_pitch, slice_pitch }}); + auto src_origin = vector(p_origin); + auto src_pitch = pitch(region, {{ img.pixel_size(), + img.row_pitch(), img.slice_pitch() }}); + + validate_common(q, deps); + validate_object(q, ptr, {}, dst_pitch, region); + validate_object(q, img, src_origin, region); + validate_object_access(img, CL_MEM_HOST_READ_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_READ_IMAGE, deps, + soft_copy_op(q, ptr, {}, dst_pitch, + &img, src_origin, src_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueWriteImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + const size_t *p_origin, const size_t *p_region, + size_t row_pitch, size_t slice_pitch, const void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &img = obj<image>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto dst_origin = vector(p_origin); + auto dst_pitch = pitch(region, {{ img.pixel_size(), + img.row_pitch(), img.slice_pitch() }}); + auto src_pitch = pitch(region, {{ img.pixel_size(), + row_pitch, slice_pitch }}); + + validate_common(q, deps); + validate_object(q, img, dst_origin, region); + validate_object(q, ptr, {}, src_pitch, region); + validate_object_access(img, CL_MEM_HOST_WRITE_ONLY); + + auto hev = create<hard_event>( + q, CL_COMMAND_WRITE_IMAGE, deps, + soft_copy_op(q, &img, dst_origin, dst_pitch, + ptr, {}, src_pitch, + region)); + + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueCopyImage(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem, + const size_t *p_src_origin, const size_t *p_dst_origin, + const size_t *p_region, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &src_img = obj<image>(d_src_mem); + auto &dst_img = obj<image>(d_dst_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto dst_origin = vector(p_dst_origin); + auto src_origin = vector(p_src_origin); + + validate_common(q, deps); + validate_object(q, dst_img, dst_origin, region); + validate_object(q, src_img, src_origin, region); + validate_copy(q, dst_img, dst_origin, src_img, src_origin, region); + + auto hev = create<hard_event>( + q, CL_COMMAND_COPY_IMAGE, deps, + hard_copy_op(q, &dst_img, dst_origin, + &src_img, src_origin, + region)); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueCopyImageToBuffer(cl_command_queue d_q, + cl_mem d_src_mem, cl_mem d_dst_mem, + const size_t *p_src_origin, const size_t *p_region, + size_t dst_offset, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &src_img = obj<image>(d_src_mem); + auto &dst_mem = obj<buffer>(d_dst_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + vector_t dst_origin = { dst_offset }; + auto dst_pitch = pitch(region, {{ src_img.pixel_size() }}); + auto src_origin = vector(p_src_origin); + auto src_pitch = pitch(region, {{ src_img.pixel_size(), + src_img.row_pitch(), + src_img.slice_pitch() }}); + + validate_common(q, deps); + validate_object(q, dst_mem, dst_origin, dst_pitch, region); + validate_object(q, src_img, src_origin, region); + + auto hev = create<hard_event>( + q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, deps, + soft_copy_op(q, &dst_mem, dst_origin, dst_pitch, + &src_img, src_origin, src_pitch, + region)); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueCopyBufferToImage(cl_command_queue d_q, + cl_mem d_src_mem, cl_mem d_dst_mem, + size_t src_offset, + const size_t *p_dst_origin, const size_t *p_region, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &src_mem = obj<buffer>(d_src_mem); + auto &dst_img = obj<image>(d_dst_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto dst_origin = vector(p_dst_origin); + auto dst_pitch = pitch(region, {{ dst_img.pixel_size(), + dst_img.row_pitch(), + dst_img.slice_pitch() }}); + vector_t src_origin = { src_offset }; + auto src_pitch = pitch(region, {{ dst_img.pixel_size() }}); + + validate_common(q, deps); + validate_object(q, dst_img, dst_origin, region); + validate_object(q, src_mem, src_origin, src_pitch, region); + + auto hev = create<hard_event>( + q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, deps, + soft_copy_op(q, &dst_img, dst_origin, dst_pitch, + &src_mem, src_origin, src_pitch, + region)); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API void * +clEnqueueMapBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + cl_map_flags flags, size_t offset, size_t size, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev, cl_int *r_errcode) try { + auto &q = obj(d_q); + auto &mem = obj<buffer>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + vector_t region = { size, 1, 1 }; + vector_t obj_origin = { offset }; + auto obj_pitch = pitch(region, {{ 1 }}); + + validate_common(q, deps); + validate_object(q, mem, obj_origin, obj_pitch, region); + validate_map_flags(mem, flags); + + void *map = mem.resource(q).add_map(q, flags, blocking, obj_origin, region); + + auto hev = create<hard_event>(q, CL_COMMAND_MAP_BUFFER, deps); + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + ret_error(r_errcode, CL_SUCCESS); + return map; + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API void * +clEnqueueMapImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking, + cl_map_flags flags, + const size_t *p_origin, const size_t *p_region, + size_t *row_pitch, size_t *slice_pitch, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev, cl_int *r_errcode) try { + auto &q = obj(d_q); + auto &img = obj<image>(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + auto region = vector(p_region); + auto origin = vector(p_origin); + + validate_common(q, deps); + validate_object(q, img, origin, region); + validate_map_flags(img, flags); + + void *map = img.resource(q).add_map(q, flags, blocking, origin, region); + + auto hev = create<hard_event>(q, CL_COMMAND_MAP_IMAGE, deps); + if (blocking) + hev().wait_signalled(); + + ret_object(rd_ev, hev); + ret_error(r_errcode, CL_SUCCESS); + return map; + +} catch (error &e) { + ret_error(r_errcode, e); + return NULL; +} + +CLOVER_API cl_int +clEnqueueUnmapMemObject(cl_command_queue d_q, cl_mem d_mem, void *ptr, + cl_uint num_deps, const cl_event *d_deps, + cl_event *rd_ev) try { + auto &q = obj(d_q); + auto &mem = obj(d_mem); + auto deps = objs<wait_list_tag>(d_deps, num_deps); + + validate_common(q, deps); + + auto hev = create<hard_event>( + q, CL_COMMAND_UNMAP_MEM_OBJECT, deps, + [=, &q, &mem](event &) { + mem.resource(q).del_map(ptr); + }); + + ret_object(rd_ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueMigrateMemObjects(cl_command_queue command_queue, + 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) { + CLOVER_NOT_SUPPORTED_UNTIL("1.2"); + return CL_INVALID_VALUE; +} + +cl_int +clover::EnqueueSVMFree(cl_command_queue d_q, + 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, + cl_int cmd) try { + + if (bool(num_svm_pointers) != bool(svm_pointers)) + return CL_INVALID_VALUE; + + auto &q = obj(d_q); + bool can_emulate = q.device().has_system_svm(); + auto deps = objs<wait_list_tag>(event_wait_list, num_events_in_wait_list); + + validate_common(q, deps); + + std::vector<void *> svm_pointers_cpy(svm_pointers, + svm_pointers + num_svm_pointers); + if (!pfn_free_func) { + if (!can_emulate) { + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + } + pfn_free_func = [](cl_command_queue, cl_uint num_svm_pointers, + void *svm_pointers[], void *) { + for (void *p : range(svm_pointers, num_svm_pointers)) + free(p); + }; + } + + auto hev = create<hard_event>(q, cmd, deps, + [=](clover::event &) mutable { + pfn_free_func(d_q, num_svm_pointers, svm_pointers_cpy.data(), + user_data); + }); + + ret_object(event, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueSVMFree(cl_command_queue d_q, + 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) { + + return EnqueueSVMFree(d_q, num_svm_pointers, svm_pointers, + pfn_free_func, user_data, num_events_in_wait_list, + event_wait_list, event, CL_COMMAND_SVM_FREE); +} + +cl_int +clover::EnqueueSVMMemcpy(cl_command_queue d_q, + 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, + cl_int cmd) try { + + if (dst_ptr == nullptr || src_ptr == nullptr) + return CL_INVALID_VALUE; + + if (static_cast<size_t>(abs(reinterpret_cast<ptrdiff_t>(dst_ptr) - + reinterpret_cast<ptrdiff_t>(src_ptr))) < size) + return CL_MEM_COPY_OVERLAP; + + auto &q = obj(d_q); + bool can_emulate = q.device().has_system_svm(); + auto deps = objs<wait_list_tag>(event_wait_list, num_events_in_wait_list); + + validate_common(q, deps); + + if (can_emulate) { + auto hev = create<hard_event>(q, cmd, deps, + [=](clover::event &) { + memcpy(dst_ptr, src_ptr, size); + }); + + if (blocking_copy) + hev().wait(); + ret_object(event, hev); + return CL_SUCCESS; + } + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueSVMMemcpy(cl_command_queue d_q, + 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) { + + return EnqueueSVMMemcpy(d_q, blocking_copy, dst_ptr, src_ptr, + size, num_events_in_wait_list, event_wait_list, + event, CL_COMMAND_SVM_MEMCPY); +} + +cl_int +clover::EnqueueSVMMemFill(cl_command_queue d_q, + 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, + cl_int cmd) try { + + if (svm_ptr == nullptr || pattern == nullptr || + !util_is_power_of_two_nonzero(pattern_size) || + pattern_size > 128 || + !ptr_is_aligned(svm_ptr, pattern_size) || + size % pattern_size) + return CL_INVALID_VALUE; + + auto &q = obj(d_q); + bool can_emulate = q.device().has_system_svm(); + auto deps = objs<wait_list_tag>(event_wait_list, num_events_in_wait_list); + + validate_common(q, deps); + + if (can_emulate) { + auto hev = create<hard_event>(q, cmd, deps, + [=](clover::event &) { + void *ptr = svm_ptr; + for (size_t s = size; s; s -= pattern_size) { + memcpy(ptr, pattern, pattern_size); + ptr = static_cast<uint8_t*>(ptr) + pattern_size; + } + }); + + ret_object(event, hev); + return CL_SUCCESS; + } + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueSVMMemFill(cl_command_queue d_q, + 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) { + + return EnqueueSVMMemFill(d_q, svm_ptr, pattern, pattern_size, + size, num_events_in_wait_list, event_wait_list, + event, CL_COMMAND_SVM_MEMFILL); +} + +cl_int +clover::EnqueueSVMMap(cl_command_queue d_q, + cl_bool blocking_map, + cl_map_flags map_flags, + void *svm_ptr, + size_t size, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event, + cl_int cmd) try { + + if (svm_ptr == nullptr || size == 0) + return CL_INVALID_VALUE; + + auto &q = obj(d_q); + bool can_emulate = q.device().has_system_svm(); + auto deps = objs<wait_list_tag>(event_wait_list, num_events_in_wait_list); + + validate_common(q, deps); + + if (can_emulate) { + auto hev = create<hard_event>(q, cmd, deps, + [](clover::event &) { }); + + ret_object(event, hev); + return CL_SUCCESS; + } + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueSVMMap(cl_command_queue d_q, + cl_bool blocking_map, + cl_map_flags map_flags, + void *svm_ptr, + size_t size, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + + return EnqueueSVMMap(d_q, blocking_map, map_flags, svm_ptr, size, + num_events_in_wait_list, event_wait_list, event, + CL_COMMAND_SVM_MAP); +} + +cl_int +clover::EnqueueSVMUnmap(cl_command_queue d_q, + void *svm_ptr, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event, + cl_int cmd) try { + + if (svm_ptr == nullptr) + return CL_INVALID_VALUE; + + auto &q = obj(d_q); + bool can_emulate = q.device().has_system_svm(); + auto deps = objs<wait_list_tag>(event_wait_list, num_events_in_wait_list); + + validate_common(q, deps); + + if (can_emulate) { + auto hev = create<hard_event>(q, cmd, deps, + [](clover::event &) { }); + + ret_object(event, hev); + return CL_SUCCESS; + } + + CLOVER_NOT_SUPPORTED_UNTIL("2.0"); + return CL_INVALID_VALUE; + +} catch (error &e) { + return e.get(); +} + +CLOVER_API cl_int +clEnqueueSVMUnmap(cl_command_queue d_q, + void *svm_ptr, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + + return EnqueueSVMUnmap(d_q, svm_ptr, num_events_in_wait_list, + event_wait_list, event, CL_COMMAND_SVM_UNMAP); +} + +CLOVER_API cl_int +clEnqueueSVMMigrateMem(cl_command_queue d_q, + cl_uint num_svm_pointers, + const void **svm_pointers, + const size_t *sizes, + const cl_mem_migration_flags flags, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + CLOVER_NOT_SUPPORTED_UNTIL("2.1"); + return CL_INVALID_VALUE; +} diff --git a/src/gallium/frontends/clover/api/util.hpp b/src/gallium/frontends/clover/api/util.hpp new file mode 100644 index 00000000000..66bd12597c6 --- /dev/null +++ b/src/gallium/frontends/clover/api/util.hpp @@ -0,0 +1,84 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +// OTHER DEALINGS IN THE SOFTWARE. +// + +#ifndef CLOVER_API_UTIL_HPP +#define CLOVER_API_UTIL_HPP + +#include <cassert> +#include <iostream> + +#include "core/error.hpp" +#include "core/property.hpp" +#include "util/algorithm.hpp" + +#ifdef HAVE_CLOVER_ICD +#define CLOVER_API +#define CLOVER_ICD_API PUBLIC +#else +#define CLOVER_API PUBLIC +#define CLOVER_ICD_API PUBLIC +#endif + +#define CLOVER_NOT_SUPPORTED_UNTIL(version) \ + do { \ + std::cerr << "CL user error: " << __func__ \ + << "() requires OpenCL version " << (version) \ + << " or greater." << std::endl; \ + } while (0) + +namespace clover { + /// + /// Return an error code in \a p if non-zero. + /// + inline void + ret_error(cl_int *p, const clover::error &e) { + if (p) + *p = e.get(); + } + + /// + /// Return a clover object in \a p if non-zero incrementing the + /// reference count of the object. + /// + template<typename T> + void + ret_object(typename T::descriptor_type **p, + const intrusive_ref<T> &v) { + if (p) { + v().retain(); + *p = desc(v()); + } + } + + /// + /// Return an API object from an intrusive reference to a Clover object, + /// incrementing the reference count of the object. + /// + template<typename T> + typename T::descriptor_type * + ret_object(const intrusive_ref<T> &v) { + v().retain(); + return desc(v()); + } +} + +#endif |