indent_style = space
indent_size = 2
-
[*.ps1]
indent_style = space
indent_size = 2
+
+[*.rs]
+indent_style = space
+indent_size = 4
)
endif
+opencl_headers = files(
+ 'CL/cl.h',
+ 'CL/cl.hpp',
+ 'CL/cl2.hpp',
+ 'CL/cl_d3d10.h',
+ 'CL/cl_d3d11.h',
+ 'CL/cl_dx9_media_sharing.h',
+ 'CL/cl_dx9_media_sharing_intel.h',
+ 'CL/cl_egl.h',
+ 'CL/cl_ext.h',
+ 'CL/cl_ext_intel.h',
+ 'CL/cl_gl.h',
+ 'CL/cl_gl_ext.h',
+ 'CL/cl_icd.h',
+ 'CL/cl_platform.h',
+ 'CL/cl_va_api_media_sharing_intel.h',
+ 'CL/cl_version.h',
+ 'CL/opencl.h',
+)
# Only install the headers if we are building a stand alone implementation and
# not an ICD enabled implementation
if with_gallium_opencl and not with_opencl_icd
install_headers(
- 'CL/cl.h',
- 'CL/cl.hpp',
- 'CL/cl2.hpp',
- 'CL/cl_d3d10.h',
- 'CL/cl_d3d11.h',
- 'CL/cl_dx9_media_sharing.h',
- 'CL/cl_dx9_media_sharing_intel.h',
- 'CL/cl_egl.h',
- 'CL/cl_ext.h',
- 'CL/cl_ext_intel.h',
- 'CL/cl_gl.h',
- 'CL/cl_gl_ext.h',
- 'CL/cl_icd.h',
- 'CL/cl_platform.h',
- 'CL/cl_va_api_media_sharing_intel.h',
- 'CL/cl_version.h',
- 'CL/opencl.h',
+ opencl_headers,
subdir: 'CL'
)
endif
).stdout(),
license : 'MIT',
meson_version : '>= 0.53',
- default_options : ['buildtype=debugoptimized', 'b_ndebug=if-release', 'c_std=c11', 'cpp_std=c++17']
+ default_options : ['buildtype=debugoptimized', 'b_ndebug=if-release', 'c_std=c11', 'cpp_std=c++17', 'rust_std=2021']
)
# In recent versions, meson can inject some extra arguments to get richer
with_opencl_icd = false
endif
+with_gallium_rusticl = get_option('gallium-rusticl')
+if with_gallium_rusticl
+ if not with_gallium
+ error('rusticl requires at least one gallium driver.')
+ endif
+
+ if meson.version().version_compare('< 0.61.4')
+ error('rusticl requires meson 0.61.4 or newer')
+ endif
+
+ add_languages('rust', required: true)
+
+ with_clc = true
+endif
+
dep_clc = null_dep
if with_libclc
dep_clc = dependency('libclc')
description : 'build gallium "clover" OpenCL frontend.',
)
option(
+ 'gallium-rusticl',
+ type : 'boolean',
+ value : false,
+ description : 'build gallium "rusticl" OpenCL frontend.',
+)
+option(
'gallium-windows-dll-name',
type : 'string',
value : 'libgallium_wgl',
--- /dev/null
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::device::get_devs_for_type;
+use crate::api::icd::*;
+use crate::api::platform::*;
+use crate::api::types::*;
+use crate::api::util::*;
+use crate::core::context::*;
+
+use self::mesa_rust_util::properties::Properties;
+use self::rusticl_opencl_gen::*;
+
+use std::collections::HashSet;
+use std::iter::FromIterator;
+use std::slice;
+use std::sync::Arc;
+
+impl CLInfo<cl_context_info> for cl_context {
+ fn query(&self, q: cl_context_info) -> CLResult<Vec<u8>> {
+ let ctx = self.get_ref()?;
+ Ok(match q {
+ CL_CONTEXT_DEVICES => {
+ cl_prop::<&Vec<cl_device_id>>(
+ &ctx.devs
+ .iter()
+ .map(|d| {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ cl_device_id::from_ptr(Arc::as_ptr(d))
+ })
+ .collect(),
+ )
+ }
+ CL_CONTEXT_NUM_DEVICES => cl_prop::<cl_uint>(ctx.devs.len() as u32),
+ CL_CONTEXT_PROPERTIES => cl_prop::<&Vec<cl_context_properties>>(&ctx.properties),
+ CL_CONTEXT_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+pub fn create_context(
+ properties: *const cl_context_properties,
+ num_devices: cl_uint,
+ devices: *const cl_device_id,
+ pfn_notify: Option<CreateContextCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<cl_context> {
+ check_cb(&pfn_notify, user_data)?;
+
+ // CL_INVALID_VALUE if devices is NULL.
+ if devices.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_VALUE if num_devices is equal to zero.
+ if num_devices == 0 {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_PROPERTY [...] if the same property name is specified more than once.
+ let props = Properties::from_ptr(properties).ok_or(CL_INVALID_PROPERTY)?;
+ for p in props.props {
+ match p.0 as u32 {
+ // CL_INVALID_PLATFORM [...] if platform value specified in properties is not a valid platform.
+ CL_CONTEXT_PLATFORM => {
+ (p.1 as cl_platform_id).get_ref()?;
+ }
+ CL_CONTEXT_INTEROP_USER_SYNC => {
+ check_cl_bool(p.1).ok_or(CL_INVALID_PROPERTY)?;
+ }
+ // CL_INVALID_PROPERTY if context property name in properties is not a supported property name
+ _ => return Err(CL_INVALID_PROPERTY),
+ }
+ }
+
+ // Duplicate devices specified in devices are ignored.
+ let set: HashSet<_> =
+ HashSet::from_iter(unsafe { slice::from_raw_parts(devices, num_devices as usize) }.iter());
+ let devs: Result<_, _> = set.into_iter().map(cl_device_id::get_arc).collect();
+
+ Ok(cl_context::from_arc(Context::new(
+ devs?,
+ Properties::from_ptr_raw(properties),
+ )))
+}
+
+pub fn create_context_from_type(
+ properties: *const cl_context_properties,
+ device_type: cl_device_type,
+ pfn_notify: Option<CreateContextCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<cl_context> {
+ // CL_INVALID_DEVICE_TYPE if device_type is not a valid value.
+ check_cl_device_type(device_type)?;
+
+ let devs: Vec<_> = get_devs_for_type(device_type)
+ .iter()
+ .map(|d| cl_device_id::from_ptr(Arc::as_ptr(d)))
+ .collect();
+
+ // CL_DEVICE_NOT_FOUND if no devices that match device_type and property values specified in properties were found.
+ if devs.is_empty() {
+ return Err(CL_DEVICE_NOT_FOUND);
+ }
+
+ // errors are essentially the same and we will always pass in a valid
+ // device list, so that's fine as well.
+ create_context(
+ properties,
+ devs.len() as u32,
+ devs.as_ptr(),
+ pfn_notify,
+ user_data,
+ )
+}
--- /dev/null
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::platform::*;
+use crate::api::util::*;
+use crate::core::device::*;
+
+use self::mesa_rust_util::ptr::*;
+use self::rusticl_opencl_gen::*;
+
+use std::cmp::min;
+use std::mem::size_of;
+use std::ptr;
+use std::sync::Arc;
+use std::sync::Once;
+
+impl CLInfo<cl_device_info> for cl_device_id {
+ fn query(&self, q: cl_device_info) -> CLResult<Vec<u8>> {
+ let dev = self.get_ref()?;
+ Ok(match q {
+ CL_DEVICE_ADDRESS_BITS => cl_prop::<cl_uint>(dev.address_bits()),
+ CL_DEVICE_ATOMIC_FENCE_CAPABILITIES => cl_prop::<cl_device_atomic_capabilities>(0),
+ CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES => cl_prop::<cl_device_atomic_capabilities>(0),
+ CL_DEVICE_AVAILABLE => cl_prop::<bool>(true),
+ CL_DEVICE_BUILT_IN_KERNELS => cl_prop::<&str>(""),
+ CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION => cl_prop::<Vec<cl_name_version>>(Vec::new()),
+ CL_DEVICE_COMPILER_AVAILABLE => cl_prop::<bool>(true),
+ CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES => {
+ cl_prop::<cl_device_device_enqueue_capabilities>(0)
+ }
+ CL_DEVICE_DOUBLE_FP_CONFIG => cl_prop::<cl_device_fp_config>(0),
+ CL_DEVICE_ENDIAN_LITTLE => cl_prop::<bool>(dev.little_endian()),
+ CL_DEVICE_ERROR_CORRECTION_SUPPORT => cl_prop::<bool>(false),
+ CL_DEVICE_EXECUTION_CAPABILITIES => {
+ cl_prop::<cl_device_exec_capabilities>(CL_EXEC_KERNEL.into())
+ }
+ CL_DEVICE_EXTENSIONS => cl_prop::<&str>(&dev.extension_string),
+ CL_DEVICE_EXTENSIONS_WITH_VERSION => cl_prop::<&Vec<cl_name_version>>(&dev.extensions),
+ CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT => cl_prop::<bool>(false),
+ CL_DEVICE_GLOBAL_MEM_CACHE_TYPE => cl_prop::<cl_device_mem_cache_type>(CL_NONE),
+ CL_DEVICE_GLOBAL_MEM_CACHE_SIZE => cl_prop::<cl_ulong>(0),
+ CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_GLOBAL_MEM_SIZE => cl_prop::<cl_ulong>(dev.global_mem_size()),
+ CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE => cl_prop::<usize>(0),
+ CL_DEVICE_HALF_FP_CONFIG => cl_prop::<cl_device_fp_config>(0),
+ CL_DEVICE_HOST_UNIFIED_MEMORY => cl_prop::<bool>(dev.unified_memory()),
+ CL_DEVICE_IL_VERSION => cl_prop::<&str>(""),
+ CL_DEVICE_ILS_WITH_VERSION => cl_prop::<Vec<cl_name_version>>(Vec::new()),
+ CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT => {
+ cl_prop::<cl_uint>(dev.image_base_address_alignment())
+ }
+ CL_DEVICE_IMAGE_MAX_ARRAY_SIZE => cl_prop::<usize>(dev.image_array_size()),
+ CL_DEVICE_IMAGE_MAX_BUFFER_SIZE => cl_prop::<usize>(dev.image_buffer_size()),
+ CL_DEVICE_IMAGE_PITCH_ALIGNMENT => cl_prop::<cl_uint>(0),
+ CL_DEVICE_IMAGE_SUPPORT => cl_prop::<bool>(dev.image_supported()),
+ CL_DEVICE_IMAGE2D_MAX_HEIGHT => cl_prop::<usize>(dev.image_2d_size()),
+ CL_DEVICE_IMAGE2D_MAX_WIDTH => cl_prop::<usize>(dev.image_2d_size()),
+ CL_DEVICE_IMAGE3D_MAX_HEIGHT => cl_prop::<usize>(dev.image_3d_size()),
+ CL_DEVICE_IMAGE3D_MAX_WIDTH => cl_prop::<usize>(dev.image_3d_size()),
+ CL_DEVICE_IMAGE3D_MAX_DEPTH => cl_prop::<usize>(dev.image_3d_size()),
+ CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED => cl_prop::<&str>("v0000-01-01-00"),
+ CL_DEVICE_LINKER_AVAILABLE => cl_prop::<bool>(true),
+ CL_DEVICE_LOCAL_MEM_SIZE => cl_prop::<cl_ulong>(dev.local_mem_size()),
+ // TODO add query for CL_LOCAL vs CL_GLOBAL
+ CL_DEVICE_LOCAL_MEM_TYPE => cl_prop::<cl_device_local_mem_type>(CL_GLOBAL),
+ CL_DEVICE_MAX_CLOCK_FREQUENCY => cl_prop::<cl_uint>(dev.max_clock_freq()),
+ CL_DEVICE_MAX_COMPUTE_UNITS => cl_prop::<cl_uint>(dev.max_compute_units()),
+ // TODO atm implemented as mem_const
+ CL_DEVICE_MAX_CONSTANT_ARGS => cl_prop::<cl_uint>(1024),
+ CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE => cl_prop::<cl_ulong>(dev.const_max_size()),
+ CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE => cl_prop::<usize>(0),
+ CL_DEVICE_MAX_MEM_ALLOC_SIZE => cl_prop::<cl_ulong>(dev.max_mem_alloc()),
+ CL_DEVICE_MAX_NUM_SUB_GROUPS => cl_prop::<cl_uint>(0),
+ CL_DEVICE_MAX_ON_DEVICE_EVENTS => cl_prop::<cl_uint>(0),
+ CL_DEVICE_MAX_ON_DEVICE_QUEUES => cl_prop::<cl_uint>(0),
+ CL_DEVICE_MAX_PARAMETER_SIZE => cl_prop::<usize>(dev.param_max_size()),
+ CL_DEVICE_MAX_PIPE_ARGS => cl_prop::<cl_uint>(0),
+ CL_DEVICE_MAX_READ_IMAGE_ARGS => cl_prop::<cl_uint>(dev.image_read_count()),
+ CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS => cl_prop::<cl_uint>(0),
+ CL_DEVICE_MAX_SAMPLERS => cl_prop::<cl_uint>(dev.max_samplers()),
+ CL_DEVICE_MAX_WORK_GROUP_SIZE => cl_prop::<usize>(dev.max_threads_per_block()),
+ CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS => cl_prop::<cl_uint>(dev.max_grid_dimensions()),
+ CL_DEVICE_MAX_WORK_ITEM_SIZES => cl_prop::<Vec<usize>>(dev.max_block_sizes()),
+ CL_DEVICE_MAX_WRITE_IMAGE_ARGS => cl_prop::<cl_uint>(dev.image_write_count()),
+ // TODO proper retrival from devices
+ CL_DEVICE_MEM_BASE_ADDR_ALIGN => cl_prop::<cl_uint>(0x1000),
+ CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE => {
+ cl_prop::<cl_uint>(size_of::<cl_ulong16>() as cl_uint)
+ }
+ CL_DEVICE_NAME => cl_prop(dev.screen().name()),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR => cl_prop::<cl_uint>(1),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF => cl_prop::<cl_uint>(0),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_INT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG => cl_prop::<cl_uint>(1),
+ CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT => cl_prop::<bool>(false),
+ CL_DEVICE_NUMERIC_VERSION => cl_prop::<cl_version>(dev.cl_version as cl_version),
+ // TODO subdevice support
+ CL_DEVICE_PARENT_DEVICE => cl_prop::<cl_device_id>(cl_device_id::from_ptr(ptr::null())),
+ CL_DEVICE_PARTITION_AFFINITY_DOMAIN => cl_prop::<cl_device_affinity_domain>(0),
+ CL_DEVICE_PARTITION_MAX_SUB_DEVICES => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PARTITION_PROPERTIES => cl_prop::<Vec<cl_device_partition_property>>(vec![0]),
+ CL_DEVICE_PARTITION_TYPE => cl_prop::<Vec<cl_device_partition_property>>(Vec::new()),
+ CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PIPE_MAX_PACKET_SIZE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PIPE_SUPPORT => cl_prop::<bool>(false),
+ CL_DEVICE_PLATFORM => cl_prop::<cl_platform_id>(get_platform()),
+ CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PREFERRED_INTEROP_USER_SYNC => cl_prop::<bool>(true),
+ CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR => cl_prop::<cl_uint>(1),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF => cl_prop::<cl_uint>(0),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG => cl_prop::<cl_uint>(1),
+ CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => cl_prop::<usize>(1),
+ // TODO
+ CL_DEVICE_PRINTF_BUFFER_SIZE => cl_prop::<usize>(0),
+ // TODO
+ CL_DEVICE_PROFILING_TIMER_RESOLUTION => cl_prop::<usize>(0),
+ CL_DEVICE_OPENCL_C_FEATURES => cl_prop::<Vec<cl_name_version>>(Vec::new()),
+ CL_DEVICE_OPENCL_C_VERSION => {
+ cl_prop::<String>(format!("OpenCL C {} ", dev.clc_version.api_str()))
+ }
+ CL_DEVICE_OPENCL_C_ALL_VERSIONS => cl_prop::<&Vec<cl_name_version>>(&dev.clc_versions),
+ CL_DEVICE_PROFILE => cl_prop(if dev.embedded {
+ "EMBEDDED_PROFILE"
+ } else {
+ "FULL_PROFILE"
+ }),
+ CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE => cl_prop::<cl_uint>(0),
+ CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES => cl_prop::<cl_command_queue_properties>(0),
+ CL_DEVICE_QUEUE_ON_HOST_PROPERTIES => {
+ cl_prop::<cl_command_queue_properties>(CL_QUEUE_PROFILING_ENABLE.into())
+ }
+ CL_DEVICE_REFERENCE_COUNT => cl_prop::<cl_uint>(1),
+ CL_DEVICE_SINGLE_FP_CONFIG => cl_prop::<cl_device_fp_config>(
+ (CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN) as cl_device_fp_config,
+ ),
+ CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS => cl_prop::<bool>(false),
+ CL_DEVICE_SVM_CAPABILITIES => cl_prop::<cl_device_svm_capabilities>(0),
+ CL_DEVICE_TYPE => cl_prop::<cl_device_type>(dev.device_type()),
+ CL_DEVICE_VENDOR => cl_prop(dev.screen().device_vendor()),
+ CL_DEVICE_VENDOR_ID => cl_prop::<cl_uint>(dev.vendor_id()),
+ CL_DEVICE_VERSION => cl_prop::<String>(format!("OpenCL {}", dev.cl_version.api_str())),
+ CL_DRIVER_VERSION => cl_prop("0.1"),
+ CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT => cl_prop::<bool>(false),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ // CL_INVALID_VALUE [...] if param_name is a value that is available as an extension and the corresponding extension is not supported by the device.
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+// TODO replace with const new container
+static mut DEVICES: Vec<Arc<Device>> = Vec::new();
+static INIT: Once = Once::new();
+
+fn load_devices() {
+ Device::all()
+ .into_iter()
+ .for_each(|d| unsafe { DEVICES.push(d) });
+}
+
+fn devs() -> &'static Vec<Arc<Device>> {
+ INIT.call_once(load_devices);
+ unsafe { &DEVICES }
+}
+
+pub fn get_devs_for_type(device_type: cl_device_type) -> Vec<&'static Arc<Device>> {
+ devs()
+ .iter()
+ .filter(|d| device_type & d.device_type() != 0)
+ .collect()
+}
+
+pub fn get_device_ids(
+ platform: cl_platform_id,
+ device_type: cl_device_type,
+ num_entries: cl_uint,
+ devices: *mut cl_device_id,
+ num_devices: *mut cl_uint,
+) -> CLResult<()> {
+ // CL_INVALID_PLATFORM if platform is not a valid platform.
+ platform.get_ref()?;
+
+ // CL_INVALID_DEVICE_TYPE if device_type is not a valid value.
+ check_cl_device_type(device_type)?;
+
+ // CL_INVALID_VALUE if num_entries is equal to zero and devices is not NULL
+ if num_entries == 0 && !devices.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_VALUE [...] if both num_devices and devices are NULL.
+ if num_devices.is_null() && devices.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let devs = get_devs_for_type(device_type);
+ // CL_DEVICE_NOT_FOUND if no OpenCL devices that matched device_type were found
+ if devs.is_empty() {
+ return Err(CL_DEVICE_NOT_FOUND);
+ }
+
+ // num_devices returns the number of OpenCL devices available that match device_type. If
+ // num_devices is NULL, this argument is ignored.
+ num_devices.write_checked(devs.len() as cl_uint);
+
+ if !devices.is_null() {
+ let n = min(num_entries as usize, devs.len());
+
+ #[allow(clippy::needless_range_loop)]
+ for i in 0..n {
+ unsafe {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ *devices.add(i) = cl_device_id::from_ptr(Arc::as_ptr(devs[i]));
+ }
+ }
+ }
+
+ Ok(())
+}
--- /dev/null
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::util::*;
+use crate::core::event::*;
+use crate::core::queue::*;
+
+use self::rusticl_opencl_gen::*;
+
+use std::ptr;
+use std::sync::Arc;
+
+impl CLInfo<cl_event_info> for cl_event {
+ fn query(&self, q: cl_event_info) -> CLResult<Vec<u8>> {
+ let event = self.get_ref()?;
+ Ok(match q {
+ CL_EVENT_COMMAND_EXECUTION_STATUS => cl_prop::<cl_int>(event.status()),
+ CL_EVENT_CONTEXT => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&event.context);
+ cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+ }
+ CL_EVENT_COMMAND_QUEUE => {
+ let ptr = match event.queue.as_ref() {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ Some(queue) => Arc::as_ptr(queue),
+ None => ptr::null_mut(),
+ };
+ cl_prop::<cl_command_queue>(cl_command_queue::from_ptr(ptr))
+ }
+ CL_EVENT_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ CL_EVENT_COMMAND_TYPE => cl_prop::<cl_command_type>(event.cmd_type),
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+pub fn create_user_event(context: cl_context) -> CLResult<cl_event> {
+ let c = context.get_arc()?;
+ Ok(cl_event::from_arc(Event::new_user(c)))
+}
+
+pub fn create_and_queue(
+ q: Arc<Queue>,
+ cmd_type: cl_command_type,
+ deps: Vec<Arc<Event>>,
+ event: *mut cl_event,
+ block: bool,
+ work: EventSig,
+) -> CLResult<()> {
+ let e = Event::new(&q, cmd_type, deps, work);
+ cl_event::leak_ref(event, &e);
+ q.queue(&e);
+ if block {
+ q.flush(true)?;
+ }
+ Ok(())
+}
--- /dev/null
+#![allow(non_snake_case)]
+
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::context::*;
+use crate::api::device::*;
+use crate::api::event::*;
+use crate::api::kernel::*;
+use crate::api::memory::*;
+use crate::api::platform::*;
+use crate::api::program::*;
+use crate::api::queue::*;
+use crate::api::types::*;
+use crate::api::util::*;
+
+use self::mesa_rust_util::ptr::*;
+use self::rusticl_opencl_gen::*;
+
+use std::ffi::CStr;
+use std::ptr;
+use std::sync::Arc;
+
+pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
+ clGetPlatformIDs: Some(cl_get_platform_ids),
+ clGetPlatformInfo: Some(cl_get_platform_info),
+ clGetDeviceIDs: Some(cl_get_device_ids),
+ clGetDeviceInfo: Some(cl_get_device_info),
+ clCreateContext: Some(cl_create_context),
+ clCreateContextFromType: Some(cl_create_context_from_type),
+ clRetainContext: Some(cl_retain_context),
+ clReleaseContext: Some(cl_release_context),
+ clGetContextInfo: Some(cl_get_context_info),
+ clCreateCommandQueue: Some(cl_create_command_queue),
+ clRetainCommandQueue: Some(cl_retain_command_queue),
+ clReleaseCommandQueue: Some(cl_release_command_queue),
+ clGetCommandQueueInfo: Some(cl_get_command_queue_info),
+ clSetCommandQueueProperty: None,
+ clCreateBuffer: Some(cl_create_buffer),
+ clCreateImage2D: Some(cl_create_image_2d),
+ clCreateImage3D: Some(cl_create_image_3d),
+ clRetainMemObject: Some(cl_retain_mem_object),
+ clReleaseMemObject: Some(cl_release_mem_object),
+ clGetSupportedImageFormats: Some(cl_get_supported_image_formats),
+ clGetMemObjectInfo: Some(cl_get_mem_object_info),
+ clGetImageInfo: Some(cl_get_image_info),
+ clCreateSampler: Some(cl_create_sampler),
+ clRetainSampler: Some(cl_retain_sampler),
+ clReleaseSampler: Some(cl_release_sampler),
+ clGetSamplerInfo: Some(cl_get_sampler_info),
+ clCreateProgramWithSource: Some(cl_create_program_with_source),
+ clCreateProgramWithBinary: None,
+ clRetainProgram: Some(cl_retain_program),
+ clReleaseProgram: Some(cl_release_program),
+ clBuildProgram: Some(cl_build_program),
+ clUnloadCompiler: None,
+ clGetProgramInfo: Some(cl_get_program_info),
+ clGetProgramBuildInfo: Some(cl_get_program_build_info),
+ clCreateKernel: Some(cl_create_kernel),
+ clCreateKernelsInProgram: Some(cl_create_kernels_in_program),
+ clRetainKernel: None,
+ clReleaseKernel: None,
+ clSetKernelArg: None,
+ clGetKernelInfo: None,
+ clGetKernelWorkGroupInfo: None,
+ clWaitForEvents: Some(cl_wait_for_events),
+ clGetEventInfo: Some(cl_get_event_info),
+ clRetainEvent: None,
+ clReleaseEvent: Some(cl_release_event),
+ clGetEventProfilingInfo: Some(cl_get_event_profiling_info),
+ clFlush: None,
+ clFinish: Some(cl_finish),
+ clEnqueueReadBuffer: Some(cl_enqueue_read_buffer),
+ clEnqueueWriteBuffer: Some(cl_enqueue_write_buffer),
+ clEnqueueCopyBuffer: Some(cl_enqueue_copy_buffer),
+ clEnqueueReadImage: Some(cl_enqueue_read_image),
+ clEnqueueWriteImage: Some(cl_enqueue_write_image),
+ clEnqueueCopyImage: Some(cl_enqueue_copy_image),
+ clEnqueueCopyImageToBuffer: Some(cl_enqueue_copy_image_to_buffer),
+ clEnqueueCopyBufferToImage: Some(cl_enqueue_copy_buffer_to_image),
+ clEnqueueMapBuffer: Some(cl_enqueue_map_buffer),
+ clEnqueueMapImage: Some(cl_enqueue_map_image),
+ clEnqueueUnmapMemObject: Some(cl_enqueue_unmap_mem_object),
+ clEnqueueNDRangeKernel: None,
+ clEnqueueTask: None,
+ clEnqueueNativeKernel: None,
+ clEnqueueMarker: None,
+ clEnqueueWaitForEvents: None,
+ clEnqueueBarrier: None,
+ clGetExtensionFunctionAddress: Some(cl_get_extension_function_address),
+ clCreateFromGLBuffer: None,
+ clCreateFromGLTexture2D: None,
+ clCreateFromGLTexture3D: None,
+ clCreateFromGLRenderbuffer: None,
+ clGetGLObjectInfo: None,
+ clGetGLTextureInfo: None,
+ clEnqueueAcquireGLObjects: None,
+ clEnqueueReleaseGLObjects: None,
+ clGetGLContextInfoKHR: None,
+ clGetDeviceIDsFromD3D10KHR: ptr::null_mut(),
+ clCreateFromD3D10BufferKHR: ptr::null_mut(),
+ clCreateFromD3D10Texture2DKHR: ptr::null_mut(),
+ clCreateFromD3D10Texture3DKHR: ptr::null_mut(),
+ clEnqueueAcquireD3D10ObjectsKHR: ptr::null_mut(),
+ clEnqueueReleaseD3D10ObjectsKHR: ptr::null_mut(),
+ clSetEventCallback: Some(cl_set_event_callback),
+ clCreateSubBuffer: Some(cl_create_sub_buffer),
+ clSetMemObjectDestructorCallback: Some(cl_set_mem_object_destructor_callback),
+ clCreateUserEvent: Some(cl_create_user_event),
+ clSetUserEventStatus: Some(cl_set_user_event_status),
+ clEnqueueReadBufferRect: Some(cl_enqueue_read_buffer_rect),
+ clEnqueueWriteBufferRect: Some(cl_enqueue_write_buffer_rect),
+ clEnqueueCopyBufferRect: Some(cl_enqueue_copy_buffer_rect),
+ clCreateSubDevicesEXT: None,
+ clRetainDeviceEXT: None,
+ clReleaseDeviceEXT: None,
+ clCreateEventFromGLsyncKHR: None,
+ clCreateSubDevices: None,
+ clRetainDevice: None,
+ clReleaseDevice: None,
+ clCreateImage: Some(cl_create_image),
+ clCreateProgramWithBuiltInKernels: None,
+ clCompileProgram: Some(cl_compile_program),
+ clLinkProgram: Some(cl_link_program),
+ clUnloadPlatformCompiler: Some(cl_unload_platform_compiler),
+ clGetKernelArgInfo: None,
+ clEnqueueFillBuffer: None,
+ clEnqueueFillImage: Some(cl_enqueue_fill_image),
+ clEnqueueMigrateMemObjects: None,
+ clEnqueueMarkerWithWaitList: None,
+ clEnqueueBarrierWithWaitList: None,
+ clGetExtensionFunctionAddressForPlatform: None,
+ clCreateFromGLTexture: None,
+ clGetDeviceIDsFromD3D11KHR: ptr::null_mut(),
+ clCreateFromD3D11BufferKHR: ptr::null_mut(),
+ clCreateFromD3D11Texture2DKHR: ptr::null_mut(),
+ clCreateFromD3D11Texture3DKHR: ptr::null_mut(),
+ clCreateFromDX9MediaSurfaceKHR: ptr::null_mut(),
+ clEnqueueAcquireD3D11ObjectsKHR: ptr::null_mut(),
+ clEnqueueReleaseD3D11ObjectsKHR: ptr::null_mut(),
+ clGetDeviceIDsFromDX9MediaAdapterKHR: ptr::null_mut(),
+ clEnqueueAcquireDX9MediaSurfacesKHR: ptr::null_mut(),
+ clEnqueueReleaseDX9MediaSurfacesKHR: ptr::null_mut(),
+ clCreateFromEGLImageKHR: None,
+ clEnqueueAcquireEGLObjectsKHR: None,
+ clEnqueueReleaseEGLObjectsKHR: None,
+ clCreateEventFromEGLSyncKHR: None,
+ clCreateCommandQueueWithProperties: None,
+ clCreatePipe: None,
+ clGetPipeInfo: None,
+ clSVMAlloc: None,
+ clSVMFree: None,
+ clEnqueueSVMFree: None,
+ clEnqueueSVMMemcpy: None,
+ clEnqueueSVMMemFill: None,
+ clEnqueueSVMMap: None,
+ clEnqueueSVMUnmap: None,
+ clCreateSamplerWithProperties: None,
+ clSetKernelArgSVMPointer: None,
+ clSetKernelExecInfo: None,
+ clGetKernelSubGroupInfoKHR: None,
+ clCloneKernel: None,
+ clCreateProgramWithIL: None,
+ clEnqueueSVMMigrateMem: None,
+ clGetDeviceAndHostTimer: None,
+ clGetHostTimer: None,
+ clGetKernelSubGroupInfo: None,
+ clSetDefaultDeviceCommandQueue: None,
+ clSetProgramReleaseCallback: None,
+ clSetProgramSpecializationConstant: None,
+ clCreateBufferWithProperties: None,
+ clCreateImageWithProperties: None,
+ clSetContextDestructorCallback: None,
+};
+
+pub type CLError = cl_int;
+pub type CLResult<T> = Result<T, CLError>;
+
+#[repr(C)]
+pub struct CLObjectBase<const ERR: i32> {
+ dispatch: &'static cl_icd_dispatch,
+ type_err: i32,
+}
+
+impl<const ERR: i32> Default for CLObjectBase<ERR> {
+ fn default() -> Self {
+ Self::new()
+ }
+}
+
+impl<const ERR: i32> CLObjectBase<ERR> {
+ pub fn new() -> Self {
+ Self {
+ dispatch: &DISPATCH,
+ type_err: ERR,
+ }
+ }
+
+ pub fn check_ptr(ptr: *const Self) -> CLResult<()> {
+ if ptr.is_null() {
+ return Err(ERR);
+ }
+
+ unsafe {
+ if !::std::ptr::eq((*ptr).dispatch, &DISPATCH) {
+ return Err(ERR);
+ }
+
+ if (*ptr).type_err != ERR {
+ return Err(ERR);
+ }
+
+ Ok(())
+ }
+ }
+}
+
+pub trait ReferenceCountedAPIPointer<T, const ERR: i32> {
+ fn get_ptr(&self) -> CLResult<*const T>;
+
+ // TODO: I can't find a trait that would let me say T: pointer so that
+ // I can do the cast in the main trait implementation. So we need to
+ // implement that as part of the macro where we know the real type.
+ fn from_ptr(ptr: *const T) -> Self;
+
+ fn leak_ref(ptr: *mut Self, r: &std::sync::Arc<T>)
+ where
+ Self: Sized,
+ {
+ if !ptr.is_null() {
+ unsafe {
+ ptr.write(Self::from_arc(r.clone()));
+ }
+ }
+ }
+
+ fn get_ref(&self) -> CLResult<&'static T> {
+ unsafe { Ok(self.get_ptr()?.as_ref().unwrap()) }
+ }
+
+ fn get_arc(&self) -> CLResult<Arc<T>> {
+ unsafe {
+ let ptr = self.get_ptr()?;
+ Arc::increment_strong_count(ptr);
+ Ok(Arc::from_raw(ptr))
+ }
+ }
+
+ fn from_arc(arc: Arc<T>) -> Self
+ where
+ Self: Sized,
+ {
+ Self::from_ptr(Arc::into_raw(arc))
+ }
+
+ fn get_arc_vec_from_arr(objs: *const Self, count: u32) -> CLResult<Vec<Arc<T>>>
+ where
+ Self: Sized,
+ {
+ // CL spec requires validation for obj arrays, both values have to make sense
+ if objs.is_null() && count > 0 || !objs.is_null() && count == 0 {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let mut res = Vec::new();
+ if objs.is_null() || count == 0 {
+ return Ok(res);
+ }
+
+ for i in 0..count as usize {
+ unsafe {
+ res.push((*objs.add(i)).get_arc()?);
+ }
+ }
+ Ok(res)
+ }
+
+ fn retain(&self) -> CLResult<()> {
+ unsafe {
+ Arc::increment_strong_count(self.get_ptr()?);
+ Ok(())
+ }
+ }
+
+ fn release(&self) -> CLResult<Arc<T>> {
+ unsafe { Ok(Arc::from_raw(self.get_ptr()?)) }
+ }
+
+ fn refcnt(&self) -> CLResult<u32> {
+ Ok((Arc::strong_count(&self.get_arc()?) - 1) as u32)
+ }
+}
+
+#[macro_export]
+macro_rules! impl_cl_type_trait {
+ ($cl: ident, $t: ty, $err: ident) => {
+ impl $crate::api::icd::ReferenceCountedAPIPointer<$t, $err> for $cl {
+ fn get_ptr(&self) -> CLResult<*const $t> {
+ type Base = $crate::api::icd::CLObjectBase<$err>;
+ Base::check_ptr(self.cast())?;
+
+ // Now that we've verified the object, it should be safe to
+ // dereference it. As one more double check, make sure that
+ // the CLObjectBase is at the start of the object
+ let obj_ptr: *const $t = self.cast();
+ unsafe {
+ let base_ptr = ::std::ptr::addr_of!((*obj_ptr).base);
+ assert!((obj_ptr as usize) == (base_ptr as usize));
+ }
+
+ Ok(obj_ptr)
+ }
+
+ fn from_ptr(ptr: *const $t) -> Self {
+ ptr as Self
+ }
+ }
+
+ // there are two reason to implement those traits for all objects
+ // 1. it speeds up operations
+ // 2. we want to check for real equality more explicit to stay conformant with the API
+ // and to not break in subtle ways e.g. using CL objects as keys in HashMaps.
+ impl std::cmp::Eq for $t {}
+ impl std::cmp::PartialEq for $t {
+ fn eq(&self, other: &Self) -> bool {
+ (self as *const Self) == (other as *const Self)
+ }
+ }
+
+ impl std::hash::Hash for $t {
+ fn hash<H: std::hash::Hasher>(&self, state: &mut H) {
+ (self as *const Self).hash(state);
+ }
+ }
+ };
+}
+
+// We need those functions exported
+
+#[no_mangle]
+extern "C" fn clGetPlatformInfo(
+ platform: cl_platform_id,
+ param_name: cl_platform_info,
+ param_value_size: usize,
+ param_value: *mut ::std::ffi::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ cl_get_platform_info(
+ platform,
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ )
+}
+
+#[no_mangle]
+extern "C" fn clGetExtensionFunctionAddress(
+ function_name: *const ::std::os::raw::c_char,
+) -> *mut ::std::ffi::c_void {
+ cl_get_extension_function_address(function_name)
+}
+
+#[no_mangle]
+extern "C" fn clIcdGetPlatformIDsKHR(
+ num_entries: cl_uint,
+ platforms: *mut cl_platform_id,
+ num_platforms: *mut cl_uint,
+) -> cl_int {
+ cl_icd_get_platform_ids_khr(num_entries, platforms, num_platforms)
+}
+
+// helper macros to make it less painful
+
+macro_rules! match_err {
+ ($exp: expr) => {
+ match $exp {
+ Ok(_) => CL_SUCCESS as cl_int,
+ Err(e) => e,
+ }
+ };
+}
+
+macro_rules! match_obj {
+ ($exp: expr, $err: ident) => {
+ match $exp {
+ Ok(o) => {
+ $err.write_checked(CL_SUCCESS as cl_int);
+ o
+ }
+ Err(e) => {
+ $err.write_checked(e);
+ ptr::null_mut()
+ }
+ }
+ };
+}
+
+macro_rules! match_obj_expl {
+ ($exp: expr, $err: ident) => {
+ match $exp {
+ Ok((o, c)) => {
+ $err.write_checked(c as cl_int);
+ o
+ }
+ Err(e) => {
+ $err.write_checked(e);
+ ptr::null_mut()
+ }
+ }
+ };
+}
+
+// extern "C" function stubs in ICD and extension order
+
+extern "C" fn cl_get_platform_ids(
+ num_entries: cl_uint,
+ platforms: *mut cl_platform_id,
+ num_platforms: *mut cl_uint,
+) -> cl_int {
+ match_err!(get_platform_ids(num_entries, platforms, num_platforms))
+}
+
+extern "C" fn cl_get_platform_info(
+ platform: cl_platform_id,
+ param_name: cl_platform_info,
+ param_value_size: usize,
+ param_value: *mut ::std::ffi::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(platform.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_get_device_ids(
+ platform: cl_platform_id,
+ device_type: cl_device_type,
+ num_entries: cl_uint,
+ devices: *mut cl_device_id,
+ num_devices: *mut cl_uint,
+) -> cl_int {
+ match_err!(get_device_ids(
+ platform,
+ device_type,
+ num_entries,
+ devices,
+ num_devices
+ ))
+}
+
+extern "C" fn cl_get_device_info(
+ device: cl_device_id,
+ param_name: cl_device_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(device.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_context(
+ properties: *const cl_context_properties,
+ num_devices: cl_uint,
+ devices: *const cl_device_id,
+ pfn_notify: Option<CreateContextCB>,
+ user_data: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_context {
+ match_obj!(
+ create_context(properties, num_devices, devices, pfn_notify, user_data),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_create_context_from_type(
+ properties: *const cl_context_properties,
+ device_type: cl_device_type,
+ pfn_notify: Option<CreateContextCB>,
+ user_data: *mut ::std::ffi::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_context {
+ match_obj!(
+ create_context_from_type(properties, device_type, pfn_notify, user_data),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_retain_context(context: cl_context) -> cl_int {
+ match_err!(context.retain())
+}
+
+extern "C" fn cl_release_context(context: cl_context) -> cl_int {
+ match_err!(context.release())
+}
+
+extern "C" fn cl_get_context_info(
+ context: cl_context,
+ param_name: cl_context_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(context.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_command_queue(
+ context: cl_context,
+ device: cl_device_id,
+ properties: cl_command_queue_properties,
+ errcode_ret: *mut cl_int,
+) -> cl_command_queue {
+ match_obj!(
+ create_command_queue(context, device, properties),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_retain_command_queue(command_queue: cl_command_queue) -> cl_int {
+ match_err!(command_queue.retain())
+}
+
+extern "C" fn cl_release_command_queue(command_queue: cl_command_queue) -> cl_int {
+ match_err!(command_queue.release())
+}
+
+extern "C" fn cl_get_command_queue_info(
+ command_queue: cl_command_queue,
+ param_name: cl_command_queue_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(command_queue.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_buffer(
+ context: cl_context,
+ flags: cl_mem_flags,
+ size: usize,
+ host_ptr: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_mem {
+ match_obj!(create_buffer(context, flags, size, host_ptr,), errcode_ret)
+}
+
+extern "C" fn cl_create_image_2d(
+ _context: cl_context,
+ _flags: cl_mem_flags,
+ _image_format: *const cl_image_format,
+ _image_width: usize,
+ _image_height: usize,
+ _image_row_pitch: usize,
+ _host_ptr: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_mem {
+ println!("cl_create_image_2d not implemented");
+ errcode_ret.write_checked(CL_OUT_OF_HOST_MEMORY);
+ ptr::null_mut()
+}
+
+extern "C" fn cl_create_image_3d(
+ _context: cl_context,
+ _flags: cl_mem_flags,
+ _image_format: *const cl_image_format,
+ _image_width: usize,
+ _image_height: usize,
+ _image_depth: usize,
+ _image_row_pitch: usize,
+ _image_slice_pitch: usize,
+ _host_ptr: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_mem {
+ println!("cl_create_image_3d not implemented");
+ errcode_ret.write_checked(CL_OUT_OF_HOST_MEMORY);
+ ptr::null_mut()
+}
+
+extern "C" fn cl_retain_mem_object(mem: cl_mem) -> cl_int {
+ match_err!(mem.retain())
+}
+
+extern "C" fn cl_release_mem_object(mem: cl_mem) -> cl_int {
+ match_err!(mem.release())
+}
+
+extern "C" fn cl_get_supported_image_formats(
+ context: cl_context,
+ flags: cl_mem_flags,
+ image_type: cl_mem_object_type,
+ num_entries: cl_uint,
+ image_formats: *mut cl_image_format,
+ num_image_formats: *mut cl_uint,
+) -> cl_int {
+ match_err!(get_supported_image_formats(
+ context,
+ flags,
+ image_type,
+ num_entries,
+ image_formats,
+ num_image_formats
+ ))
+}
+
+extern "C" fn cl_get_mem_object_info(
+ memobj: cl_mem,
+ param_name: cl_mem_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(memobj.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_get_image_info(
+ image: cl_mem,
+ param_name: cl_image_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(image.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_sampler(
+ context: cl_context,
+ normalized_coords: cl_bool,
+ addressing_mode: cl_addressing_mode,
+ filter_mode: cl_filter_mode,
+ errcode_ret: *mut cl_int,
+) -> cl_sampler {
+ match_obj!(
+ create_sampler(context, normalized_coords, addressing_mode, filter_mode),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_retain_sampler(sampler: cl_sampler) -> cl_int {
+ match_err!(sampler.retain())
+}
+
+extern "C" fn cl_release_sampler(sampler: cl_sampler) -> cl_int {
+ match_err!(sampler.release())
+}
+
+extern "C" fn cl_get_sampler_info(
+ sampler: cl_sampler,
+ param_name: cl_sampler_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(sampler.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_program_with_source(
+ context: cl_context,
+ count: cl_uint,
+ strings: *mut *const ::std::os::raw::c_char,
+ lengths: *const usize,
+ errcode_ret: *mut cl_int,
+) -> cl_program {
+ match_obj!(
+ create_program_with_source(context, count, strings, lengths),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_retain_program(program: cl_program) -> cl_int {
+ match_err!(program.retain())
+}
+
+extern "C" fn cl_release_program(program: cl_program) -> cl_int {
+ match_err!(program.release())
+}
+
+extern "C" fn cl_build_program(
+ program: cl_program,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ options: *const ::std::os::raw::c_char,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> cl_int {
+ match_err!(build_program(
+ program,
+ num_devices,
+ device_list,
+ options,
+ pfn_notify,
+ user_data,
+ ))
+}
+
+extern "C" fn cl_get_program_info(
+ program: cl_program,
+ param_name: cl_program_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(program.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_get_program_build_info(
+ program: cl_program,
+ device: cl_device_id,
+ param_name: cl_program_build_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(program.get_info_obj(
+ device,
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_create_kernel(
+ program: cl_program,
+ kernel_name: *const ::std::os::raw::c_char,
+ errcode_ret: *mut cl_int,
+) -> cl_kernel {
+ match_obj!(create_kernel(program, kernel_name), errcode_ret)
+}
+
+extern "C" fn cl_create_kernels_in_program(
+ _program: cl_program,
+ _num_kernels: cl_uint,
+ _kernels: *mut cl_kernel,
+ _num_kernels_ret: *mut cl_uint,
+) -> cl_int {
+ println!("cl_create_kernels_in_program not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_wait_for_events(_num_events: cl_uint, _event_list: *const cl_event) -> cl_int {
+ println!("cl_wait_for_events not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_get_event_info(
+ event: cl_event,
+ param_name: cl_event_info,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+) -> cl_int {
+ match_err!(event.get_info(
+ param_name,
+ param_value_size,
+ param_value,
+ param_value_size_ret,
+ ))
+}
+
+extern "C" fn cl_release_event(event: cl_event) -> cl_int {
+ match_err!(event.release())
+}
+
+extern "C" fn cl_get_event_profiling_info(
+ _event: cl_event,
+ _param_name: cl_profiling_info,
+ _param_value_size: usize,
+ _param_value: *mut ::std::os::raw::c_void,
+ _param_value_size_ret: *mut usize,
+) -> cl_int {
+ println!("cl_get_event_profiling_info not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_finish(command_queue: cl_command_queue) -> cl_int {
+ match_err!(finish_queue(command_queue))
+}
+
+extern "C" fn cl_enqueue_read_buffer(
+ _command_queue: cl_command_queue,
+ _buffer: cl_mem,
+ _blocking_read: cl_bool,
+ _offset: usize,
+ _cb: usize,
+ _ptr: *mut ::std::os::raw::c_void,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_read_buffer not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_write_buffer(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_write: cl_bool,
+ offset: usize,
+ cb: usize,
+ ptr: *const ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> cl_int {
+ match_err!(enqueue_write_buffer(
+ command_queue,
+ buffer,
+ blocking_write,
+ offset,
+ cb,
+ ptr,
+ num_events_in_wait_list,
+ event_wait_list,
+ event
+ ))
+}
+
+extern "C" fn cl_enqueue_copy_buffer(
+ _command_queue: cl_command_queue,
+ _src_buffer: cl_mem,
+ _dst_buffer: cl_mem,
+ _src_offset: usize,
+ _dst_offset: usize,
+ _cb: usize,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_copy_buffer not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_read_image(
+ _command_queue: cl_command_queue,
+ _image: cl_mem,
+ _blocking_read: cl_bool,
+ _origin: *const usize,
+ _region: *const usize,
+ _row_pitch: usize,
+ _slice_pitch: usize,
+ _ptr: *mut ::std::os::raw::c_void,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_read_image not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_write_image(
+ _command_queue: cl_command_queue,
+ _image: cl_mem,
+ _blocking_write: cl_bool,
+ _origin: *const usize,
+ _region: *const usize,
+ _input_row_pitch: usize,
+ _input_slice_pitch: usize,
+ _ptr: *const ::std::os::raw::c_void,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_write_image not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_copy_image(
+ _command_queue: cl_command_queue,
+ _src_image: cl_mem,
+ _dst_image: cl_mem,
+ _src_origin: *const usize,
+ _dst_origin: *const usize,
+ _region: *const usize,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_copy_image not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_copy_image_to_buffer(
+ _command_queue: cl_command_queue,
+ _src_image: cl_mem,
+ _dst_buffer: cl_mem,
+ _src_origin: *const usize,
+ _region: *const usize,
+ _dst_offset: usize,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_copy_image_to_buffer not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_copy_buffer_to_image(
+ _command_queue: cl_command_queue,
+ _src_buffer: cl_mem,
+ _dst_image: cl_mem,
+ _src_offset: usize,
+ _dst_origin: *const usize,
+ _region: *const usize,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_copy_buffer_to_image not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_map_buffer(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_map: cl_bool,
+ map_flags: cl_map_flags,
+ offset: usize,
+ cb: usize,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+ errcode_ret: *mut cl_int,
+) -> *mut ::std::os::raw::c_void {
+ match_obj!(
+ enqueue_map_buffer(
+ command_queue,
+ buffer,
+ blocking_map,
+ map_flags,
+ offset,
+ cb,
+ num_events_in_wait_list,
+ event_wait_list,
+ event,
+ ),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_enqueue_map_image(
+ _command_queue: cl_command_queue,
+ _image: cl_mem,
+ _blocking_map: cl_bool,
+ _map_flags: cl_map_flags,
+ _origin: *const usize,
+ _region: *const usize,
+ _image_row_pitch: *mut usize,
+ _image_slice_pitch: *mut usize,
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+ errcode_ret: *mut cl_int,
+) -> *mut ::std::os::raw::c_void {
+ println!("cl_enqueue_map_image not implemented");
+ errcode_ret.write_checked(CL_OUT_OF_HOST_MEMORY);
+ ptr::null_mut()
+}
+
+extern "C" fn cl_enqueue_unmap_mem_object(
+ command_queue: cl_command_queue,
+ memobj: cl_mem,
+ mapped_ptr: *mut ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> cl_int {
+ match_err!(enqueue_unmap_mem_object(
+ command_queue,
+ memobj,
+ mapped_ptr,
+ num_events_in_wait_list,
+ event_wait_list,
+ event
+ ))
+}
+
+extern "C" fn cl_get_extension_function_address(
+ function_name: *const ::std::os::raw::c_char,
+) -> *mut ::std::ffi::c_void {
+ if function_name.is_null() {
+ return ptr::null_mut();
+ }
+ match unsafe { CStr::from_ptr(function_name) }.to_str().unwrap() {
+ "clGetPlatformInfo" => cl_get_platform_info as *mut std::ffi::c_void,
+ "clIcdGetPlatformIDsKHR" => cl_icd_get_platform_ids_khr as *mut std::ffi::c_void,
+ _ => ptr::null_mut(),
+ }
+}
+
+extern "C" fn cl_set_event_callback(
+ _event: cl_event,
+ _command_exec_callback_type: cl_int,
+ _pfn_notify: Option<EventCB>,
+ _user_data: *mut ::std::os::raw::c_void,
+) -> cl_int {
+ println!("cl_set_event_callback not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_create_sub_buffer(
+ buffer: cl_mem,
+ flags: cl_mem_flags,
+ buffer_create_type: cl_buffer_create_type,
+ buffer_create_info: *const ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_mem {
+ match_obj!(
+ create_sub_buffer(buffer, flags, buffer_create_type, buffer_create_info,),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_set_mem_object_destructor_callback(
+ memobj: cl_mem,
+ pfn_notify: Option<MemCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> cl_int {
+ match_err!(set_mem_object_destructor_callback(
+ memobj, pfn_notify, user_data,
+ ))
+}
+
+extern "C" fn cl_create_user_event(context: cl_context, errcode_ret: *mut cl_int) -> cl_event {
+ match_obj!(create_user_event(context), errcode_ret)
+}
+
+extern "C" fn cl_set_user_event_status(_event: cl_event, _execution_status: cl_int) -> cl_int {
+ println!("cl_set_user_event_status not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_read_buffer_rect(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_read: cl_bool,
+ buffer_origin: *const usize,
+ host_origin: *const usize,
+ region: *const usize,
+ buffer_row_pitch: usize,
+ buffer_slice_pitch: usize,
+ host_row_pitch: usize,
+ host_slice_pitch: usize,
+ ptr: *mut ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> cl_int {
+ match_err!(enqueue_read_buffer_rect(
+ command_queue,
+ buffer,
+ blocking_read,
+ buffer_origin,
+ host_origin,
+ region,
+ buffer_row_pitch,
+ buffer_slice_pitch,
+ host_row_pitch,
+ host_slice_pitch,
+ ptr,
+ num_events_in_wait_list,
+ event_wait_list,
+ event,
+ ))
+}
+
+extern "C" fn cl_enqueue_write_buffer_rect(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_write: cl_bool,
+ buffer_origin: *const usize,
+ host_origin: *const usize,
+ region: *const usize,
+ buffer_row_pitch: usize,
+ buffer_slice_pitch: usize,
+ host_row_pitch: usize,
+ host_slice_pitch: usize,
+ ptr: *const ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> cl_int {
+ match_err!(enqueue_write_buffer_rect(
+ command_queue,
+ buffer,
+ blocking_write,
+ buffer_origin,
+ host_origin,
+ region,
+ buffer_row_pitch,
+ buffer_slice_pitch,
+ host_row_pitch,
+ host_slice_pitch,
+ ptr,
+ num_events_in_wait_list,
+ event_wait_list,
+ event,
+ ))
+}
+
+extern "C" fn cl_enqueue_copy_buffer_rect(
+ command_queue: cl_command_queue,
+ src_buffer: cl_mem,
+ dst_buffer: cl_mem,
+ src_origin: *const usize,
+ dst_origin: *const usize,
+ region: *const usize,
+ src_row_pitch: usize,
+ src_slice_pitch: usize,
+ dst_row_pitch: usize,
+ dst_slice_pitch: usize,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> cl_int {
+ match_err!(enqueue_copy_buffer_rect(
+ command_queue,
+ src_buffer,
+ dst_buffer,
+ src_origin,
+ dst_origin,
+ region,
+ src_row_pitch,
+ src_slice_pitch,
+ dst_row_pitch,
+ dst_slice_pitch,
+ num_events_in_wait_list,
+ event_wait_list,
+ event,
+ ))
+}
+
+extern "C" fn cl_create_image(
+ context: cl_context,
+ flags: cl_mem_flags,
+ image_format: *const cl_image_format,
+ image_desc: *const cl_image_desc,
+ host_ptr: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_mem {
+ match_obj!(
+ create_image(context, flags, image_format, image_desc, host_ptr),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_compile_program(
+ program: cl_program,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ options: *const ::std::os::raw::c_char,
+ num_input_headers: cl_uint,
+ input_headers: *const cl_program,
+ header_include_names: *mut *const ::std::os::raw::c_char,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> cl_int {
+ match_err!(compile_program(
+ program,
+ num_devices,
+ device_list,
+ options,
+ num_input_headers,
+ input_headers,
+ header_include_names,
+ pfn_notify,
+ user_data,
+ ))
+}
+
+extern "C" fn cl_link_program(
+ context: cl_context,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ options: *const ::std::os::raw::c_char,
+ num_input_programs: cl_uint,
+ input_programs: *const cl_program,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+ errcode_ret: *mut cl_int,
+) -> cl_program {
+ match_obj_expl!(
+ link_program(
+ context,
+ num_devices,
+ device_list,
+ options,
+ num_input_programs,
+ input_programs,
+ pfn_notify,
+ user_data,
+ ),
+ errcode_ret
+ )
+}
+
+extern "C" fn cl_unload_platform_compiler(_platform: cl_platform_id) -> cl_int {
+ println!("cl_unload_platform_compiler not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+extern "C" fn cl_enqueue_fill_image(
+ _command_queue: cl_command_queue,
+ _image: cl_mem,
+ _fill_color: *const ::std::os::raw::c_void,
+ _origin: *const [usize; 3usize],
+ _region: *const [usize; 3usize],
+ _num_events_in_wait_list: cl_uint,
+ _event_wait_list: *const cl_event,
+ _event: *mut cl_event,
+) -> cl_int {
+ println!("cl_enqueue_fill_image not implemented");
+ CL_OUT_OF_HOST_MEMORY
+}
+
+// cl_khr_icd
+extern "C" fn cl_icd_get_platform_ids_khr(
+ num_entries: cl_uint,
+ platforms: *mut cl_platform_id,
+ num_platforms: *mut cl_uint,
+) -> cl_int {
+ match_err!(get_platform_ids(num_entries, platforms, num_platforms))
+}
--- /dev/null
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+
+use self::rusticl_opencl_gen::*;
+
+pub fn create_kernel(
+ program: cl_program,
+ kernel_name: *const ::std::os::raw::c_char,
+) -> CLResult<cl_kernel> {
+ let _p = program.get_ref()?;
+
+ // CL_INVALID_VALUE if kernel_name is NULL.
+ if kernel_name.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ println!("create_kernel not implemented");
+ Err(CL_OUT_OF_HOST_MEMORY)
+
+ //• CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
+ //• CL_INVALID_KERNEL_NAME if kernel_name is not found in program.
+ //• CL_INVALID_KERNEL_DEFINITION if the function definition for __kernel function given by kernel_name such as the number of arguments, the argument types are not the same for all devices for which the program executable has been built.
+}
--- /dev/null
+#![allow(non_upper_case_globals)]
+
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::event::create_and_queue;
+use crate::api::icd::*;
+use crate::api::types::*;
+use crate::api::util::*;
+use crate::core::device::*;
+use crate::core::memory::*;
+use crate::*;
+
+use self::mesa_rust_util::ptr::*;
+use self::rusticl_opencl_gen::*;
+
+use std::cmp::Ordering;
+use std::os::raw::c_void;
+use std::ptr;
+use std::sync::Arc;
+
+fn validate_mem_flags(flags: cl_mem_flags, images: bool) -> CLResult<()> {
+ let mut valid_flags = cl_bitfield::from(
+ CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | CL_MEM_KERNEL_READ_AND_WRITE,
+ );
+
+ if !images {
+ valid_flags |= cl_bitfield::from(
+ CL_MEM_USE_HOST_PTR
+ | CL_MEM_ALLOC_HOST_PTR
+ | CL_MEM_COPY_HOST_PTR
+ | CL_MEM_HOST_WRITE_ONLY
+ | CL_MEM_HOST_READ_ONLY
+ | CL_MEM_HOST_NO_ACCESS,
+ );
+ }
+
+ let read_write_group =
+ cl_bitfield::from(CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY);
+
+ let alloc_host_group = cl_bitfield::from(CL_MEM_ALLOC_HOST_PTR | CL_MEM_USE_HOST_PTR);
+
+ let copy_host_group = cl_bitfield::from(CL_MEM_COPY_HOST_PTR | CL_MEM_USE_HOST_PTR);
+
+ let host_read_write_group =
+ cl_bitfield::from(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS);
+
+ if (flags & !valid_flags != 0)
+ || (flags & read_write_group).count_ones() > 1
+ || (flags & alloc_host_group).count_ones() > 1
+ || (flags & copy_host_group).count_ones() > 1
+ || (flags & host_read_write_group).count_ones() > 1
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+ Ok(())
+}
+
+fn filter_image_access_flags(flags: cl_mem_flags) -> cl_mem_flags {
+ flags
+ & (CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | CL_MEM_KERNEL_READ_AND_WRITE)
+ as cl_mem_flags
+}
+
+fn inherit_mem_flags(mut flags: cl_mem_flags, mem: &Mem) -> cl_mem_flags {
+ let read_write_mask = cl_bitfield::from(
+ CL_MEM_READ_WRITE |
+ CL_MEM_WRITE_ONLY |
+ CL_MEM_READ_ONLY |
+ // not in spec, but...
+ CL_MEM_KERNEL_READ_AND_WRITE,
+ );
+ let host_ptr_mask =
+ cl_bitfield::from(CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR);
+ let host_mask =
+ cl_bitfield::from(CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS);
+
+ // For CL_MEM_OBJECT_IMAGE1D_BUFFER image type, or an image created from another memory object
+ // (image or buffer)...
+ //
+ // ... if the CL_MEM_READ_WRITE, CL_MEM_READ_ONLY or CL_MEM_WRITE_ONLY values are not
+ // specified in flags, they are inherited from the corresponding memory access qualifiers
+ // associated with mem_object. ...
+ if flags & read_write_mask == 0 {
+ flags |= mem.flags & read_write_mask;
+ }
+
+ // ... The CL_MEM_USE_HOST_PTR, CL_MEM_ALLOC_HOST_PTR and CL_MEM_COPY_HOST_PTR values cannot
+ // be specified in flags but are inherited from the corresponding memory access qualifiers
+ // associated with mem_object. ...
+ flags &= !host_ptr_mask;
+ flags |= mem.flags & host_ptr_mask;
+
+ // ... If the CL_MEM_HOST_WRITE_ONLY, CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS values
+ // are not specified in flags, they are inherited from the corresponding memory access
+ // qualifiers associated with mem_object.
+ if flags & host_mask == 0 {
+ flags |= mem.flags & host_mask;
+ }
+
+ flags
+}
+
+fn image_type_valid(image_type: cl_mem_object_type) -> bool {
+ CL_IMAGE_TYPES.contains(&image_type)
+}
+
+fn validate_addressing_mode(addressing_mode: cl_addressing_mode) -> CLResult<()> {
+ match addressing_mode {
+ CL_ADDRESS_NONE
+ | CL_ADDRESS_CLAMP_TO_EDGE
+ | CL_ADDRESS_CLAMP
+ | CL_ADDRESS_REPEAT
+ | CL_ADDRESS_MIRRORED_REPEAT => Ok(()),
+ _ => Err(CL_INVALID_VALUE),
+ }
+}
+
+fn validate_filter_mode(filter_mode: cl_filter_mode) -> CLResult<()> {
+ match filter_mode {
+ CL_FILTER_NEAREST | CL_FILTER_LINEAR => Ok(()),
+ _ => Err(CL_INVALID_VALUE),
+ }
+}
+
+fn validate_host_ptr(host_ptr: *mut ::std::os::raw::c_void, flags: cl_mem_flags) -> CLResult<()> {
+ // CL_INVALID_HOST_PTR if host_ptr is NULL and CL_MEM_USE_HOST_PTR or CL_MEM_COPY_HOST_PTR are
+ // set in flags
+ if host_ptr.is_null()
+ && flags & (cl_mem_flags::from(CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)) != 0
+ {
+ return Err(CL_INVALID_HOST_PTR);
+ }
+
+ // or if host_ptr is not NULL but CL_MEM_COPY_HOST_PTR or CL_MEM_USE_HOST_PTR are not set in
+ // flags.
+ if !host_ptr.is_null()
+ && flags & (cl_mem_flags::from(CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)) == 0
+ {
+ return Err(CL_INVALID_HOST_PTR);
+ }
+
+ Ok(())
+}
+
+fn validate_matching_buffer_flags(mem: &Mem, flags: cl_mem_flags) -> CLResult<()> {
+ // CL_INVALID_VALUE if an image is being created from another memory object (buffer or image)
+ // under one of the following circumstances:
+ //
+ // 1) mem_object was created with CL_MEM_WRITE_ONLY and
+ // flags specifies CL_MEM_READ_WRITE or CL_MEM_READ_ONLY,
+ if bit_check(mem.flags, CL_MEM_WRITE_ONLY) && bit_check(flags, CL_MEM_READ_WRITE | CL_MEM_READ_ONLY) ||
+ // 2) mem_object was created with CL_MEM_READ_ONLY and
+ // flags specifies CL_MEM_READ_WRITE or CL_MEM_WRITE_ONLY,
+ bit_check(mem.flags, CL_MEM_READ_ONLY) && bit_check(flags, CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY) ||
+ // 3) flags specifies CL_MEM_USE_HOST_PTR or CL_MEM_ALLOC_HOST_PTR or CL_MEM_COPY_HOST_PTR.
+ bit_check(flags, CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR) ||
+ // CL_INVALID_VALUE if an image is being created from another memory object (buffer or image)
+ // and mem_object was created with CL_MEM_HOST_WRITE_ONLY and flags specifies CL_MEM_HOST_READ_ONLY
+ bit_check(mem.flags, CL_MEM_HOST_WRITE_ONLY) && bit_check(flags, CL_MEM_HOST_READ_ONLY) ||
+ // or if mem_object was created with CL_MEM_HOST_READ_ONLY and flags specifies CL_MEM_HOST_WRITE_ONLY
+ bit_check(mem.flags, CL_MEM_HOST_READ_ONLY) && bit_check(flags, CL_MEM_HOST_WRITE_ONLY) ||
+ // or if mem_object was created with CL_MEM_HOST_NO_ACCESS and_flags_ specifies CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_WRITE_ONLY.
+ bit_check(mem.flags, CL_MEM_HOST_NO_ACCESS) && bit_check(flags, CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_WRITE_ONLY)
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ Ok(())
+}
+
+impl CLInfo<cl_mem_info> for cl_mem {
+ fn query(&self, q: cl_mem_info) -> CLResult<Vec<u8>> {
+ let mem = self.get_ref()?;
+ Ok(match *q {
+ CL_MEM_ASSOCIATED_MEMOBJECT => {
+ let ptr = match mem.parent.as_ref() {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ Some(parent) => Arc::as_ptr(parent),
+ None => ptr::null(),
+ };
+ cl_prop::<cl_mem>(cl_mem::from_ptr(ptr))
+ }
+ CL_MEM_CONTEXT => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&mem.context);
+ cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+ }
+ CL_MEM_FLAGS => cl_prop::<cl_mem_flags>(mem.flags),
+ // TODO debugging feature
+ CL_MEM_MAP_COUNT => cl_prop::<cl_uint>(0),
+ CL_MEM_HOST_PTR => cl_prop::<*mut c_void>(mem.host_ptr),
+ CL_MEM_OFFSET => cl_prop::<usize>(mem.offset),
+ CL_MEM_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ CL_MEM_SIZE => cl_prop::<usize>(mem.size),
+ CL_MEM_TYPE => cl_prop::<cl_mem_object_type>(mem.mem_type),
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+pub fn create_buffer(
+ context: cl_context,
+ flags: cl_mem_flags,
+ size: usize,
+ host_ptr: *mut ::std::os::raw::c_void,
+) -> CLResult<cl_mem> {
+ let c = context.get_arc()?;
+
+ // CL_INVALID_VALUE if values specified in flags are not valid as defined in the Memory Flags table.
+ validate_mem_flags(flags, false)?;
+
+ // CL_INVALID_BUFFER_SIZE if size is 0
+ if size == 0 {
+ return Err(CL_INVALID_BUFFER_SIZE);
+ }
+
+ // ... or if size is greater than CL_DEVICE_MAX_MEM_ALLOC_SIZE for all devices in context.
+ for dev in &c.devs {
+ if checked_compare(size, Ordering::Greater, dev.max_mem_alloc()) {
+ return Err(CL_INVALID_BUFFER_SIZE);
+ }
+ }
+
+ validate_host_ptr(host_ptr, flags)?;
+
+ Ok(cl_mem::from_arc(Mem::new_buffer(c, flags, size, host_ptr)?))
+}
+
+pub fn create_sub_buffer(
+ buffer: cl_mem,
+ mut flags: cl_mem_flags,
+ buffer_create_type: cl_buffer_create_type,
+ buffer_create_info: *const ::std::os::raw::c_void,
+) -> CLResult<cl_mem> {
+ let b = buffer.get_arc()?;
+
+ // CL_INVALID_MEM_OBJECT if buffer ... is a sub-buffer object.
+ if b.parent.is_some() {
+ return Err(CL_INVALID_MEM_OBJECT);
+ }
+
+ validate_matching_buffer_flags(&b, flags)?;
+
+ flags = inherit_mem_flags(flags, &b);
+ validate_mem_flags(flags, false)?;
+
+ let (offset, size) = match buffer_create_type {
+ CL_BUFFER_CREATE_TYPE_REGION => {
+ // buffer_create_info is a pointer to a cl_buffer_region structure specifying a region of
+ // the buffer.
+ // CL_INVALID_VALUE if value(s) specified in buffer_create_info (for a given
+ // buffer_create_type) is not valid or if buffer_create_info is NULL.
+ let region = unsafe { buffer_create_info.cast::<cl_buffer_region>().as_ref() }
+ .ok_or(CL_INVALID_VALUE)?;
+
+ // CL_INVALID_BUFFER_SIZE if the size field of the cl_buffer_region structure passed in
+ // buffer_create_info is 0.
+ if region.size == 0 {
+ return Err(CL_INVALID_BUFFER_SIZE);
+ }
+
+ // CL_INVALID_VALUE if the region specified by the cl_buffer_region structure passed in
+ // buffer_create_info is out of bounds in buffer.
+ if region.origin + region.size > b.size {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ (region.origin, region.size)
+ }
+ // CL_INVALID_VALUE if the value specified in buffer_create_type is not valid.
+ _ => return Err(CL_INVALID_VALUE),
+ };
+
+ Ok(cl_mem::from_arc(Mem::new_sub_buffer(
+ b, flags, offset, size,
+ )))
+
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if there are no devices in context associated with buffer for which the origin field of the cl_buffer_region structure passed in buffer_create_info is aligned to the CL_DEVICE_MEM_BASE_ADDR_ALIGN value.
+}
+
+pub fn set_mem_object_destructor_callback(
+ memobj: cl_mem,
+ pfn_notify: Option<MemCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<()> {
+ let m = memobj.get_ref()?;
+
+ // CL_INVALID_VALUE if pfn_notify is NULL.
+ if pfn_notify.is_none() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ m.cbs
+ .lock()
+ .unwrap()
+ .push(cl_closure!(|m| pfn_notify(m, user_data)));
+ Ok(())
+}
+
+fn validate_image_format<'a>(
+ image_format: *const cl_image_format,
+) -> CLResult<(&'a cl_image_format, u8)> {
+ // CL_INVALID_IMAGE_FORMAT_DESCRIPTOR ... if image_format is NULL.
+ let format = unsafe { image_format.as_ref() }.ok_or(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR)?;
+
+ let channels = match format.image_channel_order {
+ CL_R | CL_A | CL_DEPTH | CL_LUMINANCE | CL_INTENSITY => 1,
+
+ CL_RG | CL_RA | CL_Rx => 2,
+
+ CL_RGB | CL_RGx | CL_sRGB => 3,
+
+ CL_RGBA | CL_ARGB | CL_BGRA | CL_ABGR | CL_RGBx | CL_sRGBA | CL_sBGRA | CL_sRGBx => 4,
+
+ _ => return Err(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR),
+ };
+
+ let channel_size = match format.image_channel_data_type {
+ CL_SNORM_INT8 | CL_UNORM_INT8 | CL_SIGNED_INT8 | CL_UNSIGNED_INT8 => 1,
+
+ CL_SNORM_INT16 | CL_UNORM_INT16 | CL_SIGNED_INT16 | CL_UNSIGNED_INT16 | CL_HALF_FLOAT
+ | CL_UNORM_SHORT_565 | CL_UNORM_SHORT_555 => 2,
+
+ CL_SIGNED_INT32
+ | CL_UNSIGNED_INT32
+ | CL_FLOAT
+ | CL_UNORM_INT_101010
+ | CL_UNORM_INT_101010_2 => 4,
+
+ _ => return Err(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR),
+ };
+
+ let packed = [
+ CL_UNORM_SHORT_565,
+ CL_UNORM_SHORT_555,
+ CL_UNORM_INT_101010,
+ CL_UNORM_INT_101010,
+ ]
+ .contains(&format.image_channel_data_type);
+
+ // special validation
+ let valid_combination = match format.image_channel_data_type {
+ CL_UNORM_SHORT_565 | CL_UNORM_SHORT_555 | CL_UNORM_INT_101010 => {
+ [CL_RGB, CL_RGBx].contains(&format.image_channel_data_type)
+ }
+ CL_UNORM_INT_101010_2 => format.image_channel_data_type == CL_RGBA,
+ _ => true,
+ };
+ if !valid_combination {
+ return Err(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
+ }
+
+ Ok((
+ format,
+ if packed {
+ channel_size
+ } else {
+ channels * channel_size
+ },
+ ))
+}
+
+fn validate_image_desc(
+ image_desc: *const cl_image_desc,
+ host_ptr: *mut ::std::os::raw::c_void,
+ elem_size: usize,
+ devs: &[Arc<Device>],
+) -> CLResult<cl_image_desc> {
+ // CL_INVALID_IMAGE_DESCRIPTOR if values specified in image_desc are not valid
+ const err: cl_int = CL_INVALID_IMAGE_DESCRIPTOR;
+
+ // CL_INVALID_IMAGE_DESCRIPTOR ... if image_desc is NULL.
+ let mut desc = *unsafe { image_desc.as_ref() }.ok_or(err)?;
+
+ // image_type describes the image type and must be either CL_MEM_OBJECT_IMAGE1D,
+ // CL_MEM_OBJECT_IMAGE1D_BUFFER, CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D,
+ // CL_MEM_OBJECT_IMAGE2D_ARRAY, or CL_MEM_OBJECT_IMAGE3D.
+ let (dims, array) = match desc.image_type {
+ CL_MEM_OBJECT_IMAGE1D | CL_MEM_OBJECT_IMAGE1D_BUFFER => (1, false),
+ CL_MEM_OBJECT_IMAGE1D_ARRAY => (1, true),
+ CL_MEM_OBJECT_IMAGE2D => (2, false),
+ CL_MEM_OBJECT_IMAGE2D_ARRAY => (2, true),
+ CL_MEM_OBJECT_IMAGE3D => (3, false),
+ _ => return Err(err),
+ };
+
+ // image_width is the width of the image in pixels. For a 2D image and image array, the image
+ // width must be a value ≥ 1 and ≤ CL_DEVICE_IMAGE2D_MAX_WIDTH. For a 3D image, the image width
+ // must be a value ≥ 1 and ≤ CL_DEVICE_IMAGE3D_MAX_WIDTH. For a 1D image buffer, the image width
+ // must be a value ≥ 1 and ≤ CL_DEVICE_IMAGE_MAX_BUFFER_SIZE. For a 1D image and 1D image array,
+ // the image width must be a value ≥ 1 and ≤ CL_DEVICE_IMAGE2D_MAX_WIDTH.
+ //
+ // image_height is the height of the image in pixels. This is only used if the image is a 2D or
+ // 3D image, or a 2D image array. For a 2D image or image array, the image height must be a
+ // value ≥ 1 and ≤ CL_DEVICE_IMAGE2D_MAX_HEIGHT. For a 3D image, the image height must be a
+ // value ≥ 1 and ≤ CL_DEVICE_IMAGE3D_MAX_HEIGHT.
+ //
+ // image_depth is the depth of the image in pixels. This is only used if the image is a 3D image
+ // and must be a value ≥ 1 and ≤ CL_DEVICE_IMAGE3D_MAX_DEPTH.
+ if desc.image_width < 1
+ || desc.image_height < 1 && dims >= 2
+ || desc.image_depth < 1 && dims >= 3
+ || desc.image_array_size < 1 && array
+ {
+ return Err(err);
+ }
+
+ let max_size = if dims == 3 {
+ devs.iter().map(|d| d.image_3d_size()).min()
+ } else if desc.image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER {
+ devs.iter().map(|d| d.image_buffer_size()).min()
+ } else {
+ devs.iter().map(|d| d.image_2d_size()).min()
+ }
+ .unwrap();
+ let max_array = devs.iter().map(|d| d.image_array_size()).min().unwrap();
+
+ // CL_INVALID_IMAGE_SIZE if image dimensions specified in image_desc exceed the maximum image
+ // dimensions described in the Device Queries table for all devices in context.
+ if desc.image_width > max_size
+ || desc.image_height > max_size && dims >= 2
+ || desc.image_depth > max_size && dims >= 3
+ || desc.image_array_size > max_array && array
+ {
+ return Err(CL_INVALID_IMAGE_SIZE);
+ }
+
+ // num_mip_levels and num_samples must be 0.
+ if desc.num_mip_levels != 0 || desc.num_samples != 0 {
+ return Err(err);
+ }
+
+ // mem_object may refer to a valid buffer or image memory object. mem_object can be a buffer
+ // memory object if image_type is CL_MEM_OBJECT_IMAGE1D_BUFFER or CL_MEM_OBJECT_IMAGE2D.
+ // mem_object can be an image object if image_type is CL_MEM_OBJECT_IMAGE2D. Otherwise it must
+ // be NULL.
+ //
+ // TODO: cl_khr_image2d_from_buffer is an optional feature
+ let p = unsafe { &desc.anon_1.mem_object };
+ if !p.is_null() {
+ let p = p.get_ref()?;
+ if !match desc.image_type {
+ CL_MEM_OBJECT_IMAGE1D_BUFFER => p.is_buffer(),
+ CL_MEM_OBJECT_IMAGE2D => !p.is_buffer(),
+ _ => false,
+ } {
+ return Err(CL_INVALID_OPERATION);
+ }
+ }
+
+ // image_row_pitch is the scan-line pitch in bytes. This must be 0 if host_ptr is NULL and can
+ // be either 0 or ≥ image_width × size of element in bytes if host_ptr is not NULL. If host_ptr
+ // is not NULL and image_row_pitch = 0, image_row_pitch is calculated as image_width × size of
+ // element in bytes. If image_row_pitch is not 0, it must be a multiple of the image element
+ // size in bytes. For a 2D image created from a buffer, the pitch specified (or computed if
+ // pitch specified is 0) must be a multiple of the maximum of the
+ // CL_DEVICE_IMAGE_PITCH_ALIGNMENT value for all devices in the context associated with the
+ // buffer specified by mem_object that support images.
+ //
+ // image_slice_pitch is the size in bytes of each 2D slice in the 3D image or the size in bytes
+ // of each image in a 1D or 2D image array. This must be 0 if host_ptr is NULL. If host_ptr is
+ // not NULL, image_slice_pitch can be either 0 or ≥ image_row_pitch × image_height for a 2D
+ // image array or 3D image and can be either 0 or ≥ image_row_pitch for a 1D image array. If
+ // host_ptr is not NULL and image_slice_pitch = 0, image_slice_pitch is calculated as
+ // image_row_pitch × image_height for a 2D image array or 3D image and image_row_pitch for a 1D
+ // image array. If image_slice_pitch is not 0, it must be a multiple of the image_row_pitch.
+ if host_ptr.is_null() {
+ if desc.image_row_pitch != 0 || desc.image_slice_pitch != 0 {
+ return Err(err);
+ }
+ } else {
+ if desc.image_row_pitch == 0 {
+ desc.image_row_pitch = desc.image_width * elem_size;
+ } else if desc.image_row_pitch % elem_size != 0 {
+ return Err(err);
+ }
+
+ if dims == 3 || array {
+ let valid_slice_pitch =
+ desc.image_row_pitch * if dims == 1 { 1 } else { desc.image_height };
+ if desc.image_slice_pitch == 0 {
+ desc.image_slice_pitch = valid_slice_pitch;
+ } else if desc.image_slice_pitch < valid_slice_pitch
+ || desc.image_slice_pitch % desc.image_row_pitch != 0
+ {
+ return Err(err);
+ }
+ }
+ }
+
+ Ok(desc)
+}
+
+fn desc_eq_no_buffer(a: &cl_image_desc, b: &cl_image_desc) -> bool {
+ a.image_type == b.image_type
+ && a.image_width == b.image_width
+ && a.image_height == b.image_height
+ && a.image_depth == b.image_depth
+ && a.image_array_size == b.image_array_size
+ && a.image_row_pitch == b.image_row_pitch
+ && a.image_slice_pitch == b.image_slice_pitch
+ && a.num_mip_levels == b.num_mip_levels
+ && a.num_samples == b.num_samples
+}
+
+fn validate_buffer(
+ desc: &cl_image_desc,
+ mut flags: cl_mem_flags,
+ format: &cl_image_format,
+ host_ptr: *mut ::std::os::raw::c_void,
+ elem_size: usize,
+) -> CLResult<cl_mem_flags> {
+ // CL_INVALID_IMAGE_DESCRIPTOR if values specified in image_desc are not valid
+ const err: cl_int = CL_INVALID_IMAGE_DESCRIPTOR;
+ let mem_object = unsafe { desc.anon_1.mem_object };
+
+ // mem_object may refer to a valid buffer or image memory object. mem_object can be a buffer
+ // memory object if image_type is CL_MEM_OBJECT_IMAGE1D_BUFFER or CL_MEM_OBJECT_IMAGE2D
+ // mem_object can be an image object if image_type is CL_MEM_OBJECT_IMAGE2D. Otherwise it must
+ // be NULL. The image pixels are taken from the memory objects data store. When the contents of
+ // the specified memory objects data store are modified, those changes are reflected in the
+ // contents of the image object and vice-versa at corresponding synchronization points.
+ if !mem_object.is_null() {
+ let mem = mem_object.get_ref()?;
+
+ match mem.mem_type {
+ CL_MEM_OBJECT_BUFFER => {
+ match desc.image_type {
+ // For a 1D image buffer created from a buffer object, the image_width × size of
+ // element in bytes must be ≤ size of the buffer object.
+ CL_MEM_OBJECT_IMAGE1D_BUFFER => {
+ if desc.image_width * elem_size > mem.size {
+ return Err(err);
+ }
+ }
+ // For a 2D image created from a buffer object, the image_row_pitch × image_height
+ // must be ≤ size of the buffer object specified by mem_object.
+ CL_MEM_OBJECT_IMAGE2D => {
+ //TODO
+ //• CL_INVALID_IMAGE_FORMAT_DESCRIPTOR if a 2D image is created from a buffer and the row pitch and base address alignment does not follow the rules described for creating a 2D image from a buffer.
+ if desc.image_row_pitch * desc.image_height > mem.size {
+ return Err(err);
+ }
+ }
+ _ => return Err(err),
+ }
+ }
+ // For an image object created from another image object, the values specified in the
+ // image descriptor except for mem_object must match the image descriptor information
+ // associated with mem_object.
+ CL_MEM_OBJECT_IMAGE2D => {
+ if desc.image_type != mem.mem_type || !desc_eq_no_buffer(desc, &mem.image_desc) {
+ return Err(err);
+ }
+
+ // CL_INVALID_IMAGE_FORMAT_DESCRIPTOR if a 2D image is created from a 2D image object
+ // and the rules described above are not followed.
+
+ // Creating a 2D image object from another 2D image object creates a new 2D image
+ // object that shares the image data store with mem_object but views the pixels in the
+ // image with a different image channel order. Restrictions are:
+ //
+ // The image channel data type specified in image_format must match the image channel
+ // data type associated with mem_object.
+ if format.image_channel_data_type != mem.image_format.image_channel_data_type {
+ return Err(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
+ }
+
+ // The image channel order specified in image_format must be compatible with the image
+ // channel order associated with mem_object. Compatible image channel orders are:
+ if format.image_channel_order != mem.image_format.image_channel_order {
+ // in image_format | in mem_object:
+ // CL_sBGRA | CL_BGRA
+ // CL_BGRA | CL_sBGRA
+ // CL_sRGBA | CL_RGBA
+ // CL_RGBA | CL_sRGBA
+ // CL_sRGB | CL_RGB
+ // CL_RGB | CL_sRGB
+ // CL_sRGBx | CL_RGBx
+ // CL_RGBx | CL_sRGBx
+ // CL_DEPTH | CL_R
+ match (
+ format.image_channel_order,
+ mem.image_format.image_channel_order,
+ ) {
+ (CL_sBGRA, CL_BGRA)
+ | (CL_BGRA, CL_sBGRA)
+ | (CL_sRGBA, CL_RGBA)
+ | (CL_RGBA, CL_sRGBA)
+ | (CL_sRGB, CL_RGB)
+ | (CL_RGB, CL_sRGB)
+ | (CL_sRGBx, CL_RGBx)
+ | (CL_RGBx, CL_sRGBx)
+ | (CL_DEPTH, CL_R) => (),
+ _ => return Err(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR),
+ }
+ }
+ }
+ _ => return Err(err),
+ }
+
+ // If the buffer object specified by mem_object was created with CL_MEM_USE_HOST_PTR, the
+ // host_ptr specified to clCreateBuffer or clCreateBufferWithProperties must be aligned to
+ // the maximum of the CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT value for all devices in the
+ // context associated with the buffer specified by mem_object that support images.
+ if mem.flags & CL_MEM_USE_HOST_PTR as cl_mem_flags != 0 {
+ for dev in &mem.context.devs {
+ let addr_alignment = dev.image_base_address_alignment();
+ if addr_alignment == 0 {
+ return Err(CL_INVALID_OPERATION);
+ } else if !is_alligned(host_ptr, addr_alignment as usize) {
+ return Err(err);
+ }
+ }
+ }
+
+ validate_matching_buffer_flags(mem, flags)?;
+
+ flags = inherit_mem_flags(flags, mem);
+ // implied by spec
+ } else if desc.image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER {
+ return Err(err);
+ }
+
+ Ok(flags)
+}
+
+impl CLInfo<cl_image_info> for cl_mem {
+ fn query(&self, q: cl_image_info) -> CLResult<Vec<u8>> {
+ let mem = self.get_ref()?;
+ Ok(match *q {
+ CL_IMAGE_ARRAY_SIZE => cl_prop::<usize>(mem.image_desc.image_array_size),
+ CL_IMAGE_BUFFER => cl_prop::<cl_mem>(unsafe { mem.image_desc.anon_1.buffer }),
+ CL_IMAGE_DEPTH => cl_prop::<usize>(mem.image_desc.image_depth),
+ CL_IMAGE_ELEMENT_SIZE => cl_prop::<usize>(mem.image_elem_size.into()),
+ CL_IMAGE_FORMAT => cl_prop::<cl_image_format>(mem.image_format),
+ CL_IMAGE_HEIGHT => cl_prop::<usize>(mem.image_desc.image_height),
+ CL_IMAGE_NUM_MIP_LEVELS => cl_prop::<cl_uint>(mem.image_desc.num_mip_levels),
+ CL_IMAGE_NUM_SAMPLES => cl_prop::<cl_uint>(mem.image_desc.num_samples),
+ CL_IMAGE_ROW_PITCH => cl_prop::<usize>(mem.image_desc.image_row_pitch),
+ CL_IMAGE_SLICE_PITCH => cl_prop::<usize>(mem.image_desc.image_slice_pitch),
+ CL_IMAGE_WIDTH => cl_prop::<usize>(mem.image_desc.image_width),
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+pub fn create_image(
+ context: cl_context,
+ mut flags: cl_mem_flags,
+ image_format: *const cl_image_format,
+ image_desc: *const cl_image_desc,
+ host_ptr: *mut ::std::os::raw::c_void,
+) -> CLResult<cl_mem> {
+ let c = context.get_arc()?;
+
+ // CL_INVALID_OPERATION if there are no devices in context that support images (i.e.
+ // CL_DEVICE_IMAGE_SUPPORT specified in the Device Queries table is CL_FALSE).
+ c.devs
+ .iter()
+ .find(|d| d.image_supported())
+ .ok_or(CL_INVALID_OPERATION)?;
+
+ let (format, elem_size) = validate_image_format(image_format)?;
+ let desc = validate_image_desc(image_desc, host_ptr, elem_size.into(), &c.devs)?;
+ flags = validate_buffer(&desc, flags, format, host_ptr, elem_size.into())?;
+
+ // For all image types except CL_MEM_OBJECT_IMAGE1D_BUFFER, if the value specified for flags is 0, the
+ // default is used which is CL_MEM_READ_WRITE.
+ if flags == 0 && desc.image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER {
+ flags = CL_MEM_READ_WRITE.into();
+ }
+
+ validate_mem_flags(flags, false)?;
+ validate_host_ptr(host_ptr, flags)?;
+
+ let filtered_flags = filter_image_access_flags(flags);
+ // CL_IMAGE_FORMAT_NOT_SUPPORTED if there are no devices in context that support image_format.
+ c.devs
+ .iter()
+ .filter_map(|d| d.formats.get(format))
+ .filter_map(|f| f.get(&desc.image_type))
+ .find(|f| *f & filtered_flags == filtered_flags)
+ .ok_or(CL_IMAGE_FORMAT_NOT_SUPPORTED)?;
+
+ Ok(cl_mem::from_arc(Mem::new_image(
+ c,
+ desc.image_type,
+ flags,
+ format,
+ desc,
+ elem_size,
+ host_ptr,
+ )))
+}
+
+pub fn get_supported_image_formats(
+ context: cl_context,
+ flags: cl_mem_flags,
+ image_type: cl_mem_object_type,
+ num_entries: cl_uint,
+ image_formats: *mut cl_image_format,
+ num_image_formats: *mut cl_uint,
+) -> CLResult<()> {
+ let c = context.get_ref()?;
+
+ // CL_INVALID_VALUE if flags
+ validate_mem_flags(flags, true)?;
+
+ // or image_type are not valid
+ if !image_type_valid(image_type) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_VALUE ... if num_entries is 0 and image_formats is not NULL.
+ if num_entries == 0 && !image_formats.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let mut res = Vec::<cl_image_format>::new();
+ let filtered_flags = filter_image_access_flags(flags);
+ for dev in &c.devs {
+ for f in &dev.formats {
+ let s = f.1.get(&image_type).unwrap_or(&0);
+
+ if filtered_flags & s == filtered_flags {
+ res.push(*f.0);
+ }
+ }
+ }
+
+ res.sort();
+ res.dedup();
+
+ num_image_formats.write_checked(res.len() as cl_uint);
+ unsafe { image_formats.copy_checked(res.as_ptr(), res.len()) };
+
+ Ok(())
+}
+
+impl CLInfo<cl_sampler_info> for cl_sampler {
+ fn query(&self, q: cl_sampler_info) -> CLResult<Vec<u8>> {
+ let sampler = self.get_ref()?;
+ Ok(match q {
+ CL_SAMPLER_ADDRESSING_MODE => cl_prop::<cl_addressing_mode>(sampler.addressing_mode),
+ CL_SAMPLER_CONTEXT => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&sampler.context);
+ cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+ }
+ CL_SAMPLER_FILTER_MODE => cl_prop::<cl_filter_mode>(sampler.filter_mode),
+ CL_SAMPLER_NORMALIZED_COORDS => cl_prop::<bool>(sampler.normalized_coords),
+ CL_SAMPLER_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+pub fn create_sampler(
+ context: cl_context,
+ normalized_coords: cl_bool,
+ addressing_mode: cl_addressing_mode,
+ filter_mode: cl_filter_mode,
+) -> CLResult<cl_sampler> {
+ let c = context.get_arc()?;
+
+ // CL_INVALID_OPERATION if images are not supported by any device associated with context (i.e.
+ // CL_DEVICE_IMAGE_SUPPORT specified in the Device Queries table is CL_FALSE).
+ c.devs
+ .iter()
+ .find(|d| d.image_supported())
+ .ok_or(CL_INVALID_OPERATION)?;
+
+ // CL_INVALID_VALUE if addressing_mode, filter_mode, normalized_coords or a combination of these
+ // arguements are not valid.
+ validate_addressing_mode(addressing_mode)?;
+ validate_filter_mode(filter_mode)?;
+
+ let sampler = Sampler::new(
+ c,
+ check_cl_bool(normalized_coords).ok_or(CL_INVALID_VALUE)?,
+ addressing_mode,
+ filter_mode,
+ );
+ Ok(cl_sampler::from_arc(sampler))
+}
+
+pub fn enqueue_write_buffer(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_write: cl_bool,
+ offset: usize,
+ cb: usize,
+ ptr: *const ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<()> {
+ let q = command_queue.get_arc()?;
+ let b = buffer.get_arc()?;
+ let block = check_cl_bool(blocking_write).ok_or(CL_INVALID_VALUE)?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_VALUE if the region being read or written specified by (offset, size) is out of
+ // bounds or if ptr is a NULL value.
+ if offset + cb > b.size || ptr.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_CONTEXT if the context associated with command_queue and buffer are not the same
+ if b.context != q.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ // CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the read and write operations are blocking
+ // and the execution status of any of the events in event_wait_list is a negative integer value.
+ if block && evs.iter().any(|e| e.is_error()) {
+ return Err(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
+ }
+
+ // CL_INVALID_OPERATION if clEnqueueWriteBuffer is called on buffer which has been created with
+ // CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS.
+ if bit_check(b.flags, CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS) {
+ return Err(CL_INVALID_OPERATION);
+ }
+
+ create_and_queue(
+ q,
+ CL_COMMAND_WRITE_BUFFER,
+ evs,
+ event,
+ block,
+ Box::new(move |q, ctx| b.write_from_user(q, ctx, offset, ptr, cb)),
+ )
+
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if buffer is a sub-buffer object and offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
+}
+
+pub fn enqueue_read_buffer_rect(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_read: cl_bool,
+ buffer_origin: *const usize,
+ host_origin: *const usize,
+ region: *const usize,
+ mut buffer_row_pitch: usize,
+ mut buffer_slice_pitch: usize,
+ mut host_row_pitch: usize,
+ mut host_slice_pitch: usize,
+ ptr: *mut ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<()> {
+ let block = check_cl_bool(blocking_read).ok_or(CL_INVALID_VALUE)?;
+ let q = command_queue.get_arc()?;
+ let buf = buffer.get_arc()?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_OPERATION if clEnqueueReadBufferRect is called on buffer which has been created
+ // with CL_MEM_HOST_WRITE_ONLY or CL_MEM_HOST_NO_ACCESS.
+ if bit_check(buf.flags, CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS) {
+ return Err(CL_INVALID_OPERATION);
+ }
+
+ // CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the read and write operations are blocking
+ // and the execution status of any of the events in event_wait_list is a negative integer value.
+ if block && evs.iter().any(|e| e.is_error()) {
+ return Err(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
+ }
+
+ // CL_INVALID_VALUE if buffer_origin, host_origin, or region is NULL.
+ if buffer_origin.is_null() ||
+ host_origin.is_null() ||
+ region.is_null() ||
+ // CL_INVALID_VALUE if ptr is NULL.
+ ptr.is_null()
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let r = unsafe { CLVec::from_raw(region) };
+ let buf_ori = unsafe { CLVec::from_raw(buffer_origin) };
+ let host_ori = unsafe { CLVec::from_raw(host_origin) };
+
+ // CL_INVALID_VALUE if any region array element is 0.
+ if r.contains(&0) ||
+ // CL_INVALID_VALUE if buffer_row_pitch is not 0 and is less than region[0].
+ buffer_row_pitch != 0 && buffer_row_pitch < r[0] ||
+ // CL_INVALID_VALUE if host_row_pitch is not 0 and is less than region[0].
+ host_row_pitch != 0 && host_row_pitch < r[0]
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If buffer_row_pitch is 0, buffer_row_pitch is computed as region[0].
+ if buffer_row_pitch == 0 {
+ buffer_row_pitch = r[0];
+ }
+
+ // If host_row_pitch is 0, host_row_pitch is computed as region[0].
+ if host_row_pitch == 0 {
+ host_row_pitch = r[0];
+ }
+
+ // CL_INVALID_VALUE if buffer_slice_pitch is not 0 and is less than region[1] × buffer_row_pitch and not a multiple of buffer_row_pitch.
+ if buffer_slice_pitch != 0 && buffer_slice_pitch < r[1] * buffer_row_pitch && buffer_slice_pitch % buffer_row_pitch != 0 ||
+ // CL_INVALID_VALUE if host_slice_pitch is not 0 and is less than region[1] × host_row_pitch and not a multiple of host_row_pitch.
+ host_slice_pitch != 0 && host_slice_pitch < r[1] * host_row_pitch && host_slice_pitch % host_row_pitch != 0
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If buffer_slice_pitch is 0, buffer_slice_pitch is computed as region[1] × buffer_row_pitch.
+ if buffer_slice_pitch == 0 {
+ buffer_slice_pitch = r[1] * buffer_row_pitch;
+ }
+
+ // If host_slice_pitch is 0, host_slice_pitch is computed as region[1] × host_row_pitch.
+ if host_slice_pitch == 0 {
+ host_slice_pitch = r[1] * host_row_pitch
+ }
+
+ // CL_INVALID_VALUE if the region being read or written specified by (buffer_origin, region,
+ // buffer_row_pitch, buffer_slice_pitch) is out of bounds.
+ if !CLVec::is_in_bound(
+ r,
+ buf_ori,
+ [1, buffer_row_pitch, buffer_slice_pitch],
+ buf.size,
+ ) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_CONTEXT if the context associated with command_queue and buffer are not the same
+ if q.context != buf.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ create_and_queue(
+ q,
+ CL_COMMAND_READ_BUFFER_RECT,
+ evs,
+ event,
+ block,
+ Box::new(move |q, ctx| {
+ buf.read_to_user_rect(
+ ptr,
+ q,
+ ctx,
+ &r,
+ &buf_ori,
+ buffer_row_pitch,
+ buffer_slice_pitch,
+ &host_ori,
+ host_row_pitch,
+ host_slice_pitch,
+ )
+ }),
+ )
+
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if buffer is a sub-buffer object and offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
+}
+
+pub fn enqueue_write_buffer_rect(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_write: cl_bool,
+ buffer_origin: *const usize,
+ host_origin: *const usize,
+ region: *const usize,
+ mut buffer_row_pitch: usize,
+ mut buffer_slice_pitch: usize,
+ mut host_row_pitch: usize,
+ mut host_slice_pitch: usize,
+ ptr: *const ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<()> {
+ let block = check_cl_bool(blocking_write).ok_or(CL_INVALID_VALUE)?;
+ let q = command_queue.get_arc()?;
+ let buf = buffer.get_arc()?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_OPERATION if clEnqueueWriteBufferRect is called on buffer which has been created
+ // with CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS.
+ if bit_check(buf.flags, CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS) {
+ return Err(CL_INVALID_OPERATION);
+ }
+
+ // CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the read and write operations are blocking
+ // and the execution status of any of the events in event_wait_list is a negative integer value.
+ if block && evs.iter().any(|e| e.is_error()) {
+ return Err(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
+ }
+
+ // CL_INVALID_VALUE if buffer_origin, host_origin, or region is NULL.
+ if buffer_origin.is_null() ||
+ host_origin.is_null() ||
+ region.is_null() ||
+ // CL_INVALID_VALUE if ptr is NULL.
+ ptr.is_null()
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let r = unsafe { CLVec::from_raw(region) };
+ let buf_ori = unsafe { CLVec::from_raw(buffer_origin) };
+ let host_ori = unsafe { CLVec::from_raw(host_origin) };
+
+ // CL_INVALID_VALUE if any region array element is 0.
+ if r.contains(&0) ||
+ // CL_INVALID_VALUE if buffer_row_pitch is not 0 and is less than region[0].
+ buffer_row_pitch != 0 && buffer_row_pitch < r[0] ||
+ // CL_INVALID_VALUE if host_row_pitch is not 0 and is less than region[0].
+ host_row_pitch != 0 && host_row_pitch < r[0]
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If buffer_row_pitch is 0, buffer_row_pitch is computed as region[0].
+ if buffer_row_pitch == 0 {
+ buffer_row_pitch = r[0];
+ }
+
+ // If host_row_pitch is 0, host_row_pitch is computed as region[0].
+ if host_row_pitch == 0 {
+ host_row_pitch = r[0];
+ }
+
+ // CL_INVALID_VALUE if buffer_slice_pitch is not 0 and is less than region[1] × buffer_row_pitch and not a multiple of buffer_row_pitch.
+ if buffer_slice_pitch != 0 && buffer_slice_pitch < r[1] * buffer_row_pitch && buffer_slice_pitch % buffer_row_pitch != 0 ||
+ // CL_INVALID_VALUE if host_slice_pitch is not 0 and is less than region[1] × host_row_pitch and not a multiple of host_row_pitch.
+ host_slice_pitch != 0 && host_slice_pitch < r[1] * host_row_pitch && host_slice_pitch % host_row_pitch != 0
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If buffer_slice_pitch is 0, buffer_slice_pitch is computed as region[1] × buffer_row_pitch.
+ if buffer_slice_pitch == 0 {
+ buffer_slice_pitch = r[1] * buffer_row_pitch;
+ }
+
+ // If host_slice_pitch is 0, host_slice_pitch is computed as region[1] × host_row_pitch.
+ if host_slice_pitch == 0 {
+ host_slice_pitch = r[1] * host_row_pitch
+ }
+
+ // CL_INVALID_VALUE if the region being read or written specified by (buffer_origin, region,
+ // buffer_row_pitch, buffer_slice_pitch) is out of bounds.
+ if !CLVec::is_in_bound(
+ r,
+ buf_ori,
+ [1, buffer_row_pitch, buffer_slice_pitch],
+ buf.size,
+ ) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_CONTEXT if the context associated with command_queue and buffer are not the same
+ if q.context != buf.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ create_and_queue(
+ q,
+ CL_COMMAND_WRITE_BUFFER_RECT,
+ evs,
+ event,
+ block,
+ Box::new(move |q, ctx| {
+ buf.write_from_user_rect(
+ ptr,
+ q,
+ ctx,
+ &r,
+ &host_ori,
+ host_row_pitch,
+ host_slice_pitch,
+ &buf_ori,
+ buffer_row_pitch,
+ buffer_slice_pitch,
+ )
+ }),
+ )
+
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if buffer is a sub-buffer object and offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
+}
+
+pub fn enqueue_copy_buffer_rect(
+ command_queue: cl_command_queue,
+ src_buffer: cl_mem,
+ dst_buffer: cl_mem,
+ src_origin: *const usize,
+ dst_origin: *const usize,
+ region: *const usize,
+ mut src_row_pitch: usize,
+ mut src_slice_pitch: usize,
+ mut dst_row_pitch: usize,
+ mut dst_slice_pitch: usize,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<()> {
+ let q = command_queue.get_arc()?;
+ let src = src_buffer.get_arc()?;
+ let dst = dst_buffer.get_arc()?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_VALUE if src_origin, dst_origin, or region is NULL.
+ if src_origin.is_null() || dst_origin.is_null() || region.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let r = unsafe { CLVec::from_raw(region) };
+ let src_ori = unsafe { CLVec::from_raw(src_origin) };
+ let dst_ori = unsafe { CLVec::from_raw(dst_origin) };
+
+ // CL_INVALID_VALUE if any region array element is 0.
+ if r.contains(&0) ||
+ // CL_INVALID_VALUE if src_row_pitch is not 0 and is less than region[0].
+ src_row_pitch != 0 && src_row_pitch < r[0] ||
+ // CL_INVALID_VALUE if dst_row_pitch is not 0 and is less than region[0].
+ dst_row_pitch != 0 && dst_row_pitch < r[0]
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If src_row_pitch is 0, src_row_pitch is computed as region[0].
+ if src_row_pitch == 0 {
+ src_row_pitch = r[0];
+ }
+
+ // If dst_row_pitch is 0, dst_row_pitch is computed as region[0].
+ if dst_row_pitch == 0 {
+ dst_row_pitch = r[0];
+ }
+
+ // CL_INVALID_VALUE if src_slice_pitch is not 0 and is less than region[1] × src_row_pitch
+ if src_slice_pitch != 0 && src_slice_pitch < r[1] * src_row_pitch ||
+ // CL_INVALID_VALUE if dst_slice_pitch is not 0 and is less than region[1] × dst_row_pitch
+ dst_slice_pitch != 0 && dst_slice_pitch < r[1] * dst_row_pitch ||
+ // if src_slice_pitch is not 0 and is not a multiple of src_row_pitch.
+ src_slice_pitch != 0 && src_slice_pitch % src_row_pitch != 0 ||
+ // if dst_slice_pitch is not 0 and is not a multiple of dst_row_pitch.
+ dst_slice_pitch != 0 && dst_slice_pitch % dst_row_pitch != 0
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // If src_slice_pitch is 0, src_slice_pitch is computed as region[1] × src_row_pitch.
+ if src_slice_pitch == 0 {
+ src_slice_pitch = r[1] * src_row_pitch;
+ }
+
+ // If dst_slice_pitch is 0, dst_slice_pitch is computed as region[1] × dst_row_pitch.
+ if dst_slice_pitch == 0 {
+ dst_slice_pitch = r[1] * dst_row_pitch;
+ }
+
+ // CL_INVALID_VALUE if src_buffer and dst_buffer are the same buffer object and src_slice_pitch
+ // is not equal to dst_slice_pitch and src_row_pitch is not equal to dst_row_pitch.
+ if src_buffer == dst_buffer
+ && src_slice_pitch != dst_slice_pitch
+ && src_row_pitch != dst_row_pitch
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_VALUE if (src_origin, region, src_row_pitch, src_slice_pitch) or (dst_origin,
+ // region, dst_row_pitch, dst_slice_pitch) require accessing elements outside the src_buffer
+ // and dst_buffer buffer objects respectively.
+ if !CLVec::is_in_bound(r, src_ori, [1, src_row_pitch, src_slice_pitch], src.size)
+ || !CLVec::is_in_bound(r, dst_ori, [1, dst_row_pitch, dst_slice_pitch], dst.size)
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_MEM_COPY_OVERLAP if src_buffer and dst_buffer are the same buffer or sub-buffer object and
+ // the source and destination regions overlap or if src_buffer and dst_buffer are different
+ // sub-buffers of the same associated buffer object and they overlap.
+ if src.has_same_parent(&dst)
+ && check_copy_overlap(
+ &src_ori,
+ src.offset,
+ &dst_ori,
+ dst.offset,
+ &r,
+ src_row_pitch,
+ src_slice_pitch,
+ )
+ {
+ return Err(CL_MEM_COPY_OVERLAP);
+ }
+
+ // CL_INVALID_CONTEXT if the context associated with command_queue, src_buffer and dst_buffer
+ // are not the same
+ if src.context != q.context || dst.context != q.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ create_and_queue(
+ q,
+ CL_COMMAND_COPY_BUFFER_RECT,
+ evs,
+ event,
+ false,
+ Box::new(move |q, ctx| {
+ src.copy_to(
+ &dst,
+ q,
+ ctx,
+ &r,
+ &src_ori,
+ src_row_pitch,
+ src_slice_pitch,
+ &dst_ori,
+ dst_row_pitch,
+ dst_slice_pitch,
+ )
+ }),
+ )
+
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if src_buffer is a sub-buffer object and offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue.
+}
+
+pub fn enqueue_map_buffer(
+ command_queue: cl_command_queue,
+ buffer: cl_mem,
+ blocking_map: cl_bool,
+ map_flags: cl_map_flags,
+ offset: usize,
+ size: usize,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<*mut c_void> {
+ let q = command_queue.get_arc()?;
+ let b = buffer.get_arc()?;
+ let block = check_cl_bool(blocking_map).ok_or(CL_INVALID_VALUE)?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_VALUE if region being mapped given by (offset, size) is out of bounds or if size
+ // is 0
+ if offset + size > b.size || size == 0 {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_VALUE ... if values specified in map_flags are not valid.
+ let valid_flags =
+ cl_bitfield::from(CL_MAP_READ | CL_MAP_WRITE | CL_MAP_WRITE_INVALIDATE_REGION);
+ let read_write_group = cl_bitfield::from(CL_MAP_READ | CL_MAP_WRITE);
+ let invalidate_group = cl_bitfield::from(CL_MAP_WRITE_INVALIDATE_REGION);
+
+ if (map_flags & !valid_flags != 0)
+ || ((map_flags & read_write_group != 0) && (map_flags & invalidate_group != 0))
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_OPERATION if buffer has been created with CL_MEM_HOST_WRITE_ONLY or
+ // CL_MEM_HOST_NO_ACCESS and CL_MAP_READ is set in map_flags
+ if bit_check(b.flags, CL_MEM_HOST_WRITE_ONLY | CL_MEM_HOST_NO_ACCESS) &&
+ bit_check(map_flags, CL_MAP_READ) ||
+ // or if buffer has been created with CL_MEM_HOST_READ_ONLY or CL_MEM_HOST_NO_ACCESS and
+ // CL_MAP_WRITE or CL_MAP_WRITE_INVALIDATE_REGION is set in map_flags.
+ bit_check(b.flags, CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS) &&
+ bit_check(map_flags, CL_MAP_WRITE | CL_MAP_WRITE_INVALIDATE_REGION)
+ {
+ return Err(CL_INVALID_OPERATION);
+ }
+
+ // CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the map operation is blocking and the
+ // execution status of any of the events in event_wait_list is a negative integer value.
+ if block && evs.iter().any(|e| e.is_error()) {
+ return Err(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
+ }
+
+ // CL_INVALID_CONTEXT if context associated with command_queue and buffer are not the same
+ if b.context != q.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ create_and_queue(
+ q.clone(),
+ CL_COMMAND_MAP_BUFFER,
+ evs,
+ event,
+ block,
+ // we don't really have anything to do here?
+ Box::new(|_, _| Ok(())),
+ )?;
+
+ Ok(b.map(&q, offset, size, block))
+ // TODO
+ // CL_MISALIGNED_SUB_BUFFER_OFFSET if buffer is a sub-buffer object and offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for the device associated with queue. This error code is missing before version 1.1.
+ // CL_MAP_FAILURE if there is a failure to map the requested region into the host address space. This error cannot occur for buffer objects created with CL_MEM_USE_HOST_PTR or CL_MEM_ALLOC_HOST_PTR.
+ // CL_INVALID_OPERATION if mapping would lead to overlapping regions being mapped for writing.
+}
+
+pub fn enqueue_unmap_mem_object(
+ command_queue: cl_command_queue,
+ memobj: cl_mem,
+ mapped_ptr: *mut ::std::os::raw::c_void,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+ event: *mut cl_event,
+) -> CLResult<()> {
+ let q = command_queue.get_arc()?;
+ let m = memobj.get_arc()?;
+ let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;
+
+ // CL_INVALID_CONTEXT if context associated with command_queue and memobj are not the same
+ if q.context != m.context {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ // CL_INVALID_VALUE if mapped_ptr is not a valid pointer returned by clEnqueueMapBuffer or
+ // clEnqueueMapImage for memobj.
+ if !m.is_mapped_ptr(mapped_ptr) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ create_and_queue(
+ q,
+ CL_COMMAND_UNMAP_MEM_OBJECT,
+ evs,
+ event,
+ false,
+ Box::new(move |q, _| {
+ m.unmap(q, mapped_ptr);
+ Ok(())
+ }),
+ )
+}
--- /dev/null
+mod context;
+mod device;
+mod event;
+pub mod icd;
+mod kernel;
+mod memory;
+mod platform;
+mod program;
+mod queue;
+pub(super) mod types;
+pub(super) mod util;
--- /dev/null
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::CLResult;
+use crate::api::icd::DISPATCH;
+use crate::api::util::*;
+use crate::core::version::*;
+
+use self::rusticl_opencl_gen::*;
+
+use self::mesa_rust_util::ptr::*;
+
+#[repr(C)]
+#[allow(non_camel_case_types)]
+pub struct _cl_platform_id {
+ dispatch: &'static cl_icd_dispatch,
+ extensions: [cl_name_version; 1],
+}
+
+impl CLInfo<cl_platform_info> for cl_platform_id {
+ fn query(&self, q: cl_platform_info) -> CLResult<Vec<u8>> {
+ let p = self.get_ref()?;
+ Ok(match q {
+ CL_PLATFORM_EXTENSIONS => cl_prop("cl_khr_icd"),
+ CL_PLATFORM_EXTENSIONS_WITH_VERSION => {
+ cl_prop::<Vec<cl_name_version>>(p.extensions.to_vec())
+ }
+ CL_PLATFORM_HOST_TIMER_RESOLUTION => cl_prop::<cl_ulong>(0),
+ CL_PLATFORM_ICD_SUFFIX_KHR => cl_prop("MESA"),
+ CL_PLATFORM_NAME => cl_prop("rusticl"),
+ CL_PLATFORM_NUMERIC_VERSION => cl_prop::<cl_version>(CLVersion::Cl3_0 as u32),
+ CL_PLATFORM_PROFILE => cl_prop("FULL_PROFILE"),
+ CL_PLATFORM_VENDOR => cl_prop("Mesa/X.org"),
+ // OpenCL<space><major_version.minor_version><space><platform-specific information>
+ CL_PLATFORM_VERSION => cl_prop("OpenCL 3.0 "),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+static PLATFORM: _cl_platform_id = _cl_platform_id {
+ dispatch: &DISPATCH,
+ extensions: [mk_cl_version_ext(1, 0, 0, "cl_khr_icd")],
+};
+
+pub fn get_platform() -> cl_platform_id {
+ &PLATFORM as *const crate::api::platform::_cl_platform_id
+ as *mut self::rusticl_opencl_gen::_cl_platform_id
+}
+
+pub trait GetPlatformRef {
+ fn get_ref(&self) -> CLResult<&'static _cl_platform_id>;
+}
+
+impl GetPlatformRef for cl_platform_id {
+ fn get_ref(&self) -> CLResult<&'static _cl_platform_id> {
+ if !self.is_null() && *self == get_platform() {
+ Ok(&PLATFORM)
+ } else {
+ Err(CL_INVALID_PLATFORM)
+ }
+ }
+}
+
+pub fn get_platform_ids(
+ num_entries: cl_uint,
+ platforms: *mut cl_platform_id,
+ num_platforms: *mut cl_uint,
+) -> CLResult<()> {
+ // CL_INVALID_VALUE if num_entries is equal to zero and platforms is not NULL
+ if num_entries == 0 && !platforms.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // or if both num_platforms and platforms are NULL."
+ if num_platforms.is_null() && platforms.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // platforms returns a list of OpenCL platforms available for access through the Khronos ICD Loader.
+ // The cl_platform_id values returned in platforms are ICD compatible and can be used to identify a
+ // specific OpenCL platform. If the platforms argument is NULL, then this argument is ignored. The
+ // number of OpenCL platforms returned is the minimum of the value specified by num_entries or the
+ // number of OpenCL platforms available.
+ platforms.write_checked(get_platform());
+
+ // num_platforms returns the number of OpenCL platforms available. If num_platforms is NULL, then
+ // this argument is ignored.
+ num_platforms.write_checked(1);
+
+ Ok(())
+}
+
+#[test]
+fn test_get_platform_info() {
+ let mut s: usize = 0;
+ let mut r = get_platform_info(
+ ptr::null(),
+ CL_PLATFORM_EXTENSIONS,
+ 0,
+ ptr::null_mut(),
+ &mut s,
+ );
+ assert!(r.is_ok());
+ assert!(s > 0);
+
+ let mut v: Vec<u8> = vec![0; s];
+ r = get_platform_info(
+ ptr::null(),
+ CL_PLATFORM_EXTENSIONS,
+ s,
+ v.as_mut_ptr().cast(),
+ &mut s,
+ );
+
+ assert!(r.is_ok());
+ assert_eq!(s, v.len());
+ assert!(!v[0..s - 2].contains(&0));
+ assert_eq!(v[s - 1], 0);
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::types::*;
+use crate::api::util::*;
+use crate::core::device::*;
+use crate::core::program::*;
+
+use self::mesa_rust::compiler::clc::*;
+use self::mesa_rust_util::string::*;
+use self::rusticl_opencl_gen::*;
+
+use std::ffi::CStr;
+use std::ffi::CString;
+use std::os::raw::c_char;
+use std::ptr;
+use std::slice;
+use std::sync::Arc;
+
+impl CLInfo<cl_program_info> for cl_program {
+ fn query(&self, q: cl_program_info) -> CLResult<Vec<u8>> {
+ let prog = self.get_ref()?;
+ Ok(match q {
+ CL_PROGRAM_CONTEXT => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&prog.context);
+ cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+ }
+ CL_PROGRAM_DEVICES => {
+ cl_prop::<&Vec<cl_device_id>>(
+ &prog
+ .devs
+ .iter()
+ .map(|d| {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ cl_device_id::from_ptr(Arc::as_ptr(d))
+ })
+ .collect(),
+ )
+ }
+ CL_PROGRAM_NUM_DEVICES => cl_prop::<cl_uint>(prog.devs.len() as cl_uint),
+ CL_PROGRAM_NUM_KERNELS => cl_prop::<usize>(prog.kernels().len()),
+ CL_PROGRAM_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ CL_PROGRAM_SOURCE => cl_prop::<&CStr>(prog.src.as_c_str()),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+impl CLInfoObj<cl_program_build_info, cl_device_id> for cl_program {
+ fn query(&self, d: cl_device_id, q: cl_program_build_info) -> CLResult<Vec<u8>> {
+ let prog = self.get_ref()?;
+ let dev = d.get_arc()?;
+ Ok(match q {
+ CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE => cl_prop::<usize>(0),
+ CL_PROGRAM_BUILD_LOG => cl_prop::<String>(prog.log(&dev)),
+ CL_PROGRAM_BUILD_OPTIONS => cl_prop::<String>(prog.options(&dev)),
+ CL_PROGRAM_BUILD_STATUS => cl_prop::<cl_build_status>(prog.status(&dev)),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+fn validate_devices(
+ device_list: *const cl_device_id,
+ num_devices: cl_uint,
+ default: &[Arc<Device>],
+) -> CLResult<Vec<Arc<Device>>> {
+ let mut devs = cl_device_id::get_arc_vec_from_arr(device_list, num_devices)?;
+
+ // If device_list is a NULL value, the compile is performed for all devices associated with
+ // program.
+ if devs.is_empty() {
+ devs = default.to_vec();
+ }
+
+ Ok(devs)
+}
+
+fn call_cb(
+ pfn_notify: Option<ProgramCB>,
+ program: cl_program,
+ user_data: *mut ::std::os::raw::c_void,
+) {
+ if let Some(cb) = pfn_notify {
+ unsafe { cb(program, user_data) };
+ }
+}
+
+pub fn create_program_with_source(
+ context: cl_context,
+ count: cl_uint,
+ strings: *mut *const c_char,
+ lengths: *const usize,
+) -> CLResult<cl_program> {
+ let c = context.get_arc()?;
+
+ // CL_INVALID_VALUE if count is zero or if strings ...
+ if count == 0 || strings.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // ... or any entry in strings is NULL.
+ let srcs = unsafe { slice::from_raw_parts(strings, count as usize) };
+ if srcs.contains(&ptr::null()) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let mut source = String::new();
+ // we don't want encoding or any other problems with the source to prevent compilations, so
+ // just use CString::from_vec_unchecked and to_string_lossy
+ for i in 0..count as usize {
+ unsafe {
+ if lengths.is_null() || *lengths.add(i) == 0 {
+ source.push_str(&CStr::from_ptr(*strings.add(i)).to_string_lossy());
+ } else {
+ let l = *lengths.add(i);
+ let arr = slice::from_raw_parts(*strings.add(i).cast(), l);
+ source.push_str(&CString::from_vec_unchecked(arr.to_vec()).to_string_lossy());
+ }
+ }
+ }
+
+ Ok(cl_program::from_arc(Program::new(
+ &c,
+ &c.devs,
+ CString::new(source).map_err(|_| CL_INVALID_VALUE)?,
+ )))
+}
+
+pub fn build_program(
+ program: cl_program,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ options: *const c_char,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<()> {
+ let mut res = true;
+ let p = program.get_ref()?;
+ let devs = validate_devices(device_list, num_devices, &p.devs)?;
+
+ check_cb(&pfn_notify, user_data)?;
+
+ // CL_BUILD_PROGRAM_FAILURE if there is a failure to build the program executable. This error
+ // will be returned if clBuildProgram does not return until the build has completed.
+ for dev in devs {
+ res &= p.compile(&dev, c_string_to_string(options), &Vec::new());
+ }
+
+ call_cb(pfn_notify, program, user_data);
+
+ // TODO link
+
+ //• CL_INVALID_BINARY if program is created with clCreateProgramWithBinary and devices listed in device_list do not have a valid program binary loaded.
+ //• CL_INVALID_BUILD_OPTIONS if the build options specified by options are invalid.
+ //• CL_INVALID_OPERATION if the build of a program executable for any of the devices listed in device_list by a previous call to clBuildProgram for program has not completed.
+ //• CL_INVALID_OPERATION if there are kernel objects attached to program.
+ //• CL_INVALID_OPERATION if program was not created with clCreateProgramWithSource, clCreateProgramWithIL or clCreateProgramWithBinary.
+
+ if res {
+ Ok(())
+ } else {
+ Err(CL_BUILD_PROGRAM_FAILURE)
+ }
+}
+
+pub fn compile_program(
+ program: cl_program,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ options: *const c_char,
+ num_input_headers: cl_uint,
+ input_headers: *const cl_program,
+ header_include_names: *mut *const c_char,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<()> {
+ let mut res = true;
+ let p = program.get_ref()?;
+ let devs = validate_devices(device_list, num_devices, &p.devs)?;
+
+ check_cb(&pfn_notify, user_data)?;
+
+ // CL_INVALID_VALUE if num_input_headers is zero and header_include_names or input_headers are
+ // not NULL or if num_input_headers is not zero and header_include_names or input_headers are
+ // NULL.
+ if num_input_headers == 0 && (!header_include_names.is_null() || !input_headers.is_null())
+ || num_input_headers != 0 && (header_include_names.is_null() || input_headers.is_null())
+ {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ let mut headers = Vec::new();
+ for h in 0..num_input_headers as usize {
+ unsafe {
+ headers.push(spirv::CLCHeader {
+ name: CStr::from_ptr(*header_include_names.add(h)).to_owned(),
+ source: &(*input_headers.add(h)).get_ref()?.src,
+ });
+ }
+ }
+
+ // CL_COMPILE_PROGRAM_FAILURE if there is a failure to compile the program source. This error
+ // will be returned if clCompileProgram does not return until the compile has completed.
+ for dev in devs {
+ res &= p.compile(&dev, c_string_to_string(options), &headers);
+ }
+
+ call_cb(pfn_notify, program, user_data);
+
+ // CL_INVALID_OPERATION if program has no source or IL available, i.e. it has not been created with clCreateProgramWithSource or clCreateProgramWithIL.
+ // • CL_INVALID_COMPILER_OPTIONS if the compiler options specified by options are invalid.
+ // • CL_INVALID_OPERATION if the compilation or build of a program executable for any of the devices listed in device_list by a previous call to clCompileProgram or clBuildProgram for program has not completed.
+ // • CL_INVALID_OPERATION if there are kernel objects attached to program.
+
+ if res {
+ Ok(())
+ } else {
+ Err(CL_COMPILE_PROGRAM_FAILURE)
+ }
+}
+
+pub fn link_program(
+ context: cl_context,
+ num_devices: cl_uint,
+ device_list: *const cl_device_id,
+ _options: *const ::std::os::raw::c_char,
+ num_input_programs: cl_uint,
+ input_programs: *const cl_program,
+ pfn_notify: Option<ProgramCB>,
+ user_data: *mut ::std::os::raw::c_void,
+) -> CLResult<(cl_program, cl_int)> {
+ let c = context.get_arc()?;
+ let devs = validate_devices(device_list, num_devices, &c.devs)?;
+ let progs = cl_program::get_arc_vec_from_arr(input_programs, num_input_programs)?;
+
+ check_cb(&pfn_notify, user_data)?;
+
+ // CL_INVALID_VALUE if num_input_programs is zero and input_programs is NULL
+ if progs.is_empty() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_DEVICE if any device in device_list is not in the list of devices associated with
+ // context.
+ if !devs.iter().all(|d| c.devs.contains(d)) {
+ return Err(CL_INVALID_DEVICE);
+ }
+
+ // CL_INVALID_OPERATION if the compilation or build of a program executable for any of the
+ // devices listed in device_list by a previous call to clCompileProgram or clBuildProgram for
+ // program has not completed.
+ for d in &devs {
+ if progs
+ .iter()
+ .map(|p| p.status(d))
+ .any(|s| s != CL_BUILD_SUCCESS as cl_build_status)
+ {
+ return Err(CL_INVALID_OPERATION);
+ }
+ }
+
+ // CL_LINK_PROGRAM_FAILURE if there is a failure to link the compiled binaries and/or libraries.
+ let res = Program::link(c, &devs, &progs);
+ let code = if devs
+ .iter()
+ .map(|d| res.status(d))
+ .all(|s| s == CL_BUILD_SUCCESS as cl_build_status)
+ {
+ CL_SUCCESS as cl_int
+ } else {
+ CL_LINK_PROGRAM_FAILURE
+ };
+
+ let res = cl_program::from_arc(res);
+
+ call_cb(pfn_notify, res, user_data);
+ Ok((res, code))
+
+ //• CL_INVALID_LINKER_OPTIONS if the linker options specified by options are invalid.
+ //• CL_INVALID_OPERATION if the rules for devices containing compiled binaries or libraries as described in input_programs argument above are not followed.
+}
--- /dev/null
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::util::*;
+use crate::core::queue::*;
+
+use self::rusticl_opencl_gen::*;
+
+use std::sync::Arc;
+
+impl CLInfo<cl_command_queue_info> for cl_command_queue {
+ fn query(&self, q: cl_command_queue_info) -> CLResult<Vec<u8>> {
+ let queue = self.get_ref()?;
+ Ok(match q {
+ CL_QUEUE_CONTEXT => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&queue.context);
+ cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+ }
+ CL_QUEUE_DEVICE => {
+ // Note we use as_ptr here which doesn't increase the reference count.
+ let ptr = Arc::as_ptr(&queue.device);
+ cl_prop::<cl_device_id>(cl_device_id::from_ptr(ptr))
+ }
+ CL_QUEUE_PROPERTIES => cl_prop::<cl_command_queue_properties>(queue.props),
+ CL_QUEUE_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+ // CL_INVALID_VALUE if param_name is not one of the supported values
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
+
+fn valid_command_queue_properties(properties: cl_command_queue_properties) -> bool {
+ let valid_flags =
+ cl_bitfield::from(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE);
+ properties & !valid_flags == 0
+}
+
+fn supported_command_queue_properties(properties: cl_command_queue_properties) -> bool {
+ let valid_flags = cl_bitfield::from(CL_QUEUE_PROFILING_ENABLE);
+ properties & !valid_flags == 0
+}
+
+pub fn create_command_queue(
+ context: cl_context,
+ device: cl_device_id,
+ properties: cl_command_queue_properties,
+) -> CLResult<cl_command_queue> {
+ // CL_INVALID_CONTEXT if context is not a valid context.
+ let c = context.get_arc()?;
+
+ // CL_INVALID_DEVICE if device is not a valid device
+ let d = device.get_arc()?;
+
+ // ... or is not associated with context.
+ if !c.devs.contains(&d) {
+ return Err(CL_INVALID_DEVICE);
+ }
+
+ // CL_INVALID_VALUE if values specified in properties are not valid.
+ if !valid_command_queue_properties(properties) {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // CL_INVALID_QUEUE_PROPERTIES if values specified in properties are valid but are not supported by the device.
+ if !supported_command_queue_properties(properties) {
+ return Err(CL_INVALID_QUEUE_PROPERTIES);
+ }
+
+ Ok(cl_command_queue::from_arc(Queue::new(c, d, properties)?))
+}
+
+pub fn finish_queue(command_queue: cl_command_queue) -> CLResult<()> {
+ // CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
+ command_queue.get_ref()?;
+ Ok(())
+}
--- /dev/null
+extern crate rusticl_opencl_gen;
+
+use self::rusticl_opencl_gen::*;
+
+#[macro_export]
+macro_rules! cl_closure {
+ (|$obj:ident| $cb:ident($($arg:ident$(,)?)*)) => {
+ Box::new(
+ unsafe {
+ move|$obj| $cb.unwrap()($($arg,)*)
+ }
+ )
+ }
+}
+
+macro_rules! cl_callback {
+ ($cb:ident {
+ $($p:ident : $ty:ty,)*
+ }) => {
+ #[allow(dead_code)]
+ pub type $cb = unsafe extern "C" fn(
+ $($p: $ty,)*
+ );
+ }
+}
+
+cl_callback!(
+ CreateContextCB {
+ errinfo: *const ::std::os::raw::c_char,
+ private_info: *const ::std::ffi::c_void,
+ cb: usize,
+ user_data: *mut ::std::ffi::c_void,
+ }
+);
+
+cl_callback!(
+ EventCB {
+ event: cl_event,
+ event_command_status: cl_int,
+ user_data: *mut ::std::os::raw::c_void,
+ }
+);
+
+cl_callback!(
+ MemCB {
+ memobj: cl_mem,
+ user_data: *mut ::std::os::raw::c_void,
+ }
+);
+
+cl_callback!(
+ ProgramCB {
+ program: cl_program,
+ user_data: *mut ::std::os::raw::c_void,
+ }
+);
+
+// a lot of APIs use 3 component vectors passed as C arrays
+#[derive(Clone, Copy, PartialEq, Eq)]
+pub struct CLVec<T> {
+ vals: [T; 3],
+}
+
+impl<T: Copy> CLVec<T> {
+ /// # Safety
+ ///
+ /// This function is intended for use around OpenCL vectors of size 3.
+ /// Most commonly for `origin` and `region` API arguments.
+ ///
+ /// Using it for anything else is undefined.
+ pub unsafe fn from_raw(v: *const T) -> Self {
+ Self { vals: *v.cast() }
+ }
+}
+
+impl CLVec<usize> {
+ pub fn is_in_bound(base: Self, offset: Self, pitch: [usize; 3], size: usize) -> bool {
+ (base + offset - [1, 1, 1]) * pitch < size
+ }
+}
+
+impl<T: Default + Copy> Default for CLVec<T> {
+ fn default() -> Self {
+ Self {
+ vals: [T::default(); 3],
+ }
+ }
+}
+
+// provides a ton of functions
+impl<T> std::ops::Deref for CLVec<T> {
+ type Target = [T; 3];
+
+ fn deref(&self) -> &Self::Target {
+ &self.vals
+ }
+}
+
+impl<T: Copy + std::ops::Add<Output = T>> std::ops::Add for CLVec<T> {
+ type Output = Self;
+
+ fn add(self, other: Self) -> Self {
+ self + other.vals
+ }
+}
+
+impl<T: Copy + std::ops::Add<Output = T>> std::ops::Add<[T; 3]> for CLVec<T> {
+ type Output = Self;
+
+ fn add(self, other: [T; 3]) -> Self {
+ Self {
+ vals: [self[0] + other[0], self[1] + other[1], self[2] + other[2]],
+ }
+ }
+}
+
+impl<T: Copy + std::ops::Sub<Output = T>> std::ops::Sub<[T; 3]> for CLVec<T> {
+ type Output = Self;
+
+ fn sub(self, other: [T; 3]) -> Self {
+ Self {
+ vals: [self[0] - other[0], self[1] - other[1], self[2] - other[2]],
+ }
+ }
+}
+
+impl<T> std::ops::Mul for CLVec<T>
+where
+ T: Copy + std::ops::Mul<Output = T> + std::ops::Add<Output = T>,
+{
+ type Output = T;
+
+ fn mul(self, other: Self) -> T {
+ self * other.vals
+ }
+}
+
+impl<T> std::ops::Mul<[T; 3]> for CLVec<T>
+where
+ T: Copy + std::ops::Mul<Output = T> + std::ops::Add<Output = T>,
+{
+ type Output = T;
+
+ fn mul(self, other: [T; 3]) -> T {
+ self[0] * other[0] + self[1] * other[1] + self[2] * other[2]
+ }
+}
--- /dev/null
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::CLResult;
+use crate::api::types::*;
+use crate::core::event::*;
+use crate::core::queue::*;
+
+use self::mesa_rust_util::ptr::CheckedPtr;
+use self::rusticl_opencl_gen::*;
+
+use std::cmp;
+use std::convert::TryInto;
+use std::ffi::CStr;
+use std::ffi::CString;
+use std::mem::size_of;
+use std::ops::BitAnd;
+use std::os::raw::c_void;
+use std::slice;
+use std::sync::Arc;
+
+pub trait CLInfo<I> {
+ fn query(&self, q: I) -> CLResult<Vec<u8>>;
+
+ fn get_info(
+ &self,
+ param_name: I,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+ ) -> CLResult<()> {
+ let d = self.query(param_name)?;
+ let size: usize = d.len();
+
+ // CL_INVALID_VALUE [...] if size in bytes specified by param_value_size is < size of return
+ // type as specified in the Context Attributes table and param_value is not a NULL value.
+ if param_value_size < size && !param_value.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // param_value_size_ret returns the actual size in bytes of data being queried by param_name.
+ // If param_value_size_ret is NULL, it is ignored.
+ param_value_size_ret.write_checked(size);
+
+ // param_value is a pointer to memory where the appropriate result being queried is returned.
+ // If param_value is NULL, it is ignored.
+ unsafe {
+ param_value.copy_checked(d.as_ptr().cast(), size);
+ }
+
+ Ok(())
+ }
+}
+
+pub trait CLInfoObj<I, O> {
+ fn query(&self, o: O, q: I) -> CLResult<Vec<u8>>;
+
+ fn get_info_obj(
+ &self,
+ obj: O,
+ param_name: I,
+ param_value_size: usize,
+ param_value: *mut ::std::os::raw::c_void,
+ param_value_size_ret: *mut usize,
+ ) -> CLResult<()> {
+ let d = self.query(obj, param_name)?;
+ let size: usize = d.len();
+
+ // CL_INVALID_VALUE [...] if size in bytes specified by param_value_size is < size of return
+ // type as specified in the Context Attributes table and param_value is not a NULL value.
+ if param_value_size < size && !param_value.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ // param_value_size_ret returns the actual size in bytes of data being queried by param_name.
+ // If param_value_size_ret is NULL, it is ignored.
+ param_value_size_ret.write_checked(size);
+
+ // param_value is a pointer to memory where the appropriate result being queried is returned.
+ // If param_value is NULL, it is ignored.
+ unsafe {
+ param_value.copy_checked(d.as_ptr().cast(), size);
+ }
+
+ Ok(())
+ }
+}
+
+pub trait CLProp {
+ fn cl_vec(&self) -> Vec<u8>;
+}
+
+macro_rules! cl_prop_for_type {
+ ($ty: ty) => {
+ impl CLProp for $ty {
+ fn cl_vec(&self) -> Vec<u8> {
+ self.to_ne_bytes().to_vec()
+ }
+ }
+ };
+}
+
+macro_rules! cl_prop_for_struct {
+ ($ty: ty) => {
+ impl CLProp for $ty {
+ fn cl_vec(&self) -> Vec<u8> {
+ unsafe { slice::from_raw_parts((self as *const Self).cast(), size_of::<Self>()) }
+ .to_vec()
+ }
+ }
+ };
+}
+
+cl_prop_for_type!(cl_char);
+cl_prop_for_type!(cl_int);
+cl_prop_for_type!(cl_uint);
+cl_prop_for_type!(cl_ulong);
+cl_prop_for_type!(isize);
+cl_prop_for_type!(usize);
+
+cl_prop_for_struct!(cl_image_format);
+cl_prop_for_struct!(cl_name_version);
+
+impl CLProp for bool {
+ fn cl_vec(&self) -> Vec<u8> {
+ cl_prop::<cl_bool>(if *self { CL_TRUE } else { CL_FALSE })
+ }
+}
+
+impl CLProp for String {
+ fn cl_vec(&self) -> Vec<u8> {
+ let mut c = self.clone();
+ c.push('\0');
+ c.into_bytes()
+ }
+}
+
+impl CLProp for &str {
+ fn cl_vec(&self) -> Vec<u8> {
+ CString::new(*self)
+ .unwrap_or_default()
+ .into_bytes_with_nul()
+ }
+}
+
+impl CLProp for &CStr {
+ fn cl_vec(&self) -> Vec<u8> {
+ self.to_bytes_with_nul().to_vec()
+ }
+}
+
+impl<T> CLProp for Vec<T>
+where
+ T: CLProp,
+{
+ fn cl_vec(&self) -> Vec<u8> {
+ let mut res: Vec<u8> = Vec::new();
+ for i in self {
+ res.append(&mut i.cl_vec())
+ }
+ res
+ }
+}
+
+impl<T> CLProp for &Vec<T>
+where
+ T: CLProp,
+{
+ fn cl_vec(&self) -> Vec<u8> {
+ let mut res: Vec<u8> = Vec::new();
+ for i in *self {
+ res.append(&mut i.cl_vec())
+ }
+ res
+ }
+}
+
+impl<T> CLProp for *const T {
+ fn cl_vec(&self) -> Vec<u8> {
+ (*self as usize).cl_vec()
+ }
+}
+
+impl<T> CLProp for *mut T {
+ fn cl_vec(&self) -> Vec<u8> {
+ (*self as usize).cl_vec()
+ }
+}
+
+pub fn cl_prop<T: CLProp>(v: T) -> Vec<u8> {
+ v.cl_vec()
+}
+
+const CL_DEVICE_TYPES: u32 = CL_DEVICE_TYPE_ACCELERATOR
+ | CL_DEVICE_TYPE_CPU
+ | CL_DEVICE_TYPE_GPU
+ | CL_DEVICE_TYPE_CUSTOM
+ | CL_DEVICE_TYPE_DEFAULT;
+
+pub fn check_cl_device_type(val: cl_device_type) -> CLResult<()> {
+ let v: u32 = val.try_into().or(Err(CL_INVALID_DEVICE_TYPE))?;
+ if v == CL_DEVICE_TYPE_ALL || v & CL_DEVICE_TYPES == v {
+ return Ok(());
+ }
+ Err(CL_INVALID_DEVICE_TYPE)
+}
+
+pub const CL_IMAGE_TYPES: [cl_mem_object_type; 6] = [
+ CL_MEM_OBJECT_IMAGE1D,
+ CL_MEM_OBJECT_IMAGE2D,
+ CL_MEM_OBJECT_IMAGE3D,
+ CL_MEM_OBJECT_IMAGE1D_ARRAY,
+ CL_MEM_OBJECT_IMAGE2D_ARRAY,
+ CL_MEM_OBJECT_IMAGE1D_BUFFER,
+];
+
+pub const fn cl_image_format(
+ order: cl_channel_order,
+ data_type: cl_channel_type,
+) -> cl_image_format {
+ cl_image_format {
+ image_channel_order: order,
+ image_channel_data_type: data_type,
+ }
+}
+
+pub fn check_cl_bool<T: PartialEq + TryInto<cl_uint>>(val: T) -> Option<bool> {
+ let c: u32 = val.try_into().ok()?;
+ if c != CL_TRUE && c != CL_FALSE {
+ return None;
+ }
+ Some(c == CL_TRUE)
+}
+
+pub fn event_list_from_cl(
+ q: &Arc<Queue>,
+ num_events_in_wait_list: cl_uint,
+ event_wait_list: *const cl_event,
+) -> CLResult<Vec<Arc<Event>>> {
+ // CL_INVALID_EVENT_WAIT_LIST if event_wait_list is NULL and num_events_in_wait_list > 0, or
+ // event_wait_list is not NULL and num_events_in_wait_list is 0, or if event objects in
+ // event_wait_list are not valid events.
+ if event_wait_list.is_null() && num_events_in_wait_list > 0
+ || !event_wait_list.is_null() && num_events_in_wait_list == 0
+ {
+ return Err(CL_INVALID_EVENT_WAIT_LIST);
+ }
+
+ let res = Event::from_cl_arr(event_wait_list, num_events_in_wait_list)
+ .map_err(|_| CL_INVALID_EVENT_WAIT_LIST)?;
+
+ // CL_INVALID_CONTEXT if context associated with command_queue and events in event_list are not
+ // the same.
+ if res.iter().any(|e| e.context != q.context) {
+ return Err(CL_INVALID_CONTEXT);
+ }
+
+ Ok(res)
+}
+
+pub fn check_cb<T>(cb: &Option<T>, user_data: *mut c_void) -> CLResult<()> {
+ // CL_INVALID_VALUE if pfn_notify is NULL but user_data is not NULL.
+ if cb.is_none() && !user_data.is_null() {
+ return Err(CL_INVALID_VALUE);
+ }
+
+ Ok(())
+}
+
+pub fn checked_compare(a: usize, o: cmp::Ordering, b: u64) -> bool {
+ if usize::BITS > u64::BITS {
+ a.cmp(&(b as usize)) == o
+ } else {
+ (a as u64).cmp(&b) == o
+ }
+}
+
+pub fn is_alligned<T>(ptr: *const T, alignment: usize) -> bool {
+ ptr as usize & (alignment - 1) == 0
+}
+
+pub fn bit_check<A: BitAnd<Output = A> + PartialEq + Default, B: Into<A>>(a: A, b: B) -> bool {
+ a & b.into() != A::default()
+}
+
+// Taken from "Appendix D: Checking for Memory Copy Overlap"
+// src_offset and dst_offset are additions to support sub-buffers
+pub fn check_copy_overlap(
+ src_origin: &CLVec<usize>,
+ src_offset: usize,
+ dst_origin: &CLVec<usize>,
+ dst_offset: usize,
+ region: &CLVec<usize>,
+ row_pitch: usize,
+ slice_pitch: usize,
+) -> bool {
+ let slice_size = (region[1] - 1) * row_pitch + region[0];
+ let block_size = (region[2] - 1) * slice_pitch + slice_size;
+ let src_start =
+ src_origin[2] * slice_pitch + src_origin[1] * row_pitch + src_origin[0] + src_offset;
+ let src_end = src_start + block_size;
+ let dst_start =
+ dst_origin[2] * slice_pitch + dst_origin[1] * row_pitch + dst_origin[0] + dst_offset;
+ let dst_end = dst_start + block_size;
+
+ /* No overlap if dst ends before src starts or if src ends
+ * before dst starts.
+ */
+ if (dst_end <= src_start) || (src_end <= dst_start) {
+ return false;
+ }
+
+ /* No overlap if region[0] for dst or src fits in the gap
+ * between region[0] and row_pitch.
+ */
+ {
+ let src_dx = (src_origin[0] + src_offset) % row_pitch;
+ let dst_dx = (dst_origin[0] + dst_offset) % row_pitch;
+ if ((dst_dx >= src_dx + region[0]) && (dst_dx + region[0] <= src_dx + row_pitch))
+ || ((src_dx >= dst_dx + region[0]) && (src_dx + region[0] <= dst_dx + row_pitch))
+ {
+ return false;
+ }
+ }
+
+ /* No overlap if region[1] for dst or src fits in the gap
+ * between region[1] and slice_pitch.
+ */
+ {
+ let src_dy = (src_origin[1] * row_pitch + src_origin[0] + src_offset) % slice_pitch;
+ let dst_dy = (dst_origin[1] * row_pitch + dst_origin[0] + dst_offset) % slice_pitch;
+ if ((dst_dy >= src_dy + slice_size) && (dst_dy + slice_size <= src_dy + slice_pitch))
+ || ((src_dy >= dst_dy + slice_size) && (src_dy + slice_size <= dst_dy + slice_pitch))
+ {
+ return false;
+ }
+ }
+
+ /* Otherwise src and dst overlap. */
+ true
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::core::device::*;
+use crate::impl_cl_type_trait;
+
+use self::mesa_rust::pipe::resource::*;
+use self::rusticl_opencl_gen::*;
+
+use std::collections::HashMap;
+use std::convert::TryInto;
+use std::os::raw::c_void;
+use std::sync::Arc;
+
+pub struct Context {
+ pub base: CLObjectBase<CL_INVALID_CONTEXT>,
+ pub devs: Vec<Arc<Device>>,
+ pub properties: Vec<cl_context_properties>,
+}
+
+impl_cl_type_trait!(cl_context, Context, CL_INVALID_CONTEXT);
+
+impl Context {
+ pub fn new(devs: Vec<Arc<Device>>, properties: Vec<cl_context_properties>) -> Arc<Context> {
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ devs: devs,
+ properties: properties,
+ })
+ }
+
+ pub fn create_buffer(&self, size: usize) -> CLResult<HashMap<Arc<Device>, PipeResource>> {
+ let adj_size: u32 = size.try_into().map_err(|_| CL_OUT_OF_HOST_MEMORY)?;
+ let mut res = HashMap::new();
+ for dev in &self.devs {
+ let resource = dev
+ .screen()
+ .resource_create_buffer(adj_size)
+ .ok_or(CL_OUT_OF_RESOURCES);
+ res.insert(Arc::clone(dev), resource?);
+ }
+ Ok(res)
+ }
+
+ pub fn create_buffer_from_user(
+ &self,
+ size: usize,
+ user_ptr: *mut c_void,
+ ) -> CLResult<HashMap<Arc<Device>, PipeResource>> {
+ let adj_size: u32 = size.try_into().map_err(|_| CL_OUT_OF_HOST_MEMORY)?;
+ let mut res = HashMap::new();
+ for dev in &self.devs {
+ let resource = dev
+ .screen()
+ .resource_create_buffer_from_user(adj_size, user_ptr)
+ .ok_or(CL_OUT_OF_RESOURCES);
+ res.insert(Arc::clone(dev), resource?);
+ }
+ Ok(res)
+ }
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate mesa_rust_gen;
+extern crate mesa_rust_util;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::util::*;
+use crate::core::format::*;
+use crate::core::util::*;
+use crate::core::version::*;
+use crate::impl_cl_type_trait;
+
+use self::mesa_rust::pipe::context::*;
+use self::mesa_rust::pipe::device::load_screens;
+use self::mesa_rust::pipe::screen::*;
+use self::mesa_rust_gen::*;
+use self::rusticl_opencl_gen::*;
+
+use std::cmp::max;
+use std::cmp::min;
+use std::collections::HashMap;
+use std::convert::TryInto;
+use std::env;
+use std::sync::Arc;
+use std::sync::Mutex;
+use std::sync::MutexGuard;
+
+pub struct Device {
+ pub base: CLObjectBase<CL_INVALID_DEVICE>,
+ screen: Arc<PipeScreen>,
+ pub cl_version: CLVersion,
+ pub clc_version: CLVersion,
+ pub clc_versions: Vec<cl_name_version>,
+ pub custom: bool,
+ pub embedded: bool,
+ pub extension_string: String,
+ pub extensions: Vec<cl_name_version>,
+ pub formats: HashMap<cl_image_format, HashMap<cl_mem_object_type, cl_mem_flags>>,
+ helper_ctx: Mutex<Arc<PipeContext>>,
+}
+
+impl_cl_type_trait!(cl_device_id, Device, CL_INVALID_DEVICE);
+
+impl Device {
+ fn new(screen: Arc<PipeScreen>) -> Option<Arc<Device>> {
+ let mut d = Self {
+ base: CLObjectBase::new(),
+ helper_ctx: Mutex::new(screen.create_context().unwrap()),
+ screen: screen,
+ cl_version: CLVersion::Cl3_0,
+ clc_version: CLVersion::Cl3_0,
+ clc_versions: Vec::new(),
+ custom: false,
+ embedded: false,
+ extension_string: String::from(""),
+ extensions: Vec::new(),
+ formats: HashMap::new(),
+ };
+
+ if !d.check_valid() {
+ return None;
+ }
+
+ d.fill_format_tables();
+
+ // check if we are embedded or full profile first
+ d.embedded = d.check_embedded_profile();
+
+ // check if we have to report it as a custom device
+ d.custom = d.check_custom();
+
+ // query supported extensions
+ d.fill_extensions();
+
+ // now figure out what version we are
+ d.check_version();
+
+ Some(Arc::new(d))
+ }
+
+ fn fill_format_tables(&mut self) {
+ for f in FORMATS {
+ let mut fs = HashMap::new();
+ for t in CL_IMAGE_TYPES {
+ let mut flags: cl_uint = 0;
+ if self.screen.is_format_supported(
+ f.pipe,
+ cl_mem_type_to_texture_target(t),
+ PIPE_BIND_SAMPLER_VIEW,
+ ) {
+ flags |= CL_MEM_READ_ONLY;
+ }
+ if self.screen.is_format_supported(
+ f.pipe,
+ cl_mem_type_to_texture_target(t),
+ PIPE_BIND_SHADER_IMAGE,
+ ) {
+ flags |= CL_MEM_WRITE_ONLY;
+ // TODO: enable once we support it
+ // flags |= CL_MEM_KERNEL_READ_AND_WRITE;
+ }
+ if self.screen.is_format_supported(
+ f.pipe,
+ cl_mem_type_to_texture_target(t),
+ PIPE_BIND_SAMPLER_VIEW | PIPE_BIND_SHADER_IMAGE,
+ ) {
+ flags |= CL_MEM_READ_WRITE;
+ }
+ fs.insert(t, flags as cl_mem_flags);
+ }
+ self.formats.insert(f.cl_image_format, fs);
+ }
+ }
+
+ fn check_valid(&self) -> bool {
+ if self.screen.param(pipe_cap::PIPE_CAP_COMPUTE) == 0 ||
+ // even though we use PIPE_SHADER_IR_NIR, PIPE_SHADER_IR_NIR_SERIALIZED marks CL support by the driver
+ self.shader_param(pipe_shader_cap::PIPE_SHADER_CAP_SUPPORTED_IRS) & (1 << (pipe_shader_ir::PIPE_SHADER_IR_NIR_SERIALIZED as i32)) == 0
+ {
+ return false;
+ }
+
+ // CL_DEVICE_MAX_PARAMETER_SIZE
+ // For this minimum value, only a maximum of 128 arguments can be passed to a kernel
+ if self.param_max_size() < 128 {
+ return false;
+ }
+ true
+ }
+
+ fn check_custom(&self) -> bool {
+ // Max size of memory object allocation in bytes. The minimum value is
+ // max(min(1024 × 1024 × 1024, 1/4th of CL_DEVICE_GLOBAL_MEM_SIZE), 32 × 1024 × 1024)
+ // for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ let mut limit = min(1024 * 1024 * 1024, self.global_mem_size());
+ limit = max(limit, 32 * 1024 * 1024);
+ if self.max_mem_alloc() < limit {
+ return true;
+ }
+
+ // CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
+ // The minimum value is 3 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.max_grid_dimensions() < 3 {
+ return true;
+ }
+
+ if self.embedded {
+ // CL_DEVICE_MAX_PARAMETER_SIZE
+ // The minimum value is 256 bytes for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.param_max_size() < 256 {
+ return true;
+ }
+
+ // CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
+ // The minimum value is 1 KB for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.const_max_size() < 1024 {
+ return true;
+ }
+
+ // TODO
+ // CL_DEVICE_MAX_CONSTANT_ARGS
+ // The minimum value is 4 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+
+ // CL_DEVICE_LOCAL_MEM_SIZE
+ // The minimum value is 1 KB for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.local_mem_size() < 1024 {
+ return true;
+ }
+ } else {
+ // CL 1.0 spec:
+ // CL_DEVICE_MAX_PARAMETER_SIZE
+ // The minimum value is 256 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.param_max_size() < 256 {
+ return true;
+ }
+
+ // CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
+ // The minimum value is 64 KB for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.const_max_size() < 64 * 1024 {
+ return true;
+ }
+
+ // TODO
+ // CL_DEVICE_MAX_CONSTANT_ARGS
+ // The minimum value is 8 for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+
+ // CL 1.0 spec:
+ // CL_DEVICE_LOCAL_MEM_SIZE
+ // The minimum value is 16 KB for devices that are not of type CL_DEVICE_TYPE_CUSTOM.
+ if self.local_mem_size() < 16 * 1024 {
+ return true;
+ }
+ }
+
+ false
+ }
+
+ fn check_embedded_profile(&self) -> bool {
+ if self.image_supported() {
+ // The minimum value is 16 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ if self.max_samplers() < 16 ||
+ // The minimum value is 128 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_read_count() < 128 ||
+ // The minimum value is 64 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_write_count() < 64 ||
+ // The minimum value is 16384 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_2d_size() < 16384 ||
+ // The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_array_size() < 2048 ||
+ // The minimum value is 65536 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_buffer_size() < 65536
+ {
+ return true;
+ }
+
+ // TODO check req formats
+ }
+ false
+ }
+
+ fn parse_env_version() -> Option<CLVersion> {
+ let val = env::var("RUSTICL_CL_VERSION").ok()?;
+ let (major, minor) = val.split_once('.')?;
+ let major = major.parse().ok()?;
+ let minor = minor.parse().ok()?;
+ mk_cl_version(major, minor, 0).try_into().ok()
+ }
+
+ // TODO add CLC checks
+ fn check_version(&mut self) {
+ let exts: Vec<&str> = self.extension_string.split(' ').collect();
+ let mut res = CLVersion::Cl3_0;
+
+ if self.embedded {
+ if self.image_supported() {
+ let supports_array_writes = !FORMATS
+ .iter()
+ .filter(|f| f.req_for_embeded_read_or_write)
+ .map(|f| self.formats.get(&f.cl_image_format).unwrap())
+ .map(|f| f.get(&CL_MEM_OBJECT_IMAGE2D_ARRAY).unwrap())
+ .any(|f| *f & cl_mem_flags::from(CL_MEM_WRITE_ONLY) == 0);
+ if self.image_3d_size() < 2048 || !supports_array_writes {
+ res = CLVersion::Cl1_2;
+ }
+ }
+ }
+
+ // TODO: check image 1D, 1Dbuffer, 1Darray and 2Darray support explicitly
+ if self.image_supported() {
+ // The minimum value is 256 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ if self.image_array_size() < 256 ||
+ // The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_buffer_size() < 2048
+ {
+ res = CLVersion::Cl1_1;
+ }
+ }
+
+ if !exts.contains(&"cl_khr_byte_addressable_store")
+ || !exts.contains(&"cl_khr_global_int32_base_atomics")
+ || !exts.contains(&"cl_khr_global_int32_extended_atomics")
+ || !exts.contains(&"cl_khr_local_int32_base_atomics")
+ || !exts.contains(&"cl_khr_local_int32_extended_atomics")
+ // The following modifications are made to the OpenCL 1.1 platform layer and runtime (sections 4 and 5):
+ // The minimum FULL_PROFILE value for CL_DEVICE_MAX_PARAMETER_SIZE increased from 256 to 1024 bytes
+ || self.param_max_size() < 1024
+ // The minimum FULL_PROFILE value for CL_DEVICE_LOCAL_MEM_SIZE increased from 16 KB to 32 KB.
+ || self.local_mem_size() < 32 * 1024
+ {
+ res = CLVersion::Cl1_0;
+ }
+
+ if let Some(val) = Self::parse_env_version() {
+ res = val;
+ }
+
+ if res >= CLVersion::Cl3_0 {
+ self.clc_versions
+ .push(mk_cl_version_ext(3, 0, 0, "OpenCL C"));
+ }
+
+ if res >= CLVersion::Cl1_2 {
+ self.clc_versions
+ .push(mk_cl_version_ext(1, 2, 0, "OpenCL C"));
+ }
+
+ if res >= CLVersion::Cl1_1 {
+ self.clc_versions
+ .push(mk_cl_version_ext(1, 1, 0, "OpenCL C"));
+ }
+
+ if res >= CLVersion::Cl1_0 {
+ self.clc_versions
+ .push(mk_cl_version_ext(1, 0, 0, "OpenCL C"));
+ }
+
+ self.cl_version = res;
+ self.clc_version = min(CLVersion::Cl1_2, res);
+ }
+
+ fn fill_extensions(&mut self) {
+ let mut exts_str: Vec<String> = Vec::new();
+ let mut exts = Vec::new();
+ let mut add_ext = |major, minor, patch, ext| {
+ exts.push(mk_cl_version_ext(major, minor, patch, ext));
+ exts_str.push(ext.to_owned());
+ };
+
+ add_ext(1, 0, 0, "cl_khr_byte_addressable_store");
+
+ self.extensions = exts;
+ self.extension_string = exts_str.join(" ");
+ }
+
+ fn shader_param(&self, cap: pipe_shader_cap) -> i32 {
+ self.screen
+ .shader_param(pipe_shader_type::PIPE_SHADER_COMPUTE, cap)
+ }
+
+ pub fn all() -> Vec<Arc<Device>> {
+ load_screens().into_iter().filter_map(Device::new).collect()
+ }
+
+ pub fn address_bits(&self) -> cl_uint {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_ADDRESS_BITS)
+ }
+
+ pub fn const_max_size(&self) -> cl_ulong {
+ self.screen
+ .param(pipe_cap::PIPE_CAP_MAX_SHADER_BUFFER_SIZE_UINT) as u64
+ }
+
+ pub fn device_type(&self) -> cl_device_type {
+ if self.custom {
+ return CL_DEVICE_TYPE_CUSTOM as cl_device_type;
+ }
+ (match self.screen.device_type() {
+ pipe_loader_device_type::PIPE_LOADER_DEVICE_SOFTWARE => CL_DEVICE_TYPE_CPU,
+ pipe_loader_device_type::PIPE_LOADER_DEVICE_PCI => {
+ CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_DEFAULT
+ }
+ pipe_loader_device_type::PIPE_LOADER_DEVICE_PLATFORM => {
+ CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_DEFAULT
+ }
+ pipe_loader_device_type::NUM_PIPE_LOADER_DEVICE_TYPES => CL_DEVICE_TYPE_CUSTOM,
+ }) as cl_device_type
+ }
+
+ pub fn global_mem_size(&self) -> cl_ulong {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE)
+ }
+
+ pub fn image_2d_size(&self) -> usize {
+ self.screen.param(pipe_cap::PIPE_CAP_MAX_TEXTURE_2D_SIZE) as usize
+ }
+
+ pub fn image_3d_size(&self) -> usize {
+ 1 << (self.screen.param(pipe_cap::PIPE_CAP_MAX_TEXTURE_3D_LEVELS) - 1)
+ }
+
+ pub fn image_3d_supported(&self) -> bool {
+ self.screen.param(pipe_cap::PIPE_CAP_MAX_TEXTURE_3D_LEVELS) != 0
+ }
+
+ pub fn image_array_size(&self) -> usize {
+ self.screen
+ .param(pipe_cap::PIPE_CAP_MAX_TEXTURE_ARRAY_LAYERS) as usize
+ }
+
+ pub fn image_base_address_alignment(&self) -> cl_uint {
+ 0
+ }
+
+ pub fn image_buffer_size(&self) -> usize {
+ self.screen
+ .param(pipe_cap::PIPE_CAP_MAX_TEXEL_BUFFER_ELEMENTS_UINT) as usize
+ }
+
+ pub fn image_read_count(&self) -> cl_uint {
+ self.shader_param(pipe_shader_cap::PIPE_SHADER_CAP_MAX_SAMPLER_VIEWS) as cl_uint
+ }
+
+ pub fn image_supported(&self) -> bool {
+ // TODO check CL_DEVICE_IMAGE_SUPPORT reqs
+ self.shader_param(pipe_shader_cap::PIPE_SHADER_CAP_MAX_SHADER_IMAGES) != 0 &&
+ // The minimum value is 8 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_read_count() >= 8 &&
+ // The minimum value is 8 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_write_count() >= 8 &&
+ // The minimum value is 2048 if CL_DEVICE_IMAGE_SUPPORT is CL_TRUE
+ self.image_2d_size() >= 2048
+ }
+
+ pub fn image_write_count(&self) -> cl_uint {
+ self.shader_param(pipe_shader_cap::PIPE_SHADER_CAP_MAX_SHADER_IMAGES) as cl_uint
+ }
+
+ pub fn little_endian(&self) -> bool {
+ let endianness = self.screen.param(pipe_cap::PIPE_CAP_ENDIANNESS);
+ endianness == (pipe_endian::PIPE_ENDIAN_LITTLE as i32)
+ }
+
+ pub fn local_mem_size(&self) -> cl_ulong {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE)
+ }
+
+ pub fn max_block_sizes(&self) -> Vec<usize> {
+ let v: Vec<u64> = self
+ .screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE);
+ v.into_iter().map(|v| v as usize).collect()
+ }
+
+ pub fn max_clock_freq(&self) -> cl_uint {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_CLOCK_FREQUENCY)
+ }
+
+ pub fn max_compute_units(&self) -> cl_uint {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS)
+ }
+
+ pub fn max_grid_dimensions(&self) -> cl_uint {
+ ComputeParam::<u64>::compute_param(
+ self.screen.as_ref(),
+ pipe_compute_cap::PIPE_COMPUTE_CAP_GRID_DIMENSION,
+ ) as cl_uint
+ }
+
+ pub fn max_mem_alloc(&self) -> cl_ulong {
+ self.screen
+ .compute_param(pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE)
+ }
+
+ pub fn max_samplers(&self) -> cl_uint {
+ self.shader_param(pipe_shader_cap::PIPE_SHADER_CAP_MAX_TEXTURE_SAMPLERS) as cl_uint
+ }
+
+ pub fn max_threads_per_block(&self) -> usize {
+ ComputeParam::<u64>::compute_param(
+ self.screen.as_ref(),
+ pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK,
+ ) as usize
+ }
+
+ pub fn param_max_size(&self) -> usize {
+ ComputeParam::<u64>::compute_param(
+ self.screen.as_ref(),
+ pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_INPUT_SIZE,
+ ) as usize
+ }
+
+ pub fn screen(&self) -> &Arc<PipeScreen> {
+ &self.screen
+ }
+
+ pub fn unified_memory(&self) -> bool {
+ self.screen.param(pipe_cap::PIPE_CAP_UMA) == 1
+ }
+
+ pub fn vendor_id(&self) -> cl_uint {
+ let id = self.screen.param(pipe_cap::PIPE_CAP_VENDOR_ID);
+ if id == -1 {
+ return 0;
+ }
+ id as u32
+ }
+
+ pub fn helper_ctx(&self) -> MutexGuard<Arc<PipeContext>> {
+ self.helper_ctx.lock().unwrap()
+ }
+
+ pub fn cl_features(&self) -> clc_optional_features {
+ clc_optional_features {
+ fp16: false,
+ fp64: false,
+ int64: false,
+ images: self.image_supported(),
+ images_read_write: false,
+ images_write_3d: false,
+ intel_subgroups: false,
+ subgroups: false,
+ }
+ }
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::core::context::*;
+use crate::core::queue::*;
+use crate::impl_cl_type_trait;
+
+use self::mesa_rust::pipe::context::*;
+use self::rusticl_opencl_gen::*;
+
+use std::slice;
+use std::sync::atomic::AtomicI32;
+use std::sync::atomic::Ordering;
+use std::sync::Arc;
+
+pub type EventSig = Box<dyn Fn(&Arc<Queue>, &Arc<PipeContext>) -> CLResult<()>>;
+
+#[repr(C)]
+pub struct Event {
+ pub base: CLObjectBase<CL_INVALID_EVENT>,
+ pub context: Arc<Context>,
+ pub queue: Option<Arc<Queue>>,
+ pub cmd_type: cl_command_type,
+ pub deps: Vec<Arc<Event>>,
+ // use AtomicI32 instead of cl_int so we can change it without a &mut reference
+ status: AtomicI32,
+ work: Option<EventSig>,
+}
+
+impl_cl_type_trait!(cl_event, Event, CL_INVALID_EVENT);
+
+// TODO shouldn't be needed, but... uff C pointers are annoying
+unsafe impl Send for Event {}
+unsafe impl Sync for Event {}
+
+impl Event {
+ pub fn new(
+ queue: &Arc<Queue>,
+ cmd_type: cl_command_type,
+ deps: Vec<Arc<Event>>,
+ work: EventSig,
+ ) -> Arc<Event> {
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: queue.context.clone(),
+ queue: Some(queue.clone()),
+ cmd_type: cmd_type,
+ deps: deps,
+ status: AtomicI32::new(CL_QUEUED as cl_int),
+ work: Some(work),
+ })
+ }
+
+ pub fn new_user(context: Arc<Context>) -> Arc<Event> {
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ queue: None,
+ cmd_type: CL_COMMAND_USER,
+ deps: Vec::new(),
+ status: AtomicI32::new(CL_SUBMITTED as cl_int),
+ work: None,
+ })
+ }
+
+ pub fn from_cl_arr(events: *const cl_event, num_events: u32) -> CLResult<Vec<Arc<Event>>> {
+ let s = unsafe { slice::from_raw_parts(events, num_events as usize) };
+ s.iter().map(|e| e.get_arc()).collect()
+ }
+
+ pub fn is_error(&self) -> bool {
+ self.status.load(Ordering::Relaxed) < 0
+ }
+
+ pub fn status(&self) -> cl_int {
+ self.status.load(Ordering::Relaxed)
+ }
+
+ // We always assume that work here simply submits stuff to the hardware even if it's just doing
+ // sw emulation or nothing at all.
+ // If anything requets waiting, we will update the status through fencing later.
+ pub fn call(&self, ctx: &Arc<PipeContext>) -> cl_int {
+ let status = self.status();
+ if status == CL_QUEUED as cl_int {
+ let new = self.work.as_ref().map_or(
+ // if there is no work
+ CL_SUBMITTED as cl_int,
+ |w| {
+ w(self.queue.as_ref().unwrap(), ctx).err().map_or(
+ // if there is an error, negate it
+ CL_SUBMITTED as cl_int,
+ |e| e,
+ )
+ },
+ );
+ self.status.store(new, Ordering::Relaxed);
+ new
+ } else {
+ status
+ }
+ }
+}
+
+// TODO worker thread per device
+// Condvar to wait on new events to work on
+// notify condvar when flushing queue events to worker
+// attach fence to flushed events on context->flush
+// store "newest" event for in-order queues per queue
+// reordering/graph building done in worker
--- /dev/null
+extern crate mesa_rust_gen;
+extern crate rusticl_opencl_gen;
+
+use crate::api::util::*;
+
+use self::mesa_rust_gen::pipe_format;
+use self::rusticl_opencl_gen::*;
+
+pub struct RusticlImageFormat {
+ pub cl_image_format: cl_image_format,
+ pub req_for_full_read_or_write: bool,
+ pub req_for_embeded_read_or_write: bool,
+ pub req_for_full_read_and_write: bool,
+ pub pipe: pipe_format,
+}
+
+pub const fn rusticl_image_format(
+ cl_image_format: cl_image_format,
+ req_for_full_read_or_write: bool,
+ req_for_embeded_read_or_write: bool,
+ req_for_full_read_and_write: bool,
+ pipe: pipe_format,
+) -> RusticlImageFormat {
+ RusticlImageFormat {
+ cl_image_format: cl_image_format,
+ req_for_full_read_or_write: req_for_full_read_or_write,
+ req_for_embeded_read_or_write: req_for_embeded_read_or_write,
+ req_for_full_read_and_write: req_for_full_read_and_write,
+ pipe: pipe,
+ }
+}
+
+pub const FORMATS: &[RusticlImageFormat] = &[
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_HALF_FLOAT),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R16_FLOAT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_FLOAT),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R32_FLOAT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_SIGNED_INT8),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R8_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_SIGNED_INT16),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R16_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_SIGNED_INT32),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R32_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_UNORM_INT8),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R8_UNORM,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_UNORM_INT16),
+ false,
+ false,
+ false,
+ pipe_format::PIPE_FORMAT_R16_UNORM,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_UNSIGNED_INT8),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R8_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_UNSIGNED_INT16),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R16_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_R, CL_UNSIGNED_INT32),
+ false,
+ false,
+ true,
+ pipe_format::PIPE_FORMAT_R32_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_HALF_FLOAT),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R16G16B16A16_FLOAT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_FLOAT),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R32G32B32A32_FLOAT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_SIGNED_INT8),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R8G8B8A8_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_SIGNED_INT16),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R16G16B16A16_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_SIGNED_INT32),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R32G32B32A32_SINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_UNORM_INT8),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R8G8B8A8_UNORM,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_UNORM_INT16),
+ true,
+ true,
+ false,
+ pipe_format::PIPE_FORMAT_R16G16B16A16_UNORM,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_UNSIGNED_INT8),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R8G8B8A8_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_UNSIGNED_INT16),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R16G16B16A16_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_RGBA, CL_UNSIGNED_INT32),
+ true,
+ true,
+ true,
+ pipe_format::PIPE_FORMAT_R32G32B32A32_UINT,
+ ),
+ rusticl_image_format(
+ cl_image_format(CL_BGRA, CL_UNORM_INT8),
+ true,
+ false,
+ false,
+ pipe_format::PIPE_FORMAT_B8G8R8A8_UNORM,
+ ),
+];
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::impl_cl_type_trait;
+
+use self::rusticl_opencl_gen::*;
+
+use std::sync::Arc;
+
+#[repr(C)]
+pub struct Kernel {
+ pub base: CLObjectBase<CL_INVALID_KERNEL>,
+}
+
+impl_cl_type_trait!(cl_kernel, Kernel, CL_INVALID_KERNEL);
+
+impl Kernel {
+ pub fn new() -> Arc<Kernel> {
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ })
+ }
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::api::types::*;
+use crate::api::util::*;
+use crate::core::context::*;
+use crate::core::device::*;
+use crate::core::queue::*;
+use crate::impl_cl_type_trait;
+
+use self::mesa_rust::pipe::context::*;
+use self::mesa_rust::pipe::resource::*;
+use self::mesa_rust::pipe::transfer::*;
+use self::rusticl_opencl_gen::*;
+
+use std::collections::HashMap;
+use std::convert::TryInto;
+use std::os::raw::c_void;
+use std::ptr;
+use std::sync::Arc;
+use std::sync::Mutex;
+
+#[repr(C)]
+pub struct Mem {
+ pub base: CLObjectBase<CL_INVALID_MEM_OBJECT>,
+ pub context: Arc<Context>,
+ pub parent: Option<Arc<Mem>>,
+ pub mem_type: cl_mem_object_type,
+ pub flags: cl_mem_flags,
+ pub size: usize,
+ pub offset: usize,
+ pub host_ptr: *mut c_void,
+ pub image_format: cl_image_format,
+ pub image_desc: cl_image_desc,
+ pub image_elem_size: u8,
+ pub cbs: Mutex<Vec<Box<dyn Fn(cl_mem)>>>,
+ res: Option<HashMap<Arc<Device>, PipeResource>>,
+ maps: Mutex<HashMap<*mut c_void, (u32, PipeTransfer)>>,
+}
+
+impl_cl_type_trait!(cl_mem, Mem, CL_INVALID_MEM_OBJECT);
+
+fn sw_copy(
+ src: *const c_void,
+ dst: *mut c_void,
+ region: &CLVec<usize>,
+ src_origin: &CLVec<usize>,
+ src_row_pitch: usize,
+ src_slice_pitch: usize,
+ dst_origin: &CLVec<usize>,
+ dst_row_pitch: usize,
+ dst_slice_pitch: usize,
+) {
+ for z in 0..region[2] {
+ for y in 0..region[1] {
+ unsafe {
+ ptr::copy_nonoverlapping(
+ src.add((*src_origin + [0, y, z]) * [1, src_row_pitch, src_slice_pitch]),
+ dst.add((*dst_origin + [0, y, z]) * [1, dst_row_pitch, dst_slice_pitch]),
+ region[0],
+ )
+ };
+ }
+ }
+}
+
+impl Mem {
+ pub fn new_buffer(
+ context: Arc<Context>,
+ flags: cl_mem_flags,
+ size: usize,
+ host_ptr: *mut c_void,
+ ) -> CLResult<Arc<Mem>> {
+ if bit_check(flags, CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR) {
+ println!("host ptr semantics not implemented!");
+ }
+
+ let buffer = if bit_check(flags, CL_MEM_USE_HOST_PTR) {
+ context.create_buffer_from_user(size, host_ptr)
+ } else {
+ context.create_buffer(size)
+ }?;
+
+ let host_ptr = if bit_check(flags, CL_MEM_USE_HOST_PTR) {
+ host_ptr
+ } else {
+ ptr::null_mut()
+ };
+
+ Ok(Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ parent: None,
+ mem_type: CL_MEM_OBJECT_BUFFER,
+ flags: flags,
+ size: size,
+ offset: 0,
+ host_ptr: host_ptr,
+ image_format: cl_image_format::default(),
+ image_desc: cl_image_desc::default(),
+ image_elem_size: 0,
+ cbs: Mutex::new(Vec::new()),
+ res: Some(buffer),
+ maps: Mutex::new(HashMap::new()),
+ }))
+ }
+
+ pub fn new_sub_buffer(
+ parent: Arc<Mem>,
+ flags: cl_mem_flags,
+ offset: usize,
+ size: usize,
+ ) -> Arc<Mem> {
+ let host_ptr = if parent.host_ptr.is_null() {
+ ptr::null_mut()
+ } else {
+ unsafe { parent.host_ptr.add(offset) }
+ };
+
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: parent.context.clone(),
+ parent: Some(parent),
+ mem_type: CL_MEM_OBJECT_BUFFER,
+ flags: flags,
+ size: size,
+ offset: offset,
+ host_ptr: host_ptr,
+ image_format: cl_image_format::default(),
+ image_desc: cl_image_desc::default(),
+ image_elem_size: 0,
+ cbs: Mutex::new(Vec::new()),
+ res: None,
+ maps: Mutex::new(HashMap::new()),
+ })
+ }
+
+ pub fn new_image(
+ context: Arc<Context>,
+ mem_type: cl_mem_object_type,
+ flags: cl_mem_flags,
+ image_format: &cl_image_format,
+ image_desc: cl_image_desc,
+ image_elem_size: u8,
+ host_ptr: *mut c_void,
+ ) -> Arc<Mem> {
+ if bit_check(
+ flags,
+ CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR,
+ ) {
+ println!("host ptr semantics not implemented!");
+ }
+
+ let host_ptr = if bit_check(flags, CL_MEM_USE_HOST_PTR) {
+ host_ptr
+ } else {
+ ptr::null_mut()
+ };
+
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ parent: None,
+ mem_type: mem_type,
+ flags: flags,
+ size: 0,
+ offset: 0,
+ host_ptr: host_ptr,
+ image_format: *image_format,
+ image_desc: image_desc,
+ image_elem_size: image_elem_size,
+ cbs: Mutex::new(Vec::new()),
+ res: None,
+ maps: Mutex::new(HashMap::new()),
+ })
+ }
+
+ pub fn is_buffer(&self) -> bool {
+ self.mem_type == CL_MEM_OBJECT_BUFFER
+ }
+
+ pub fn has_same_parent(&self, other: &Self) -> bool {
+ let a = self.parent.as_ref().map_or(self, |p| p);
+ let b = other.parent.as_ref().map_or(other, |p| p);
+ ptr::eq(a, b)
+ }
+
+ fn get_res(&self) -> &HashMap<Arc<Device>, PipeResource> {
+ self.parent
+ .as_ref()
+ .map_or(self, |p| p.as_ref())
+ .res
+ .as_ref()
+ .unwrap()
+ }
+
+ pub fn write_from_user(
+ &self,
+ q: &Arc<Queue>,
+ ctx: &PipeContext,
+ offset: usize,
+ ptr: *const c_void,
+ size: usize,
+ ) -> CLResult<()> {
+ // TODO support sub buffers
+ let r = self.get_res().get(&q.device).unwrap();
+ ctx.buffer_subdata(
+ r,
+ offset.try_into().map_err(|_| CL_OUT_OF_HOST_MEMORY)?,
+ ptr,
+ size.try_into().map_err(|_| CL_OUT_OF_HOST_MEMORY)?,
+ );
+ Ok(())
+ }
+
+ pub fn write_from_user_rect(
+ &self,
+ src: *const c_void,
+ q: &Arc<Queue>,
+ ctx: &Arc<PipeContext>,
+ region: &CLVec<usize>,
+ src_origin: &CLVec<usize>,
+ src_row_pitch: usize,
+ src_slice_pitch: usize,
+ dst_origin: &CLVec<usize>,
+ dst_row_pitch: usize,
+ dst_slice_pitch: usize,
+ ) -> CLResult<()> {
+ let r = self.res.as_ref().unwrap().get(&q.device).unwrap();
+ let tx = ctx.buffer_map(r, 0, self.size.try_into().unwrap(), true);
+
+ sw_copy(
+ src,
+ tx.ptr(),
+ region,
+ src_origin,
+ src_row_pitch,
+ src_slice_pitch,
+ dst_origin,
+ dst_row_pitch,
+ dst_slice_pitch,
+ );
+
+ drop(tx);
+ Ok(())
+ }
+
+ pub fn read_to_user_rect(
+ &self,
+ dst: *mut c_void,
+ q: &Arc<Queue>,
+ ctx: &Arc<PipeContext>,
+ region: &CLVec<usize>,
+ src_origin: &CLVec<usize>,
+ src_row_pitch: usize,
+ src_slice_pitch: usize,
+ dst_origin: &CLVec<usize>,
+ dst_row_pitch: usize,
+ dst_slice_pitch: usize,
+ ) -> CLResult<()> {
+ let r = self.res.as_ref().unwrap().get(&q.device).unwrap();
+ let tx = ctx.buffer_map(r, 0, self.size.try_into().unwrap(), true);
+
+ sw_copy(
+ tx.ptr(),
+ dst,
+ region,
+ src_origin,
+ src_row_pitch,
+ src_slice_pitch,
+ dst_origin,
+ dst_row_pitch,
+ dst_slice_pitch,
+ );
+
+ drop(tx);
+ Ok(())
+ }
+
+ pub fn copy_to(
+ &self,
+ dst: &Self,
+ q: &Arc<Queue>,
+ ctx: &Arc<PipeContext>,
+ region: &CLVec<usize>,
+ src_origin: &CLVec<usize>,
+ src_row_pitch: usize,
+ src_slice_pitch: usize,
+ dst_origin: &CLVec<usize>,
+ dst_row_pitch: usize,
+ dst_slice_pitch: usize,
+ ) -> CLResult<()> {
+ let res_src = self.res.as_ref().unwrap().get(&q.device).unwrap();
+ let res_dst = dst.res.as_ref().unwrap().get(&q.device).unwrap();
+
+ let tx_src = ctx.buffer_map(res_src, 0, self.size.try_into().unwrap(), true);
+ let tx_dst = ctx.buffer_map(res_dst, 0, dst.size.try_into().unwrap(), true);
+
+ // TODO check to use hw accelerated paths (e.g. resource_copy_region or blits)
+ sw_copy(
+ tx_src.ptr(),
+ tx_dst.ptr(),
+ region,
+ src_origin,
+ src_row_pitch,
+ src_slice_pitch,
+ dst_origin,
+ dst_row_pitch,
+ dst_slice_pitch,
+ );
+
+ drop(tx_src);
+ drop(tx_dst);
+
+ Ok(())
+ }
+
+ // TODO use PIPE_MAP_UNSYNCHRONIZED for non blocking
+ pub fn map(&self, q: &Arc<Queue>, offset: usize, size: usize, block: bool) -> *mut c_void {
+ let res = self.res.as_ref().unwrap().get(&q.device).unwrap();
+ let tx = q.device.helper_ctx().buffer_map(
+ res,
+ offset.try_into().unwrap(),
+ size.try_into().unwrap(),
+ block,
+ );
+ let ptr = tx.ptr();
+ let mut lock = self.maps.lock().unwrap();
+ let e = lock.get_mut(&ptr);
+
+ // if we already have a mapping, reuse that and increase the refcount
+ if let Some(e) = e {
+ e.0 += 1;
+ } else {
+ lock.insert(tx.ptr(), (1, tx));
+ }
+
+ ptr
+ }
+
+ pub fn is_mapped_ptr(&self, ptr: *mut c_void) -> bool {
+ self.maps.lock().unwrap().contains_key(&ptr)
+ }
+
+ pub fn unmap(&self, q: &Arc<Queue>, ptr: *mut c_void) {
+ let mut lock = self.maps.lock().unwrap();
+ let e = lock.get_mut(&ptr).unwrap();
+
+ e.0 -= 1;
+ if e.0 == 0 {
+ lock.remove(&ptr)
+ .unwrap()
+ .1
+ .with_ctx(&q.device.helper_ctx());
+ }
+ }
+}
+
+impl Drop for Mem {
+ fn drop(&mut self) {
+ let cl = cl_mem::from_ptr(self);
+ self.cbs
+ .get_mut()
+ .unwrap()
+ .iter()
+ .rev()
+ .for_each(|cb| cb(cl));
+ }
+}
+
+#[repr(C)]
+pub struct Sampler {
+ pub base: CLObjectBase<CL_INVALID_SAMPLER>,
+ pub context: Arc<Context>,
+ pub normalized_coords: bool,
+ pub addressing_mode: cl_addressing_mode,
+ pub filter_mode: cl_filter_mode,
+}
+
+impl_cl_type_trait!(cl_sampler, Sampler, CL_INVALID_SAMPLER);
+
+impl Sampler {
+ pub fn new(
+ context: Arc<Context>,
+ normalized_coords: bool,
+ addressing_mode: cl_addressing_mode,
+ filter_mode: cl_filter_mode,
+ ) -> Arc<Sampler> {
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ normalized_coords: normalized_coords,
+ addressing_mode: addressing_mode,
+ filter_mode: filter_mode,
+ })
+ }
+}
--- /dev/null
+pub mod context;
+pub mod device;
+pub mod event;
+pub mod format;
+pub mod kernel;
+pub mod memory;
+pub mod program;
+pub mod queue;
+pub mod util;
+pub mod version;
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::core::context::*;
+use crate::core::device::*;
+use crate::impl_cl_type_trait;
+
+use self::mesa_rust::compiler::clc::*;
+use self::rusticl_opencl_gen::*;
+
+use std::collections::HashMap;
+use std::collections::HashSet;
+use std::ffi::CString;
+use std::sync::Arc;
+use std::sync::Mutex;
+use std::sync::MutexGuard;
+
+#[repr(C)]
+pub struct Program {
+ pub base: CLObjectBase<CL_INVALID_PROGRAM>,
+ pub context: Arc<Context>,
+ pub devs: Vec<Arc<Device>>,
+ pub src: CString,
+ build: Mutex<ProgramBuild>,
+}
+
+impl_cl_type_trait!(cl_program, Program, CL_INVALID_PROGRAM);
+
+struct ProgramBuild {
+ builds: HashMap<Arc<Device>, ProgramDevBuild>,
+ kernels: Vec<String>,
+}
+
+struct ProgramDevBuild {
+ spirv: Option<spirv::SPIRVBin>,
+ status: cl_build_status,
+ options: String,
+ log: String,
+}
+
+fn prepare_options(options: &str) -> Vec<CString> {
+ options
+ .split_whitespace()
+ .map(|a| match a {
+ "-cl-denorms-are-zero" => "-fdenormal-fp-math=positive-zero",
+ _ => a,
+ })
+ .map(CString::new)
+ .map(Result::unwrap)
+ .collect()
+}
+
+impl Program {
+ pub fn new(context: &Arc<Context>, devs: &[Arc<Device>], src: CString) -> Arc<Program> {
+ let builds = devs
+ .iter()
+ .map(|d| {
+ (
+ d.clone(),
+ ProgramDevBuild {
+ spirv: None,
+ status: CL_BUILD_NONE,
+ log: String::from(""),
+ options: String::from(""),
+ },
+ )
+ })
+ .collect();
+
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context.clone(),
+ devs: devs.to_vec(),
+ src: src,
+ build: Mutex::new(ProgramBuild {
+ builds: builds,
+ kernels: Vec::new(),
+ }),
+ })
+ }
+
+ fn build_info(&self) -> MutexGuard<ProgramBuild> {
+ self.build.lock().unwrap()
+ }
+
+ fn dev_build_info<'a>(
+ l: &'a mut MutexGuard<ProgramBuild>,
+ dev: &Arc<Device>,
+ ) -> &'a mut ProgramDevBuild {
+ l.builds.get_mut(dev).unwrap()
+ }
+
+ pub fn status(&self, dev: &Arc<Device>) -> cl_build_status {
+ Self::dev_build_info(&mut self.build_info(), dev).status
+ }
+
+ pub fn log(&self, dev: &Arc<Device>) -> String {
+ Self::dev_build_info(&mut self.build_info(), dev)
+ .log
+ .clone()
+ }
+
+ pub fn options(&self, dev: &Arc<Device>) -> String {
+ Self::dev_build_info(&mut self.build_info(), dev)
+ .options
+ .clone()
+ }
+
+ pub fn kernels(&self) -> Vec<String> {
+ self.build_info().kernels.clone()
+ }
+
+ pub fn compile(
+ &self,
+ dev: &Arc<Device>,
+ options: String,
+ headers: &[spirv::CLCHeader],
+ ) -> bool {
+ let mut info = self.build_info();
+ let d = Self::dev_build_info(&mut info, dev);
+ let args = prepare_options(&options);
+
+ let (spirv, log) = spirv::SPIRVBin::from_clc(&self.src, &args, headers, dev.cl_features());
+
+ d.spirv = spirv;
+ d.log = log;
+ d.options = options;
+
+ if d.spirv.is_some() {
+ d.status = CL_BUILD_SUCCESS as cl_build_status;
+ true
+ } else {
+ d.status = CL_BUILD_ERROR;
+ false
+ }
+ }
+
+ pub fn link(
+ context: Arc<Context>,
+ devs: &[Arc<Device>],
+ progs: &[Arc<Program>],
+ ) -> Arc<Program> {
+ let devs: Vec<Arc<Device>> = devs.iter().map(|d| (*d).clone()).collect();
+ let mut builds = HashMap::new();
+ let mut kernels = HashSet::new();
+ let mut locks: Vec<_> = progs.iter().map(|p| p.build_info()).collect();
+
+ for d in &devs {
+ let bins: Vec<_> = locks
+ .iter_mut()
+ .map(|l| Self::dev_build_info(l, d).spirv.as_ref().unwrap())
+ .collect();
+
+ let (spirv, log) = spirv::SPIRVBin::link(&bins, false);
+ let status = if let Some(spirv) = &spirv {
+ for k in spirv.kernels() {
+ kernels.insert(k);
+ }
+ CL_BUILD_SUCCESS as cl_build_status
+ } else {
+ CL_BUILD_ERROR
+ };
+
+ builds.insert(
+ d.clone(),
+ ProgramDevBuild {
+ spirv: spirv,
+ status: status,
+ log: log,
+ options: String::from(""),
+ },
+ );
+ }
+
+ Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ devs: devs,
+ src: CString::new("").unwrap(),
+ build: Mutex::new(ProgramBuild {
+ builds: builds,
+ kernels: kernels.into_iter().collect(),
+ }),
+ })
+ }
+}
--- /dev/null
+extern crate mesa_rust;
+extern crate rusticl_opencl_gen;
+
+use crate::api::icd::*;
+use crate::core::context::*;
+use crate::core::device::*;
+use crate::core::event::*;
+use crate::impl_cl_type_trait;
+
+use self::rusticl_opencl_gen::*;
+
+use std::sync::mpsc;
+use std::sync::Arc;
+use std::sync::Mutex;
+use std::thread;
+use std::thread::JoinHandle;
+
+#[repr(C)]
+pub struct Queue {
+ pub base: CLObjectBase<CL_INVALID_COMMAND_QUEUE>,
+ pub context: Arc<Context>,
+ pub device: Arc<Device>,
+ pub props: cl_command_queue_properties,
+ pending: Mutex<Vec<Arc<Event>>>,
+ _thrd: Option<JoinHandle<()>>,
+ chan_in: mpsc::Sender<Vec<Arc<Event>>>,
+ chan_out: mpsc::Receiver<bool>,
+}
+
+impl_cl_type_trait!(cl_command_queue, Queue, CL_INVALID_COMMAND_QUEUE);
+
+impl Queue {
+ pub fn new(
+ context: Arc<Context>,
+ device: Arc<Device>,
+ props: cl_command_queue_properties,
+ ) -> CLResult<Arc<Queue>> {
+ // we assume that memory allocation is the only possible failure. Any other failure reason
+ // should be detected earlier (e.g.: checking for CAPs).
+ let pipe = device.screen().create_context().unwrap();
+ let (tx_q, rx_t) = mpsc::channel::<Vec<Arc<Event>>>();
+ let (tx_t, rx_q) = mpsc::channel::<bool>();
+ Ok(Arc::new(Self {
+ base: CLObjectBase::new(),
+ context: context,
+ device: device,
+ props: props,
+ pending: Mutex::new(Vec::new()),
+ _thrd: Some(
+ thread::Builder::new()
+ .name("rusticl queue thread".into())
+ .spawn(move || loop {
+ let r = rx_t.recv();
+ if r.is_err() {
+ break;
+ }
+ for e in r.unwrap() {
+ e.call(&pipe);
+ }
+ if tx_t.send(true).is_err() {
+ break;
+ }
+ })
+ .unwrap(),
+ ),
+ chan_in: tx_q,
+ chan_out: rx_q,
+ }))
+ }
+
+ pub fn queue(&self, e: &Arc<Event>) {
+ self.pending.lock().unwrap().push(e.clone());
+ }
+
+ // TODO: implement non blocking flush
+ pub fn flush(&self, _wait: bool) -> CLResult<()> {
+ let mut p = self.pending.lock().unwrap();
+ // This should never ever error, but if it does return an error
+ self.chan_in
+ .send((*p).drain(0..).collect())
+ .map_err(|_| CL_OUT_OF_HOST_MEMORY)?;
+ self.chan_out.recv().unwrap();
+ Ok(())
+ }
+}
+
+impl Drop for Queue {
+ fn drop(&mut self) {
+ // when deleting the application side object, we have to flush
+ // From the OpenCL spec:
+ // clReleaseCommandQueue performs an implicit flush to issue any previously queued OpenCL
+ // commands in command_queue.
+ // TODO: maybe we have to do it on every release?
+ let _ = self.flush(true);
+ }
+}
--- /dev/null
+extern crate mesa_rust_gen;
+extern crate rusticl_opencl_gen;
+
+use self::mesa_rust_gen::*;
+use self::rusticl_opencl_gen::*;
+
+pub fn cl_mem_type_to_texture_target(mem_type: cl_mem_object_type) -> pipe_texture_target {
+ match mem_type {
+ CL_MEM_OBJECT_IMAGE1D => pipe_texture_target::PIPE_TEXTURE_1D,
+ CL_MEM_OBJECT_IMAGE2D => pipe_texture_target::PIPE_TEXTURE_2D,
+ CL_MEM_OBJECT_IMAGE3D => pipe_texture_target::PIPE_TEXTURE_3D,
+ CL_MEM_OBJECT_IMAGE1D_ARRAY => pipe_texture_target::PIPE_TEXTURE_1D_ARRAY,
+ CL_MEM_OBJECT_IMAGE2D_ARRAY => pipe_texture_target::PIPE_TEXTURE_2D_ARRAY,
+ CL_MEM_OBJECT_IMAGE1D_BUFFER => pipe_texture_target::PIPE_BUFFER,
+ _ => pipe_texture_target::PIPE_TEXTURE_2D,
+ }
+}
--- /dev/null
+extern crate rusticl_opencl_gen;
+
+use self::rusticl_opencl_gen::*;
+
+use std::convert::TryFrom;
+use std::os::raw::c_char;
+
+pub const CL1_0_VER: u32 = mk_cl_version(1, 0, 0);
+pub const CL1_1_VER: u32 = mk_cl_version(1, 1, 0);
+pub const CL1_2_VER: u32 = mk_cl_version(1, 2, 0);
+pub const CL2_0_VER: u32 = mk_cl_version(2, 0, 0);
+pub const CL2_1_VER: u32 = mk_cl_version(2, 1, 0);
+pub const CL2_2_VER: u32 = mk_cl_version(2, 2, 0);
+pub const CL3_0_VER: u32 = mk_cl_version(3, 0, 0);
+
+#[allow(dead_code)]
+#[repr(u32)]
+#[derive(Copy, Clone, Eq, Hash, PartialEq, PartialOrd, Ord)]
+pub enum CLVersion {
+ Cl1_0 = CL1_0_VER,
+ Cl1_1 = CL1_1_VER,
+ Cl1_2 = CL1_2_VER,
+ Cl2_0 = CL2_0_VER,
+ Cl2_1 = CL2_1_VER,
+ Cl2_2 = CL2_2_VER,
+ Cl3_0 = CL3_0_VER,
+}
+
+pub const fn mk_cl_version_ext(major: u32, minor: u32, patch: u32, ext: &str) -> cl_name_version {
+ let mut name: [c_char; 64] = [0; 64];
+ let ext = ext.as_bytes();
+
+ let mut i = 0;
+ while i < ext.len() {
+ name[i] = ext[i] as c_char;
+ i += 1;
+ }
+
+ cl_name_version {
+ version: mk_cl_version(major, minor, patch),
+ name,
+ }
+}
+
+pub const fn mk_cl_version(major: u32, minor: u32, patch: u32) -> u32 {
+ ((major & CL_VERSION_MAJOR_MASK) << (CL_VERSION_MINOR_BITS + CL_VERSION_PATCH_BITS))
+ | ((minor & CL_VERSION_MINOR_MASK) << CL_VERSION_PATCH_BITS)
+ | (patch & CL_VERSION_PATCH_MASK)
+}
+
+impl CLVersion {
+ pub fn api_str(&self) -> &'static str {
+ match self {
+ CLVersion::Cl1_0 => "1.0",
+ CLVersion::Cl1_1 => "1.1",
+ CLVersion::Cl1_2 => "1.2",
+ CLVersion::Cl2_0 => "2.0",
+ CLVersion::Cl2_1 => "2.1",
+ CLVersion::Cl2_2 => "2.2",
+ CLVersion::Cl3_0 => "3.0",
+ }
+ }
+}
+
+impl TryFrom<u32> for CLVersion {
+ type Error = cl_int;
+
+ fn try_from(value: u32) -> Result<Self, Self::Error> {
+ Ok(match value {
+ CL1_0_VER => CLVersion::Cl1_0,
+ CL1_1_VER => CLVersion::Cl1_1,
+ CL1_2_VER => CLVersion::Cl1_2,
+ CL2_0_VER => CLVersion::Cl2_0,
+ CL2_1_VER => CLVersion::Cl2_1,
+ CL2_2_VER => CLVersion::Cl2_2,
+ CL3_0_VER => CLVersion::Cl3_0,
+ _ => return Err(CL_INVALID_VALUE),
+ })
+ }
+}
--- /dev/null
+#![allow(clippy::collapsible_if)]
+#![allow(clippy::needless_range_loop)]
+
+mod api;
+pub mod core;
--- /dev/null
+pub mod spirv;
--- /dev/null
+extern crate mesa_rust_gen;
+extern crate mesa_rust_util;
+
+use self::mesa_rust_gen::*;
+use self::mesa_rust_util::string::*;
+
+use std::ffi::CString;
+use std::os::raw::c_char;
+use std::os::raw::c_void;
+use std::ptr;
+use std::slice;
+
+const INPUT_STR: *const c_char = b"input.cl\0" as *const u8 as *const c_char;
+
+pub struct SPIRVBin {
+ spirv: clc_binary,
+ info: Option<clc_parsed_spirv>,
+}
+
+pub struct CLCHeader<'a> {
+ pub name: CString,
+ pub source: &'a CString,
+}
+
+unsafe extern "C" fn msg_callback(data: *mut std::ffi::c_void, msg: *const c_char) {
+ let msgs = (data as *mut Vec<String>).as_mut().expect("");
+ msgs.push(c_string_to_string(msg));
+}
+
+impl SPIRVBin {
+ pub fn from_clc(
+ source: &CString,
+ args: &[CString],
+ headers: &[CLCHeader],
+ features: clc_optional_features,
+ ) -> (Option<Self>, String) {
+ let c_headers: Vec<_> = headers
+ .iter()
+ .map(|h| clc_named_value {
+ name: h.name.as_ptr(),
+ value: h.source.as_ptr(),
+ })
+ .collect();
+
+ let c_args: Vec<_> = args.iter().map(|a| a.as_ptr()).collect();
+
+ let args = clc_compile_args {
+ headers: c_headers.as_ptr(),
+ num_headers: c_headers.len() as u32,
+ source: clc_named_value {
+ name: INPUT_STR,
+ value: source.as_ptr(),
+ },
+ args: c_args.as_ptr(),
+ num_args: c_args.len() as u32,
+ spirv_version: clc_spirv_version::CLC_SPIRV_VERSION_MAX,
+ features: features,
+ allowed_spirv_extensions: ptr::null(),
+ };
+ let mut msgs: Vec<String> = Vec::new();
+ let logger = clc_logger {
+ priv_: &mut msgs as *mut Vec<String> as *mut c_void,
+ error: Some(msg_callback),
+ warning: Some(msg_callback),
+ };
+ let mut out = clc_binary::default();
+
+ let res = unsafe { clc_compile_c_to_spirv(&args, &logger, &mut out) };
+
+ let res = if res {
+ Some(SPIRVBin {
+ spirv: out,
+ info: None,
+ })
+ } else {
+ None
+ };
+ (res, msgs.join("\n"))
+ }
+
+ pub fn link(spirvs: &[&SPIRVBin], library: bool) -> (Option<Self>, String) {
+ let bins: Vec<_> = spirvs.iter().map(|s| &s.spirv as *const _).collect();
+
+ let linker_args = clc_linker_args {
+ in_objs: bins.as_ptr(),
+ num_in_objs: bins.len() as u32,
+ create_library: library as u32,
+ };
+
+ let mut msgs: Vec<String> = Vec::new();
+ let logger = clc_logger {
+ priv_: &mut msgs as *mut Vec<String> as *mut c_void,
+ error: Some(msg_callback),
+ warning: Some(msg_callback),
+ };
+
+ let mut out = clc_binary::default();
+ let res = unsafe { clc_link_spirv(&linker_args, &logger, &mut out) };
+
+ let info;
+ if !library {
+ let mut pspirv = clc_parsed_spirv::default();
+ let res = unsafe { clc_parse_spirv(&out, &logger, &mut pspirv) };
+
+ if res {
+ info = Some(pspirv);
+ } else {
+ info = None;
+ }
+ } else {
+ info = None;
+ }
+
+ let res = if res {
+ Some(SPIRVBin {
+ spirv: out,
+ info: info,
+ })
+ } else {
+ None
+ };
+ (res, msgs.join("\n"))
+ }
+
+ pub fn kernels(&self) -> Vec<String> {
+ unsafe {
+ match self.info {
+ None => Vec::new(),
+ Some(info) => slice::from_raw_parts(info.kernels, info.num_kernels as usize)
+ .iter()
+ .map(|i| i.name)
+ .map(c_string_to_string)
+ .collect(),
+ }
+ }
+ }
+}
+
+impl Drop for SPIRVBin {
+ fn drop(&mut self) {
+ unsafe {
+ clc_free_spirv(&mut self.spirv);
+ if let Some(info) = &mut self.info {
+ clc_free_parsed_spirv(info);
+ }
+ }
+ }
+}
--- /dev/null
+pub mod clc;
--- /dev/null
+pub mod compiler;
+pub mod pipe;
--- /dev/null
+extern crate mesa_rust_gen;
+
+use crate::pipe::resource::*;
+use crate::pipe::transfer::*;
+
+use self::mesa_rust_gen::*;
+
+use std::os::raw::*;
+use std::ptr;
+use std::ptr::*;
+use std::sync::Arc;
+
+pub struct PipeContext {
+ pipe: NonNull<pipe_context>,
+}
+
+unsafe impl Send for PipeContext {}
+unsafe impl Sync for PipeContext {}
+
+impl PipeContext {
+ pub(super) fn new(context: *mut pipe_context) -> Option<Arc<Self>> {
+ let s = Self {
+ pipe: NonNull::new(context)?,
+ };
+
+ if !has_required_cbs(unsafe { s.pipe.as_ref() }) {
+ assert!(false, "Context missing features. This should never happen!");
+ return None;
+ }
+
+ Some(Arc::new(s))
+ }
+
+ pub fn buffer_subdata(
+ &self,
+ res: &PipeResource,
+ offset: c_uint,
+ data: *const c_void,
+ size: c_uint,
+ ) {
+ unsafe {
+ self.pipe.as_ref().buffer_subdata.unwrap()(
+ self.pipe.as_ptr(),
+ res.pipe(),
+ pipe_map_flags::PIPE_MAP_WRITE.0, // TODO PIPE_MAP_x
+ offset,
+ size,
+ data,
+ )
+ }
+ }
+
+ pub fn buffer_map(
+ &self,
+ res: &PipeResource,
+ offset: i32,
+ size: i32,
+ block: bool,
+ ) -> PipeTransfer {
+ let mut b = pipe_box::default();
+ let mut out: *mut pipe_transfer = ptr::null_mut();
+
+ b.x = offset;
+ b.width = size;
+ b.height = 1;
+ b.depth = 1;
+
+ let flags = match block {
+ false => pipe_map_flags::PIPE_MAP_UNSYNCHRONIZED,
+ true => pipe_map_flags(0),
+ } | pipe_map_flags::PIPE_MAP_READ_WRITE;
+
+ let ptr = unsafe {
+ self.pipe.as_ref().buffer_map.unwrap()(
+ self.pipe.as_ptr(),
+ res.pipe(),
+ 0,
+ flags.0,
+ &b,
+ &mut out,
+ )
+ };
+
+ PipeTransfer::new(out, ptr)
+ }
+
+ pub(super) fn buffer_unmap(&self, tx: *mut pipe_transfer) {
+ unsafe { self.pipe.as_ref().buffer_unmap.unwrap()(self.pipe.as_ptr(), tx) };
+ }
+
+ pub fn blit(&self, src: &PipeResource, dst: &PipeResource) {
+ let mut blit_info = pipe_blit_info::default();
+ blit_info.src.resource = src.pipe();
+ blit_info.dst.resource = dst.pipe();
+
+ println!("blit not implemented!");
+
+ unsafe { self.pipe.as_ref().blit.unwrap()(self.pipe.as_ptr(), &blit_info) }
+ }
+}
+
+impl Drop for PipeContext {
+ fn drop(&mut self) {
+ unsafe {
+ self.pipe.as_ref().destroy.unwrap()(self.pipe.as_ptr());
+ }
+ }
+}
+
+fn has_required_cbs(c: &pipe_context) -> bool {
+ c.destroy.is_some()
+ && c.blit.is_some()
+ && c.buffer_map.is_some()
+ && c.buffer_subdata.is_some()
+ && c.buffer_unmap.is_some()
+}
--- /dev/null
+extern crate mesa_rust_gen;
+
+use crate::pipe::screen::*;
+
+use self::mesa_rust_gen::*;
+
+use std::ptr;
+use std::sync::Arc;
+
+#[derive(PartialEq)]
+pub(super) struct PipeLoaderDevice {
+ pub(super) ldev: *mut pipe_loader_device,
+}
+
+impl PipeLoaderDevice {
+ fn new(ldev: *mut pipe_loader_device) -> Option<Self> {
+ if ldev.is_null() {
+ return None;
+ }
+ Some(Self { ldev })
+ }
+
+ fn load_screen(self) -> Option<Arc<PipeScreen>> {
+ let s = unsafe { pipe_loader_create_screen(self.ldev) };
+ PipeScreen::new(self, s)
+ }
+}
+
+impl Drop for PipeLoaderDevice {
+ fn drop(&mut self) {
+ unsafe {
+ pipe_loader_release(&mut self.ldev, 1);
+ }
+ }
+}
+
+fn load_devs() -> Vec<PipeLoaderDevice> {
+ let n = unsafe { pipe_loader_probe(ptr::null_mut(), 0) };
+ let mut devices: Vec<*mut pipe_loader_device> = vec![ptr::null_mut(); n as usize];
+ unsafe {
+ pipe_loader_probe(devices.as_mut_ptr(), n);
+ }
+
+ devices
+ .into_iter()
+ .filter_map(PipeLoaderDevice::new)
+ .collect()
+}
+
+pub fn load_screens() -> Vec<Arc<PipeScreen>> {
+ load_devs()
+ .into_iter()
+ .filter_map(PipeLoaderDevice::load_screen)
+ .collect()
+}
--- /dev/null
+pub mod context;
+pub mod device;
+pub mod resource;
+pub mod screen;
+pub mod transfer;
--- /dev/null
+extern crate mesa_rust_gen;
+
+use self::mesa_rust_gen::*;
+
+use std::ptr;
+
+pub struct PipeResource {
+ pipe: *mut pipe_resource,
+}
+
+impl PipeResource {
+ pub fn new(res: *mut pipe_resource) -> Option<Self> {
+ if res.is_null() {
+ return None;
+ }
+
+ Some(Self { pipe: res })
+ }
+
+ pub(super) fn pipe(&self) -> *mut pipe_resource {
+ self.pipe
+ }
+}
+
+impl Drop for PipeResource {
+ fn drop(&mut self) {
+ unsafe { pipe_resource_reference(&mut self.pipe, ptr::null_mut()) }
+ }
+}
--- /dev/null
+extern crate mesa_rust_gen;
+extern crate mesa_rust_util;
+
+use crate::pipe::context::*;
+use crate::pipe::device::*;
+use crate::pipe::resource::*;
+
+use self::mesa_rust_gen::*;
+use self::mesa_rust_util::string::*;
+
+use std::convert::TryInto;
+use std::mem::size_of;
+use std::os::raw::c_void;
+use std::ptr;
+use std::sync::Arc;
+
+#[derive(PartialEq)]
+pub struct PipeScreen {
+ ldev: PipeLoaderDevice,
+ screen: *mut pipe_screen,
+}
+
+// until we have a better solution
+pub trait ComputeParam<T> {
+ fn compute_param(&self, cap: pipe_compute_cap) -> T;
+}
+
+macro_rules! compute_param_impl {
+ ($ty:ty) => {
+ impl ComputeParam<$ty> for PipeScreen {
+ fn compute_param(&self, cap: pipe_compute_cap) -> $ty {
+ let size = self.compute_param_wrapped(cap, ptr::null_mut());
+ let mut d = [0; size_of::<$ty>()];
+ assert_eq!(size as usize, d.len());
+ self.compute_param_wrapped(cap, d.as_mut_ptr().cast());
+ <$ty>::from_ne_bytes(d)
+ }
+ }
+ };
+}
+
+compute_param_impl!(u32);
+compute_param_impl!(u64);
+
+impl ComputeParam<Vec<u64>> for PipeScreen {
+ fn compute_param(&self, cap: pipe_compute_cap) -> Vec<u64> {
+ let size = self.compute_param_wrapped(cap, ptr::null_mut());
+ let elems = (size / 8) as usize;
+
+ let mut res: Vec<u64> = Vec::new();
+ let mut d: Vec<u8> = vec![0; size as usize];
+
+ self.compute_param_wrapped(cap, d.as_mut_ptr().cast());
+ for i in 0..elems {
+ let offset = i * 8;
+ let slice = &d[offset..offset + 8];
+ res.push(u64::from_ne_bytes(slice.try_into().expect("")));
+ }
+ res
+ }
+}
+
+impl PipeScreen {
+ pub(super) fn new(ldev: PipeLoaderDevice, screen: *mut pipe_screen) -> Option<Arc<Self>> {
+ if screen.is_null() || !has_required_cbs(screen) {
+ return None;
+ }
+
+ Some(Arc::new(Self { ldev, screen }))
+ }
+
+ pub fn create_context(self: &Arc<Self>) -> Option<Arc<PipeContext>> {
+ PipeContext::new(unsafe {
+ (*self.screen).context_create.unwrap()(
+ self.screen,
+ ptr::null_mut(),
+ PIPE_CONTEXT_COMPUTE_ONLY,
+ )
+ })
+ }
+
+ pub fn resource_create_buffer(&self, size: u32) -> Option<PipeResource> {
+ let mut tmpl = pipe_resource::default();
+
+ tmpl.set_target(pipe_texture_target::PIPE_BUFFER);
+ tmpl.width0 = size;
+ tmpl.height0 = 1;
+ tmpl.depth0 = 1;
+ tmpl.array_size = 1;
+
+ PipeResource::new(unsafe { (*self.screen).resource_create.unwrap()(self.screen, &tmpl) })
+ }
+
+ pub fn resource_create_buffer_from_user(
+ &self,
+ size: u32,
+ mem: *mut c_void,
+ ) -> Option<PipeResource> {
+ let mut tmpl = pipe_resource::default();
+
+ tmpl.set_target(pipe_texture_target::PIPE_BUFFER);
+ tmpl.width0 = size;
+ tmpl.height0 = 1;
+ tmpl.depth0 = 1;
+ tmpl.array_size = 1;
+
+ PipeResource::new(unsafe {
+ (*self.screen).resource_from_user_memory.unwrap()(self.screen, &tmpl, mem)
+ })
+ }
+
+ pub fn param(&self, cap: pipe_cap) -> i32 {
+ unsafe { (*self.screen).get_param.unwrap()(self.screen, cap) }
+ }
+
+ pub fn shader_param(&self, t: pipe_shader_type, cap: pipe_shader_cap) -> i32 {
+ unsafe { (*self.screen).get_shader_param.unwrap()(self.screen, t, cap) }
+ }
+
+ fn compute_param_wrapped(&self, cap: pipe_compute_cap, ptr: *mut c_void) -> i32 {
+ let s = &mut unsafe { *self.screen };
+ unsafe {
+ s.get_compute_param.unwrap()(self.screen, pipe_shader_ir::PIPE_SHADER_IR_NIR, cap, ptr)
+ }
+ }
+
+ pub fn name(&self) -> String {
+ unsafe {
+ let s = *self.screen;
+ c_string_to_string(s.get_name.unwrap()(self.screen))
+ }
+ }
+
+ pub fn device_vendor(&self) -> String {
+ unsafe {
+ let s = *self.screen;
+ c_string_to_string(s.get_device_vendor.unwrap()(self.screen))
+ }
+ }
+
+ pub fn device_type(&self) -> pipe_loader_device_type {
+ unsafe { *self.ldev.ldev }.type_
+ }
+
+ pub fn is_format_supported(
+ &self,
+ format: pipe_format,
+ target: pipe_texture_target,
+ bindings: u32,
+ ) -> bool {
+ let s = &mut unsafe { *self.screen };
+ unsafe { s.is_format_supported.unwrap()(self.screen, format, target, 0, 0, bindings) }
+ }
+}
+
+impl Drop for PipeScreen {
+ fn drop(&mut self) {
+ unsafe {
+ (*self.screen).destroy.unwrap()(self.screen);
+ }
+ }
+}
+
+fn has_required_cbs(screen: *mut pipe_screen) -> bool {
+ let s = unsafe { *screen };
+ s.context_create.is_some()
+ && s.destroy.is_some()
+ && s.get_compute_param.is_some()
+ && s.get_name.is_some()
+ && s.get_param.is_some()
+ && s.get_shader_param.is_some()
+ && s.is_format_supported.is_some()
+ && s.resource_create.is_some()
+ && s.resource_from_user_memory.is_some()
+}
--- /dev/null
+extern crate mesa_rust_gen;
+
+use crate::pipe::context::*;
+
+use self::mesa_rust_gen::*;
+
+use std::ops::Deref;
+use std::os::raw::c_void;
+use std::ptr;
+
+pub struct PipeTransfer {
+ pipe: *mut pipe_transfer,
+ res: *mut pipe_resource,
+ ptr: *mut c_void,
+}
+
+pub struct GuardedPipeTransfer<'a> {
+ inner: PipeTransfer,
+ ctx: &'a PipeContext,
+}
+
+impl<'a> Deref for GuardedPipeTransfer<'a> {
+ type Target = PipeTransfer;
+
+ fn deref(&self) -> &Self::Target {
+ &self.inner
+ }
+}
+
+impl<'a> Drop for GuardedPipeTransfer<'a> {
+ fn drop(&mut self) {
+ self.ctx.buffer_unmap(self.inner.pipe);
+ unsafe { pipe_resource_reference(&mut self.inner.res, ptr::null_mut()) };
+ }
+}
+
+impl PipeTransfer {
+ pub(super) fn new(pipe: *mut pipe_transfer, ptr: *mut c_void) -> Self {
+ let mut res: *mut pipe_resource = ptr::null_mut();
+ unsafe { pipe_resource_reference(&mut res, (*pipe).resource) }
+
+ Self {
+ pipe: pipe,
+ res: res,
+ ptr: ptr,
+ }
+ }
+
+ pub fn ptr(&self) -> *mut c_void {
+ self.ptr
+ }
+
+ pub fn with_ctx(self, ctx: &PipeContext) -> GuardedPipeTransfer {
+ GuardedPipeTransfer {
+ inner: self,
+ ctx: ctx,
+ }
+ }
+}
+
+// use set_ctx before operating on the PipeTransfer inside a block where it gets droped
+impl Drop for PipeTransfer {
+ fn drop(&mut self) {
+ assert_eq!(ptr::null_mut(), self.res);
+ }
+}
--- /dev/null
+# Copyright ©
+
+# 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.
+
+rust = import('unstable-rust')
+
+libmesa_rust_util_files = files(
+ 'util/lib.rs',
+ 'util/properties.rs',
+ 'util/ptr.rs',
+ 'util/string.rs',
+)
+
+libmesa_rust_files = files(
+ 'mesa/lib.rs',
+ 'mesa/compiler/mod.rs',
+ 'mesa/compiler/clc/mod.rs',
+ 'mesa/compiler/clc/spirv.rs',
+ 'mesa/pipe/context.rs',
+ 'mesa/pipe/device.rs',
+ 'mesa/pipe/mod.rs',
+ 'mesa/pipe/screen.rs',
+ 'mesa/pipe/transfer.rs',
+)
+
+rusticl_files = files(
+ 'lib.rs',
+ 'api/context.rs',
+ 'api/device.rs',
+ 'api/event.rs',
+ 'api/icd.rs',
+ 'api/kernel.rs',
+ 'api/memory.rs',
+ 'api/mod.rs',
+ 'api/platform.rs',
+ 'api/program.rs',
+ 'api/queue.rs',
+ 'api/types.rs',
+ 'api/util.rs',
+ 'core/context.rs',
+ 'core/device.rs',
+ 'core/format.rs',
+ 'core/kernel.rs',
+ 'core/memory.rs',
+ 'core/mod.rs',
+ 'core/program.rs',
+ 'core/queue.rs',
+ 'core/util.rs',
+ 'core/version.rs',
+)
+
+rusticl_args = [
+ # we error on all clippy warnings unless they are disabled
+ '-Dclippy::all',
+ # we want to add asserts in control flow
+ '-Aclippy::assertions_on_constants',
+ # warns on Arc<_> as keys
+ '-Aclippy::mutable_key_type',
+ '-Aclippy::not_unsafe_ptr_arg_deref',
+ # dunno, kind of looks nicier being explicit
+ '-Aclippy::redundant_field_names',
+ '-Aclippy::too_many_arguments',
+ '-Aclippy::type_complexity',
+]
+
+rusticl_gen_args = [
+ # can't do anything about it anyway
+ '-Aclippy::all',
+]
+
+rusticl_bindgen_args = [
+ '--no-convert-floats',
+ '--use-array-pointers-in-arguments',
+ '--size_t-is-usize',
+ '--default-enum-style', 'rust',
+ '--with-derive-partialeq',
+ '--with-derive-eq',
+ '--with-derive-partialord',
+ '--with-derive-ord',
+ '--with-derive-hash',
+ '--with-derive-default',
+ '--anon-fields-prefix', 'anon_',
+]
+
+rusticl_opencl_bindings_rs = rust.bindgen(
+ input : [
+ 'rusticl_opencl_bindings.h',
+ opencl_headers,
+ ],
+ output : 'rusticl_opencl_bindings.rs',
+ include_directories : [
+ inc_include,
+ ],
+ c_args : [
+ '-DCL_USE_DEPRECATED_OPENCL_1_0_APIS',
+ '-DCL_USE_DEPRECATED_OPENCL_1_1_APIS',
+ '-DCL_USE_DEPRECATED_OPENCL_1_2_APIS',
+ '-DCL_USE_DEPRECATED_OPENCL_2_0_APIS',
+ '-DCL_USE_DEPRECATED_OPENCL_2_1_APIS',
+ '-DCL_USE_DEPRECATED_OPENCL_2_2_APIS',
+ '-DCL_TARGET_OPENCL_VERSION=300',
+ ],
+ args : [
+ rusticl_bindgen_args,
+ '--disable-header-comment',
+ '--ignore-functions',
+ # needed because bindgen adds *mut void fields...
+ '--raw-line', 'unsafe impl std::marker::Sync for _cl_icd_dispatch {}',
+ '--whitelist-type', 'cl_.*',
+ '--whitelist-var', 'CL_.*',
+ # some info types need to be strongly typed so we can implement various get_infos
+ '--new-type-alias-deref', 'cl_(mem|image)_info',
+ ],
+)
+
+rusticl_opencl_gen = static_library(
+ 'rusticl_opencl_gen',
+ rusticl_opencl_bindings_rs,
+ gnu_symbol_visibility : 'hidden',
+ rust_crate_type : 'rlib',
+ rust_args : [
+ rusticl_gen_args,
+ '-Anon_snake_case',
+ '-Anon_camel_case_types',
+ '-Anon_upper_case_globals',
+ ],
+)
+
+rusticl_mesa_bindings_inline_wrapper = static_library(
+ 'mesa_bindings_inline_wrapper',
+ [
+ 'rusticl_mesa_inline_bindings_wrapper.c',
+ 'rusticl_mesa_inline_bindings_wrapper.h'
+ ],
+ gnu_symbol_visibility : 'hidden',
+ include_directories : [
+ inc_gallium,
+ inc_gallium_aux,
+ inc_include,
+ inc_src,
+ ],
+ c_args : pre_args,
+ dependencies: [
+ idep_nir_headers,
+ ],
+)
+
+rusticl_mesa_bindings_rs = rust.bindgen(
+ input : 'rusticl_mesa_bindings.h',
+ output : 'rusticl_mesa_bindings.rs',
+ include_directories : [
+ inc_gallium,
+ inc_gallium_aux,
+ inc_include,
+ inc_src,
+ ],
+ c_args : pre_args,
+ args : [
+ rusticl_bindgen_args,
+ '--whitelist-function', 'clc_.*',
+ '--whitelist-function', 'pipe_.*',
+ '--whitelist-type', 'pipe_endian',
+ '--whitelist-var', 'PIPE_.*',
+ '--bitfield-enum', 'pipe_map_flags',
+ ],
+)
+
+idep_rusticl_gen = declare_dependency(
+ sources: [
+ rusticl_opencl_bindings_rs,
+ ],
+)
+
+libmesa_rust_gen = static_library(
+ 'mesa_rust_gen',
+ rusticl_mesa_bindings_rs,
+ gnu_symbol_visibility : 'hidden',
+ link_with: [
+ libgallium,
+ ],
+ dependencies: [
+ idep_clc,
+ ],
+ rust_crate_type : 'rlib',
+ rust_args : [
+ rusticl_gen_args,
+ '-Anon_snake_case',
+ '-Anon_camel_case_types',
+ '-Anon_upper_case_globals',
+ ],
+)
+
+libmesa_rust_util = static_library(
+ 'mesa_rust_util',
+ [libmesa_rust_util_files],
+ gnu_symbol_visibility : 'hidden',
+ rust_crate_type : 'rlib',
+ rust_args : [
+ rusticl_args,
+ ],
+)
+
+libmesa_rust = static_library(
+ 'mesa_rust',
+ [libmesa_rust_files],
+ gnu_symbol_visibility : 'hidden',
+ rust_crate_type : 'rlib',
+ rust_args : [
+ rusticl_args,
+ ],
+ link_with : [
+ libmesa_rust_gen,
+ libmesa_rust_util,
+ rusticl_mesa_bindings_inline_wrapper,
+ ]
+)
+
+librusticl = static_library(
+ 'rusticl',
+ [rusticl_files],
+ gnu_symbol_visibility : 'hidden',
+ rust_crate_type : 'staticlib',
+ rust_args : [
+ rusticl_args,
+ ],
+ link_with : [
+ libmesa_rust,
+ libmesa_rust_util,
+ rusticl_opencl_gen,
+ ],
+ dependencies : [
+ idep_rusticl_gen,
+ ],
+)
--- /dev/null
+#include "rusticl_mesa_inline_bindings_wrapper.h"
+
+#include "compiler/clc/clc.h"
+
+#include "pipe/p_context.h"
+#include "pipe/p_defines.h"
+#include "pipe/p_screen.h"
+#include "pipe/p_state.h"
+#include "pipe-loader/pipe_loader.h"
--- /dev/null
+#include "rusticl_mesa_inline_bindings_wrapper.h"
+
+void
+pipe_resource_reference(struct pipe_resource **dst, struct pipe_resource *src)
+{
+ __pipe_resource_reference_wraped(dst, src);
+}
--- /dev/null
+#define pipe_resource_reference __pipe_resource_reference_wraped
+#include "util/u_inlines.h"
+#undef pipe_resource_reference
+
+void pipe_resource_reference(struct pipe_resource **dst, struct pipe_resource *src);
--- /dev/null
+#include <CL/cl_icd.h>
+
+#define DECL_CL_STRUCT(name) struct name { const cl_icd_dispatch *dispatch; }
+DECL_CL_STRUCT(_cl_command_queue);
+DECL_CL_STRUCT(_cl_context);
+DECL_CL_STRUCT(_cl_device_id);
+DECL_CL_STRUCT(_cl_event);
+DECL_CL_STRUCT(_cl_kernel);
+DECL_CL_STRUCT(_cl_mem);
+DECL_CL_STRUCT(_cl_platform_id);
+DECL_CL_STRUCT(_cl_program);
+DECL_CL_STRUCT(_cl_sampler);
+#undef DECL_CL_STRUCT
--- /dev/null
+pub mod properties;
+pub mod ptr;
+pub mod string;
--- /dev/null
+pub struct Properties<T> {
+ pub props: Vec<(T, T)>,
+}
+
+impl<T: Copy + PartialEq + Default> Properties<T> {
+ #[allow(clippy::not_unsafe_ptr_arg_deref)]
+ pub fn from_ptr_raw(mut p: *const T) -> Vec<T> {
+ let mut res: Vec<T> = Vec::new();
+
+ if !p.is_null() {
+ unsafe {
+ while *p != T::default() {
+ res.push(*p);
+ res.push(*p.add(1));
+ p = p.add(2);
+ }
+ }
+ res.push(T::default());
+ }
+
+ res
+ }
+
+ #[allow(clippy::not_unsafe_ptr_arg_deref)]
+ pub fn from_ptr(mut p: *const T) -> Option<Self> {
+ let mut res = Self::default();
+
+ if !p.is_null() {
+ let mut k: Vec<T> = Vec::new();
+ let mut v: Vec<T> = Vec::new();
+
+ unsafe {
+ while *p != T::default() {
+ if k.contains(&*p) {
+ return None;
+ }
+ k.push(*p);
+ v.push(*p.add(1));
+ p = p.add(2);
+ }
+ }
+
+ res.props = k.iter().cloned().zip(v).collect();
+ }
+
+ Some(res)
+ }
+}
+
+impl<T> Default for Properties<T> {
+ fn default() -> Self {
+ Self { props: Vec::new() }
+ }
+}
--- /dev/null
+use std::ptr;
+
+pub trait CheckedPtr<T> {
+ /// # Safety
+ ///
+ /// besides a null check the function can't make sure the pointer is valid
+ /// for the entire size
+ unsafe fn copy_checked(self, val: *const T, size: usize);
+ fn write_checked(self, val: T);
+}
+
+impl<T> CheckedPtr<T> for *mut T {
+ unsafe fn copy_checked(self, val: *const T, size: usize) {
+ if !self.is_null() {
+ ptr::copy(val, self, size);
+ }
+ }
+
+ fn write_checked(self, val: T) {
+ if !self.is_null() {
+ unsafe {
+ *self = val;
+ }
+ }
+ }
+}
--- /dev/null
+use std::ffi::CStr;
+use std::os::raw::c_char;
+
+#[allow(clippy::not_unsafe_ptr_arg_deref)]
+pub fn c_string_to_string(cstr: *const c_char) -> String {
+ if cstr.is_null() {
+ return String::from("");
+ }
+
+ let res = unsafe { CStr::from_ptr(cstr).to_str() };
+ assert!(res.is_ok());
+ String::from(res.unwrap_or(""))
+}
subdir('frontends/clover')
subdir('targets/opencl')
endif
+if with_gallium_rusticl
+ subdir('frontends/rusticl')
+ subdir('targets/rusticl')
+endif
if with_dri
subdir('frontends/dri')
subdir('targets/dri')
--- /dev/null
+# Copyright © 2017 Intel Corporation
+
+# 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.
+
+opencl_version = '1'
+
+librusticl = shared_library(
+ 'RusticlOpenCL',
+ 'target.c',
+ gnu_symbol_visibility : 'hidden',
+ include_directories : [
+ inc_include,
+ inc_src,
+ inc_gallium,
+ inc_gallium_aux,
+ inc_gallium_drivers,
+ inc_gallium_winsys,
+ ],
+ link_whole : librusticl,
+ link_with : [
+ libgalliumvl,
+ libpipe_loader_static,
+ libswdri,
+ libswkmsdri,
+ libwsw,
+ libws_null,
+ ],
+ dependencies : [
+ driver_iris,
+ driver_nouveau,
+ driver_swrast,
+ idep_nir,
+ ],
+ version : '@0@.0.0'.format(opencl_version),
+ install : true,
+)
+
+_config = configuration_data()
+_config.set('OPENCL_VERSION', opencl_version)
+configure_file(
+ configuration : _config,
+ input : 'rusticl.icd.in',
+ output : 'rusticl.icd',
+ install : true,
+ install_dir : join_paths(get_option('sysconfdir'), 'OpenCL', 'vendors'),
+)
--- /dev/null
+libRusticlOpenCL.so.@OPENCL_VERSION@
--- /dev/null
+#include "target-helpers/drm_helper.h"
+#include "target-helpers/sw_helper.h"
projects / platform / upstream / mesa.git / commitdiff