From 20c90fed5a0ab0202ee1ef474c71cb816164a448 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sun, 8 Nov 2020 20:28:21 +0100 Subject: [PATCH] rusticl: added Initial code drop for Rusticl :) Signed-off-by: Karol Herbst Acked-by: Alyssa Rosenzweig Part-of: --- .editorconfig | 5 +- include/meson.build | 37 +- meson.build | 17 +- meson_options.txt | 6 + src/gallium/frontends/rusticl/api/context.rs | 117 ++ src/gallium/frontends/rusticl/api/device.rs | 230 ++++ src/gallium/frontends/rusticl/api/event.rs | 58 + src/gallium/frontends/rusticl/api/icd.rs | 1260 ++++++++++++++++++ src/gallium/frontends/rusticl/api/kernel.rs | 24 + src/gallium/frontends/rusticl/api/memory.rs | 1343 ++++++++++++++++++++ src/gallium/frontends/rusticl/api/mod.rs | 11 + src/gallium/frontends/rusticl/api/platform.rs | 121 ++ src/gallium/frontends/rusticl/api/program.rs | 287 +++++ src/gallium/frontends/rusticl/api/queue.rs | 77 ++ src/gallium/frontends/rusticl/api/types.rs | 147 +++ src/gallium/frontends/rusticl/api/util.rs | 341 +++++ src/gallium/frontends/rusticl/core/context.rs | 62 + src/gallium/frontends/rusticl/core/device.rs | 489 +++++++ src/gallium/frontends/rusticl/core/event.rs | 110 ++ src/gallium/frontends/rusticl/core/format.rs | 181 +++ src/gallium/frontends/rusticl/core/kernel.rs | 24 + src/gallium/frontends/rusticl/core/memory.rs | 398 ++++++ src/gallium/frontends/rusticl/core/mod.rs | 10 + src/gallium/frontends/rusticl/core/program.rs | 187 +++ src/gallium/frontends/rusticl/core/queue.rs | 96 ++ src/gallium/frontends/rusticl/core/util.rs | 17 + src/gallium/frontends/rusticl/core/version.rs | 80 ++ src/gallium/frontends/rusticl/lib.rs | 5 + .../frontends/rusticl/mesa/compiler/clc/mod.rs | 1 + .../frontends/rusticl/mesa/compiler/clc/spirv.rs | 148 +++ src/gallium/frontends/rusticl/mesa/compiler/mod.rs | 1 + src/gallium/frontends/rusticl/mesa/lib.rs | 2 + src/gallium/frontends/rusticl/mesa/pipe/context.rs | 116 ++ src/gallium/frontends/rusticl/mesa/pipe/device.rs | 55 + src/gallium/frontends/rusticl/mesa/pipe/mod.rs | 5 + .../frontends/rusticl/mesa/pipe/resource.rs | 29 + src/gallium/frontends/rusticl/mesa/pipe/screen.rs | 175 +++ .../frontends/rusticl/mesa/pipe/transfer.rs | 66 + src/gallium/frontends/rusticl/meson.build | 250 ++++ .../frontends/rusticl/rusticl_mesa_bindings.h | 9 + .../rusticl/rusticl_mesa_inline_bindings_wrapper.c | 7 + .../rusticl/rusticl_mesa_inline_bindings_wrapper.h | 5 + .../frontends/rusticl/rusticl_opencl_bindings.h | 13 + src/gallium/frontends/rusticl/util/lib.rs | 3 + src/gallium/frontends/rusticl/util/properties.rs | 54 + src/gallium/frontends/rusticl/util/ptr.rs | 26 + src/gallium/frontends/rusticl/util/string.rs | 13 + src/gallium/meson.build | 4 + src/gallium/targets/rusticl/meson.build | 62 + src/gallium/targets/rusticl/rusticl.icd.in | 1 + src/gallium/targets/rusticl/target.c | 2 + 51 files changed, 6768 insertions(+), 19 deletions(-) create mode 100644 src/gallium/frontends/rusticl/api/context.rs create mode 100644 src/gallium/frontends/rusticl/api/device.rs create mode 100644 src/gallium/frontends/rusticl/api/event.rs create mode 100644 src/gallium/frontends/rusticl/api/icd.rs create mode 100644 src/gallium/frontends/rusticl/api/kernel.rs create mode 100644 src/gallium/frontends/rusticl/api/memory.rs create mode 100644 src/gallium/frontends/rusticl/api/mod.rs create mode 100644 src/gallium/frontends/rusticl/api/platform.rs create mode 100644 src/gallium/frontends/rusticl/api/program.rs create mode 100644 src/gallium/frontends/rusticl/api/queue.rs create mode 100644 src/gallium/frontends/rusticl/api/types.rs create mode 100644 src/gallium/frontends/rusticl/api/util.rs create mode 100644 src/gallium/frontends/rusticl/core/context.rs create mode 100644 src/gallium/frontends/rusticl/core/device.rs create mode 100644 src/gallium/frontends/rusticl/core/event.rs create mode 100644 src/gallium/frontends/rusticl/core/format.rs create mode 100644 src/gallium/frontends/rusticl/core/kernel.rs create mode 100644 src/gallium/frontends/rusticl/core/memory.rs create mode 100644 src/gallium/frontends/rusticl/core/mod.rs create mode 100644 src/gallium/frontends/rusticl/core/program.rs create mode 100644 src/gallium/frontends/rusticl/core/queue.rs create mode 100644 src/gallium/frontends/rusticl/core/util.rs create mode 100644 src/gallium/frontends/rusticl/core/version.rs create mode 100644 src/gallium/frontends/rusticl/lib.rs create mode 100644 src/gallium/frontends/rusticl/mesa/compiler/clc/mod.rs create mode 100644 src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs create mode 100644 src/gallium/frontends/rusticl/mesa/compiler/mod.rs create mode 100644 src/gallium/frontends/rusticl/mesa/lib.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/context.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/device.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/mod.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/resource.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/screen.rs create mode 100644 src/gallium/frontends/rusticl/mesa/pipe/transfer.rs create mode 100644 src/gallium/frontends/rusticl/meson.build create mode 100644 src/gallium/frontends/rusticl/rusticl_mesa_bindings.h create mode 100644 src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.c create mode 100644 src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.h create mode 100644 src/gallium/frontends/rusticl/rusticl_opencl_bindings.h create mode 100644 src/gallium/frontends/rusticl/util/lib.rs create mode 100644 src/gallium/frontends/rusticl/util/properties.rs create mode 100644 src/gallium/frontends/rusticl/util/ptr.rs create mode 100644 src/gallium/frontends/rusticl/util/string.rs create mode 100644 src/gallium/targets/rusticl/meson.build create mode 100644 src/gallium/targets/rusticl/rusticl.icd.in create mode 100644 src/gallium/targets/rusticl/target.c diff --git a/.editorconfig b/.editorconfig index f4f0566..6909880 100644 --- a/.editorconfig +++ b/.editorconfig @@ -35,7 +35,10 @@ trim_trailing_whitespace = false indent_style = space indent_size = 2 - [*.ps1] indent_style = space indent_size = 2 + +[*.rs] +indent_style = space +indent_size = 4 diff --git a/include/meson.build b/include/meson.build index 46804f7..5cfbef5 100644 --- a/include/meson.build +++ b/include/meson.build @@ -118,27 +118,30 @@ if with_platform_haiku ) 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 diff --git a/meson.build b/meson.build index 4c6c4ff..2acc449 100644 --- a/meson.build +++ b/meson.build @@ -27,7 +27,7 @@ project( ).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 @@ -953,6 +953,21 @@ else 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') diff --git a/meson_options.txt b/meson_options.txt index 283a02b..e83d476 100644 --- a/meson_options.txt +++ b/meson_options.txt @@ -164,6 +164,12 @@ option( 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', diff --git a/src/gallium/frontends/rusticl/api/context.rs b/src/gallium/frontends/rusticl/api/context.rs new file mode 100644 index 0000000..5eb49a5 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/context.rs @@ -0,0 +1,117 @@ +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 for cl_context { + fn query(&self, q: cl_context_info) -> CLResult> { + let ctx = self.get_ref()?; + Ok(match q { + CL_CONTEXT_DEVICES => { + cl_prop::<&Vec>( + &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::(ctx.devs.len() as u32), + CL_CONTEXT_PROPERTIES => cl_prop::<&Vec>(&ctx.properties), + CL_CONTEXT_REFERENCE_COUNT => cl_prop::(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, + user_data: *mut ::std::os::raw::c_void, +) -> CLResult { + 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, + user_data: *mut ::std::os::raw::c_void, +) -> CLResult { + // 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, + ) +} diff --git a/src/gallium/frontends/rusticl/api/device.rs b/src/gallium/frontends/rusticl/api/device.rs new file mode 100644 index 0000000..7a57194 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/device.rs @@ -0,0 +1,230 @@ +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 for cl_device_id { + fn query(&self, q: cl_device_info) -> CLResult> { + let dev = self.get_ref()?; + Ok(match q { + CL_DEVICE_ADDRESS_BITS => cl_prop::(dev.address_bits()), + CL_DEVICE_ATOMIC_FENCE_CAPABILITIES => cl_prop::(0), + CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES => cl_prop::(0), + CL_DEVICE_AVAILABLE => cl_prop::(true), + CL_DEVICE_BUILT_IN_KERNELS => cl_prop::<&str>(""), + CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION => cl_prop::>(Vec::new()), + CL_DEVICE_COMPILER_AVAILABLE => cl_prop::(true), + CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES => { + cl_prop::(0) + } + CL_DEVICE_DOUBLE_FP_CONFIG => cl_prop::(0), + CL_DEVICE_ENDIAN_LITTLE => cl_prop::(dev.little_endian()), + CL_DEVICE_ERROR_CORRECTION_SUPPORT => cl_prop::(false), + CL_DEVICE_EXECUTION_CAPABILITIES => { + cl_prop::(CL_EXEC_KERNEL.into()) + } + CL_DEVICE_EXTENSIONS => cl_prop::<&str>(&dev.extension_string), + CL_DEVICE_EXTENSIONS_WITH_VERSION => cl_prop::<&Vec>(&dev.extensions), + CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT => cl_prop::(false), + CL_DEVICE_GLOBAL_MEM_CACHE_TYPE => cl_prop::(CL_NONE), + CL_DEVICE_GLOBAL_MEM_CACHE_SIZE => cl_prop::(0), + CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE => cl_prop::(0), + CL_DEVICE_GLOBAL_MEM_SIZE => cl_prop::(dev.global_mem_size()), + CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE => cl_prop::(0), + CL_DEVICE_HALF_FP_CONFIG => cl_prop::(0), + CL_DEVICE_HOST_UNIFIED_MEMORY => cl_prop::(dev.unified_memory()), + CL_DEVICE_IL_VERSION => cl_prop::<&str>(""), + CL_DEVICE_ILS_WITH_VERSION => cl_prop::>(Vec::new()), + CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT => { + cl_prop::(dev.image_base_address_alignment()) + } + CL_DEVICE_IMAGE_MAX_ARRAY_SIZE => cl_prop::(dev.image_array_size()), + CL_DEVICE_IMAGE_MAX_BUFFER_SIZE => cl_prop::(dev.image_buffer_size()), + CL_DEVICE_IMAGE_PITCH_ALIGNMENT => cl_prop::(0), + CL_DEVICE_IMAGE_SUPPORT => cl_prop::(dev.image_supported()), + CL_DEVICE_IMAGE2D_MAX_HEIGHT => cl_prop::(dev.image_2d_size()), + CL_DEVICE_IMAGE2D_MAX_WIDTH => cl_prop::(dev.image_2d_size()), + CL_DEVICE_IMAGE3D_MAX_HEIGHT => cl_prop::(dev.image_3d_size()), + CL_DEVICE_IMAGE3D_MAX_WIDTH => cl_prop::(dev.image_3d_size()), + CL_DEVICE_IMAGE3D_MAX_DEPTH => cl_prop::(dev.image_3d_size()), + CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED => cl_prop::<&str>("v0000-01-01-00"), + CL_DEVICE_LINKER_AVAILABLE => cl_prop::(true), + CL_DEVICE_LOCAL_MEM_SIZE => cl_prop::(dev.local_mem_size()), + // TODO add query for CL_LOCAL vs CL_GLOBAL + CL_DEVICE_LOCAL_MEM_TYPE => cl_prop::(CL_GLOBAL), + CL_DEVICE_MAX_CLOCK_FREQUENCY => cl_prop::(dev.max_clock_freq()), + CL_DEVICE_MAX_COMPUTE_UNITS => cl_prop::(dev.max_compute_units()), + // TODO atm implemented as mem_const + CL_DEVICE_MAX_CONSTANT_ARGS => cl_prop::(1024), + CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE => cl_prop::(dev.const_max_size()), + CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE => cl_prop::(0), + CL_DEVICE_MAX_MEM_ALLOC_SIZE => cl_prop::(dev.max_mem_alloc()), + CL_DEVICE_MAX_NUM_SUB_GROUPS => cl_prop::(0), + CL_DEVICE_MAX_ON_DEVICE_EVENTS => cl_prop::(0), + CL_DEVICE_MAX_ON_DEVICE_QUEUES => cl_prop::(0), + CL_DEVICE_MAX_PARAMETER_SIZE => cl_prop::(dev.param_max_size()), + CL_DEVICE_MAX_PIPE_ARGS => cl_prop::(0), + CL_DEVICE_MAX_READ_IMAGE_ARGS => cl_prop::(dev.image_read_count()), + CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS => cl_prop::(0), + CL_DEVICE_MAX_SAMPLERS => cl_prop::(dev.max_samplers()), + CL_DEVICE_MAX_WORK_GROUP_SIZE => cl_prop::(dev.max_threads_per_block()), + CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS => cl_prop::(dev.max_grid_dimensions()), + CL_DEVICE_MAX_WORK_ITEM_SIZES => cl_prop::>(dev.max_block_sizes()), + CL_DEVICE_MAX_WRITE_IMAGE_ARGS => cl_prop::(dev.image_write_count()), + // TODO proper retrival from devices + CL_DEVICE_MEM_BASE_ADDR_ALIGN => cl_prop::(0x1000), + CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE => { + cl_prop::(size_of::() as cl_uint) + } + CL_DEVICE_NAME => cl_prop(dev.screen().name()), + CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR => cl_prop::(1), + CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE => cl_prop::(0), + CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT => cl_prop::(1), + CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF => cl_prop::(0), + CL_DEVICE_NATIVE_VECTOR_WIDTH_INT => cl_prop::(1), + CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG => cl_prop::(1), + CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT => cl_prop::(1), + CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT => cl_prop::(false), + CL_DEVICE_NUMERIC_VERSION => cl_prop::(dev.cl_version as cl_version), + // TODO subdevice support + CL_DEVICE_PARENT_DEVICE => cl_prop::(cl_device_id::from_ptr(ptr::null())), + CL_DEVICE_PARTITION_AFFINITY_DOMAIN => cl_prop::(0), + CL_DEVICE_PARTITION_MAX_SUB_DEVICES => cl_prop::(0), + CL_DEVICE_PARTITION_PROPERTIES => cl_prop::>(vec![0]), + CL_DEVICE_PARTITION_TYPE => cl_prop::>(Vec::new()), + CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS => cl_prop::(0), + CL_DEVICE_PIPE_MAX_PACKET_SIZE => cl_prop::(0), + CL_DEVICE_PIPE_SUPPORT => cl_prop::(false), + CL_DEVICE_PLATFORM => cl_prop::(get_platform()), + CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT => cl_prop::(0), + CL_DEVICE_PREFERRED_INTEROP_USER_SYNC => cl_prop::(true), + CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT => cl_prop::(0), + CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT => cl_prop::(0), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR => cl_prop::(1), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE => cl_prop::(0), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT => cl_prop::(1), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF => cl_prop::(0), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT => cl_prop::(1), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG => cl_prop::(1), + CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT => cl_prop::(1), + CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => cl_prop::(1), + // TODO + CL_DEVICE_PRINTF_BUFFER_SIZE => cl_prop::(0), + // TODO + CL_DEVICE_PROFILING_TIMER_RESOLUTION => cl_prop::(0), + CL_DEVICE_OPENCL_C_FEATURES => cl_prop::>(Vec::new()), + CL_DEVICE_OPENCL_C_VERSION => { + cl_prop::(format!("OpenCL C {} ", dev.clc_version.api_str())) + } + CL_DEVICE_OPENCL_C_ALL_VERSIONS => cl_prop::<&Vec>(&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::(0), + CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE => cl_prop::(0), + CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES => cl_prop::(0), + CL_DEVICE_QUEUE_ON_HOST_PROPERTIES => { + cl_prop::(CL_QUEUE_PROFILING_ENABLE.into()) + } + CL_DEVICE_REFERENCE_COUNT => cl_prop::(1), + CL_DEVICE_SINGLE_FP_CONFIG => cl_prop::( + (CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN) as cl_device_fp_config, + ), + CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS => cl_prop::(false), + CL_DEVICE_SVM_CAPABILITIES => cl_prop::(0), + CL_DEVICE_TYPE => cl_prop::(dev.device_type()), + CL_DEVICE_VENDOR => cl_prop(dev.screen().device_vendor()), + CL_DEVICE_VENDOR_ID => cl_prop::(dev.vendor_id()), + CL_DEVICE_VERSION => cl_prop::(format!("OpenCL {}", dev.cl_version.api_str())), + CL_DRIVER_VERSION => cl_prop("0.1"), + CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT => cl_prop::(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> = 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> { + INIT.call_once(load_devices); + unsafe { &DEVICES } +} + +pub fn get_devs_for_type(device_type: cl_device_type) -> Vec<&'static Arc> { + 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(()) +} diff --git a/src/gallium/frontends/rusticl/api/event.rs b/src/gallium/frontends/rusticl/api/event.rs new file mode 100644 index 0000000..6532682 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/event.rs @@ -0,0 +1,58 @@ +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 for cl_event { + fn query(&self, q: cl_event_info) -> CLResult> { + let event = self.get_ref()?; + Ok(match q { + CL_EVENT_COMMAND_EXECUTION_STATUS => cl_prop::(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::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::from_ptr(ptr)) + } + CL_EVENT_REFERENCE_COUNT => cl_prop::(self.refcnt()?), + CL_EVENT_COMMAND_TYPE => cl_prop::(event.cmd_type), + _ => return Err(CL_INVALID_VALUE), + }) + } +} + +pub fn create_user_event(context: cl_context) -> CLResult { + let c = context.get_arc()?; + Ok(cl_event::from_arc(Event::new_user(c))) +} + +pub fn create_and_queue( + q: Arc, + cmd_type: cl_command_type, + deps: Vec>, + 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(()) +} diff --git a/src/gallium/frontends/rusticl/api/icd.rs b/src/gallium/frontends/rusticl/api/icd.rs new file mode 100644 index 0000000..f46a607 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/icd.rs @@ -0,0 +1,1260 @@ +#![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 = Result; + +#[repr(C)] +pub struct CLObjectBase { + dispatch: &'static cl_icd_dispatch, + type_err: i32, +} + +impl Default for CLObjectBase { + fn default() -> Self { + Self::new() + } +} + +impl CLObjectBase { + 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 { + 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) + 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> { + unsafe { + let ptr = self.get_ptr()?; + Arc::increment_strong_count(ptr); + Ok(Arc::from_raw(ptr)) + } + } + + fn from_arc(arc: Arc) -> Self + where + Self: Sized, + { + Self::from_ptr(Arc::into_raw(arc)) + } + + fn get_arc_vec_from_arr(objs: *const Self, count: u32) -> CLResult>> + 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> { + unsafe { Ok(Arc::from_raw(self.get_ptr()?)) } + } + + fn refcnt(&self) -> CLResult { + 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(&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, + 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, + 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, + 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, + _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, + 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, + 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, + 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)) +} diff --git a/src/gallium/frontends/rusticl/api/kernel.rs b/src/gallium/frontends/rusticl/api/kernel.rs new file mode 100644 index 0000000..6575474 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/kernel.rs @@ -0,0 +1,24 @@ +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 { + 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. +} diff --git a/src/gallium/frontends/rusticl/api/memory.rs b/src/gallium/frontends/rusticl/api/memory.rs new file mode 100644 index 0000000..6fade65 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/memory.rs @@ -0,0 +1,1343 @@ +#![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 for cl_mem { + fn query(&self, q: cl_mem_info) -> CLResult> { + 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::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::from_ptr(ptr)) + } + CL_MEM_FLAGS => cl_prop::(mem.flags), + // TODO debugging feature + CL_MEM_MAP_COUNT => cl_prop::(0), + CL_MEM_HOST_PTR => cl_prop::<*mut c_void>(mem.host_ptr), + CL_MEM_OFFSET => cl_prop::(mem.offset), + CL_MEM_REFERENCE_COUNT => cl_prop::(self.refcnt()?), + CL_MEM_SIZE => cl_prop::(mem.size), + CL_MEM_TYPE => cl_prop::(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 { + 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 { + 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::().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, + 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], +) -> CLResult { + // 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_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 for cl_mem { + fn query(&self, q: cl_image_info) -> CLResult> { + let mem = self.get_ref()?; + Ok(match *q { + CL_IMAGE_ARRAY_SIZE => cl_prop::(mem.image_desc.image_array_size), + CL_IMAGE_BUFFER => cl_prop::(unsafe { mem.image_desc.anon_1.buffer }), + CL_IMAGE_DEPTH => cl_prop::(mem.image_desc.image_depth), + CL_IMAGE_ELEMENT_SIZE => cl_prop::(mem.image_elem_size.into()), + CL_IMAGE_FORMAT => cl_prop::(mem.image_format), + CL_IMAGE_HEIGHT => cl_prop::(mem.image_desc.image_height), + CL_IMAGE_NUM_MIP_LEVELS => cl_prop::(mem.image_desc.num_mip_levels), + CL_IMAGE_NUM_SAMPLES => cl_prop::(mem.image_desc.num_samples), + CL_IMAGE_ROW_PITCH => cl_prop::(mem.image_desc.image_row_pitch), + CL_IMAGE_SLICE_PITCH => cl_prop::(mem.image_desc.image_slice_pitch), + CL_IMAGE_WIDTH => cl_prop::(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 { + 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::::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 for cl_sampler { + fn query(&self, q: cl_sampler_info) -> CLResult> { + let sampler = self.get_ref()?; + Ok(match q { + CL_SAMPLER_ADDRESSING_MODE => cl_prop::(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::from_ptr(ptr)) + } + CL_SAMPLER_FILTER_MODE => cl_prop::(sampler.filter_mode), + CL_SAMPLER_NORMALIZED_COORDS => cl_prop::(sampler.normalized_coords), + CL_SAMPLER_REFERENCE_COUNT => cl_prop::(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 { + 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(()) + }), + ) +} diff --git a/src/gallium/frontends/rusticl/api/mod.rs b/src/gallium/frontends/rusticl/api/mod.rs new file mode 100644 index 0000000..bba1f16 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/mod.rs @@ -0,0 +1,11 @@ +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; diff --git a/src/gallium/frontends/rusticl/api/platform.rs b/src/gallium/frontends/rusticl/api/platform.rs new file mode 100644 index 0000000..905e14c --- /dev/null +++ b/src/gallium/frontends/rusticl/api/platform.rs @@ -0,0 +1,121 @@ +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 for cl_platform_id { + fn query(&self, q: cl_platform_info) -> CLResult> { + let p = self.get_ref()?; + Ok(match q { + CL_PLATFORM_EXTENSIONS => cl_prop("cl_khr_icd"), + CL_PLATFORM_EXTENSIONS_WITH_VERSION => { + cl_prop::>(p.extensions.to_vec()) + } + CL_PLATFORM_HOST_TIMER_RESOLUTION => cl_prop::(0), + CL_PLATFORM_ICD_SUFFIX_KHR => cl_prop("MESA"), + CL_PLATFORM_NAME => cl_prop("rusticl"), + CL_PLATFORM_NUMERIC_VERSION => cl_prop::(CLVersion::Cl3_0 as u32), + CL_PLATFORM_PROFILE => cl_prop("FULL_PROFILE"), + CL_PLATFORM_VENDOR => cl_prop("Mesa/X.org"), + // OpenCL + 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 = 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); +} diff --git a/src/gallium/frontends/rusticl/api/program.rs b/src/gallium/frontends/rusticl/api/program.rs new file mode 100644 index 0000000..3f18f28 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/program.rs @@ -0,0 +1,287 @@ +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 for cl_program { + fn query(&self, q: cl_program_info) -> CLResult> { + 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::from_ptr(ptr)) + } + CL_PROGRAM_DEVICES => { + cl_prop::<&Vec>( + &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::(prog.devs.len() as cl_uint), + CL_PROGRAM_NUM_KERNELS => cl_prop::(prog.kernels().len()), + CL_PROGRAM_REFERENCE_COUNT => cl_prop::(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 for cl_program { + fn query(&self, d: cl_device_id, q: cl_program_build_info) -> CLResult> { + let prog = self.get_ref()?; + let dev = d.get_arc()?; + Ok(match q { + CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE => cl_prop::(0), + CL_PROGRAM_BUILD_LOG => cl_prop::(prog.log(&dev)), + CL_PROGRAM_BUILD_OPTIONS => cl_prop::(prog.options(&dev)), + CL_PROGRAM_BUILD_STATUS => cl_prop::(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], +) -> CLResult>> { + 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, + 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 { + 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, + 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, + 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, + 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. +} diff --git a/src/gallium/frontends/rusticl/api/queue.rs b/src/gallium/frontends/rusticl/api/queue.rs new file mode 100644 index 0000000..f44b325 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/queue.rs @@ -0,0 +1,77 @@ +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 for cl_command_queue { + fn query(&self, q: cl_command_queue_info) -> CLResult> { + 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::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::from_ptr(ptr)) + } + CL_QUEUE_PROPERTIES => cl_prop::(queue.props), + CL_QUEUE_REFERENCE_COUNT => cl_prop::(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_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(()) +} diff --git a/src/gallium/frontends/rusticl/api/types.rs b/src/gallium/frontends/rusticl/api/types.rs new file mode 100644 index 0000000..5b0391c --- /dev/null +++ b/src/gallium/frontends/rusticl/api/types.rs @@ -0,0 +1,147 @@ +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 { + vals: [T; 3], +} + +impl CLVec { + /// # 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 { + pub fn is_in_bound(base: Self, offset: Self, pitch: [usize; 3], size: usize) -> bool { + (base + offset - [1, 1, 1]) * pitch < size + } +} + +impl Default for CLVec { + fn default() -> Self { + Self { + vals: [T::default(); 3], + } + } +} + +// provides a ton of functions +impl std::ops::Deref for CLVec { + type Target = [T; 3]; + + fn deref(&self) -> &Self::Target { + &self.vals + } +} + +impl> std::ops::Add for CLVec { + type Output = Self; + + fn add(self, other: Self) -> Self { + self + other.vals + } +} + +impl> std::ops::Add<[T; 3]> for CLVec { + 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> std::ops::Sub<[T; 3]> for CLVec { + 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 std::ops::Mul for CLVec +where + T: Copy + std::ops::Mul + std::ops::Add, +{ + type Output = T; + + fn mul(self, other: Self) -> T { + self * other.vals + } +} + +impl std::ops::Mul<[T; 3]> for CLVec +where + T: Copy + std::ops::Mul + std::ops::Add, +{ + type Output = T; + + fn mul(self, other: [T; 3]) -> T { + self[0] * other[0] + self[1] * other[1] + self[2] * other[2] + } +} diff --git a/src/gallium/frontends/rusticl/api/util.rs b/src/gallium/frontends/rusticl/api/util.rs new file mode 100644 index 0000000..eb9f284 --- /dev/null +++ b/src/gallium/frontends/rusticl/api/util.rs @@ -0,0 +1,341 @@ +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 { + fn query(&self, q: I) -> CLResult>; + + 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 { + fn query(&self, o: O, q: I) -> CLResult>; + + 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; +} + +macro_rules! cl_prop_for_type { + ($ty: ty) => { + impl CLProp for $ty { + fn cl_vec(&self) -> Vec { + self.to_ne_bytes().to_vec() + } + } + }; +} + +macro_rules! cl_prop_for_struct { + ($ty: ty) => { + impl CLProp for $ty { + fn cl_vec(&self) -> Vec { + unsafe { slice::from_raw_parts((self as *const Self).cast(), size_of::()) } + .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 { + cl_prop::(if *self { CL_TRUE } else { CL_FALSE }) + } +} + +impl CLProp for String { + fn cl_vec(&self) -> Vec { + let mut c = self.clone(); + c.push('\0'); + c.into_bytes() + } +} + +impl CLProp for &str { + fn cl_vec(&self) -> Vec { + CString::new(*self) + .unwrap_or_default() + .into_bytes_with_nul() + } +} + +impl CLProp for &CStr { + fn cl_vec(&self) -> Vec { + self.to_bytes_with_nul().to_vec() + } +} + +impl CLProp for Vec +where + T: CLProp, +{ + fn cl_vec(&self) -> Vec { + let mut res: Vec = Vec::new(); + for i in self { + res.append(&mut i.cl_vec()) + } + res + } +} + +impl CLProp for &Vec +where + T: CLProp, +{ + fn cl_vec(&self) -> Vec { + let mut res: Vec = Vec::new(); + for i in *self { + res.append(&mut i.cl_vec()) + } + res + } +} + +impl CLProp for *const T { + fn cl_vec(&self) -> Vec { + (*self as usize).cl_vec() + } +} + +impl CLProp for *mut T { + fn cl_vec(&self) -> Vec { + (*self as usize).cl_vec() + } +} + +pub fn cl_prop(v: T) -> Vec { + 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>(val: T) -> Option { + 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, + num_events_in_wait_list: cl_uint, + event_wait_list: *const cl_event, +) -> CLResult>> { + // 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(cb: &Option, 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(ptr: *const T, alignment: usize) -> bool { + ptr as usize & (alignment - 1) == 0 +} + +pub fn bit_check + PartialEq + Default, B: Into>(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, + src_offset: usize, + dst_origin: &CLVec, + dst_offset: usize, + region: &CLVec, + 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 +} diff --git a/src/gallium/frontends/rusticl/core/context.rs b/src/gallium/frontends/rusticl/core/context.rs new file mode 100644 index 0000000..5e8a9d2 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/context.rs @@ -0,0 +1,62 @@ +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, + pub devs: Vec>, + pub properties: Vec, +} + +impl_cl_type_trait!(cl_context, Context, CL_INVALID_CONTEXT); + +impl Context { + pub fn new(devs: Vec>, properties: Vec) -> Arc { + Arc::new(Self { + base: CLObjectBase::new(), + devs: devs, + properties: properties, + }) + } + + pub fn create_buffer(&self, size: usize) -> CLResult, 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, 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) + } +} diff --git a/src/gallium/frontends/rusticl/core/device.rs b/src/gallium/frontends/rusticl/core/device.rs new file mode 100644 index 0000000..8ac22ce --- /dev/null +++ b/src/gallium/frontends/rusticl/core/device.rs @@ -0,0 +1,489 @@ +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, + screen: Arc, + pub cl_version: CLVersion, + pub clc_version: CLVersion, + pub clc_versions: Vec, + pub custom: bool, + pub embedded: bool, + pub extension_string: String, + pub extensions: Vec, + pub formats: HashMap>, + helper_ctx: Mutex>, +} + +impl_cl_type_trait!(cl_device_id, Device, CL_INVALID_DEVICE); + +impl Device { + fn new(screen: Arc) -> Option> { + 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 { + 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 = 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> { + 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 { + let v: Vec = 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::::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::::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::::compute_param( + self.screen.as_ref(), + pipe_compute_cap::PIPE_COMPUTE_CAP_MAX_INPUT_SIZE, + ) as usize + } + + pub fn screen(&self) -> &Arc { + &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> { + 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, + } + } +} diff --git a/src/gallium/frontends/rusticl/core/event.rs b/src/gallium/frontends/rusticl/core/event.rs new file mode 100644 index 0000000..441972d --- /dev/null +++ b/src/gallium/frontends/rusticl/core/event.rs @@ -0,0 +1,110 @@ +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, &Arc) -> CLResult<()>>; + +#[repr(C)] +pub struct Event { + pub base: CLObjectBase, + pub context: Arc, + pub queue: Option>, + pub cmd_type: cl_command_type, + pub deps: Vec>, + // use AtomicI32 instead of cl_int so we can change it without a &mut reference + status: AtomicI32, + work: Option, +} + +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, + cmd_type: cl_command_type, + deps: Vec>, + work: EventSig, + ) -> Arc { + 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) -> Arc { + 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>> { + 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) -> 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 diff --git a/src/gallium/frontends/rusticl/core/format.rs b/src/gallium/frontends/rusticl/core/format.rs new file mode 100644 index 0000000..4f3e6a7 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/format.rs @@ -0,0 +1,181 @@ +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, + ), +]; diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs new file mode 100644 index 0000000..93fd9f0 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/kernel.rs @@ -0,0 +1,24 @@ +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, +} + +impl_cl_type_trait!(cl_kernel, Kernel, CL_INVALID_KERNEL); + +impl Kernel { + pub fn new() -> Arc { + Arc::new(Self { + base: CLObjectBase::new(), + }) + } +} diff --git a/src/gallium/frontends/rusticl/core/memory.rs b/src/gallium/frontends/rusticl/core/memory.rs new file mode 100644 index 0000000..56fd286 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/memory.rs @@ -0,0 +1,398 @@ +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, + pub context: Arc, + pub parent: Option>, + 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>>, + res: Option, PipeResource>>, + maps: Mutex>, +} + +impl_cl_type_trait!(cl_mem, Mem, CL_INVALID_MEM_OBJECT); + +fn sw_copy( + src: *const c_void, + dst: *mut c_void, + region: &CLVec, + src_origin: &CLVec, + src_row_pitch: usize, + src_slice_pitch: usize, + dst_origin: &CLVec, + 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, + flags: cl_mem_flags, + size: usize, + host_ptr: *mut c_void, + ) -> CLResult> { + 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, + flags: cl_mem_flags, + offset: usize, + size: usize, + ) -> Arc { + 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, + 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 { + 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, PipeResource> { + self.parent + .as_ref() + .map_or(self, |p| p.as_ref()) + .res + .as_ref() + .unwrap() + } + + pub fn write_from_user( + &self, + q: &Arc, + 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, + ctx: &Arc, + region: &CLVec, + src_origin: &CLVec, + src_row_pitch: usize, + src_slice_pitch: usize, + dst_origin: &CLVec, + 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, + ctx: &Arc, + region: &CLVec, + src_origin: &CLVec, + src_row_pitch: usize, + src_slice_pitch: usize, + dst_origin: &CLVec, + 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, + ctx: &Arc, + region: &CLVec, + src_origin: &CLVec, + src_row_pitch: usize, + src_slice_pitch: usize, + dst_origin: &CLVec, + 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, 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, 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, + pub context: Arc, + 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, + normalized_coords: bool, + addressing_mode: cl_addressing_mode, + filter_mode: cl_filter_mode, + ) -> Arc { + Arc::new(Self { + base: CLObjectBase::new(), + context: context, + normalized_coords: normalized_coords, + addressing_mode: addressing_mode, + filter_mode: filter_mode, + }) + } +} diff --git a/src/gallium/frontends/rusticl/core/mod.rs b/src/gallium/frontends/rusticl/core/mod.rs new file mode 100644 index 0000000..e1658b6 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/mod.rs @@ -0,0 +1,10 @@ +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; diff --git a/src/gallium/frontends/rusticl/core/program.rs b/src/gallium/frontends/rusticl/core/program.rs new file mode 100644 index 0000000..799d237 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/program.rs @@ -0,0 +1,187 @@ +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, + pub context: Arc, + pub devs: Vec>, + pub src: CString, + build: Mutex, +} + +impl_cl_type_trait!(cl_program, Program, CL_INVALID_PROGRAM); + +struct ProgramBuild { + builds: HashMap, ProgramDevBuild>, + kernels: Vec, +} + +struct ProgramDevBuild { + spirv: Option, + status: cl_build_status, + options: String, + log: String, +} + +fn prepare_options(options: &str) -> Vec { + 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, devs: &[Arc], src: CString) -> Arc { + 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 { + self.build.lock().unwrap() + } + + fn dev_build_info<'a>( + l: &'a mut MutexGuard, + dev: &Arc, + ) -> &'a mut ProgramDevBuild { + l.builds.get_mut(dev).unwrap() + } + + pub fn status(&self, dev: &Arc) -> cl_build_status { + Self::dev_build_info(&mut self.build_info(), dev).status + } + + pub fn log(&self, dev: &Arc) -> String { + Self::dev_build_info(&mut self.build_info(), dev) + .log + .clone() + } + + pub fn options(&self, dev: &Arc) -> String { + Self::dev_build_info(&mut self.build_info(), dev) + .options + .clone() + } + + pub fn kernels(&self) -> Vec { + self.build_info().kernels.clone() + } + + pub fn compile( + &self, + dev: &Arc, + 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, + devs: &[Arc], + progs: &[Arc], + ) -> Arc { + let devs: Vec> = 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(), + }), + }) + } +} diff --git a/src/gallium/frontends/rusticl/core/queue.rs b/src/gallium/frontends/rusticl/core/queue.rs new file mode 100644 index 0000000..1751962 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/queue.rs @@ -0,0 +1,96 @@ +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, + pub context: Arc, + pub device: Arc, + pub props: cl_command_queue_properties, + pending: Mutex>>, + _thrd: Option>, + chan_in: mpsc::Sender>>, + chan_out: mpsc::Receiver, +} + +impl_cl_type_trait!(cl_command_queue, Queue, CL_INVALID_COMMAND_QUEUE); + +impl Queue { + pub fn new( + context: Arc, + device: Arc, + props: cl_command_queue_properties, + ) -> CLResult> { + // 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::>>(); + let (tx_t, rx_q) = mpsc::channel::(); + 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) { + 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); + } +} diff --git a/src/gallium/frontends/rusticl/core/util.rs b/src/gallium/frontends/rusticl/core/util.rs new file mode 100644 index 0000000..bbc48dd --- /dev/null +++ b/src/gallium/frontends/rusticl/core/util.rs @@ -0,0 +1,17 @@ +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, + } +} diff --git a/src/gallium/frontends/rusticl/core/version.rs b/src/gallium/frontends/rusticl/core/version.rs new file mode 100644 index 0000000..1c33756 --- /dev/null +++ b/src/gallium/frontends/rusticl/core/version.rs @@ -0,0 +1,80 @@ +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 for CLVersion { + type Error = cl_int; + + fn try_from(value: u32) -> Result { + 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), + }) + } +} diff --git a/src/gallium/frontends/rusticl/lib.rs b/src/gallium/frontends/rusticl/lib.rs new file mode 100644 index 0000000..56f592a --- /dev/null +++ b/src/gallium/frontends/rusticl/lib.rs @@ -0,0 +1,5 @@ +#![allow(clippy::collapsible_if)] +#![allow(clippy::needless_range_loop)] + +mod api; +pub mod core; diff --git a/src/gallium/frontends/rusticl/mesa/compiler/clc/mod.rs b/src/gallium/frontends/rusticl/mesa/compiler/clc/mod.rs new file mode 100644 index 0000000..e4919e7 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/compiler/clc/mod.rs @@ -0,0 +1 @@ +pub mod spirv; diff --git a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs new file mode 100644 index 0000000..ee36e55 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs @@ -0,0 +1,148 @@ +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, +} + +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).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, 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 = Vec::new(); + let logger = clc_logger { + priv_: &mut msgs as *mut Vec 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, 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 = Vec::new(); + let logger = clc_logger { + priv_: &mut msgs as *mut Vec 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 { + 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); + } + } + } +} diff --git a/src/gallium/frontends/rusticl/mesa/compiler/mod.rs b/src/gallium/frontends/rusticl/mesa/compiler/mod.rs new file mode 100644 index 0000000..7dffcb7 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/compiler/mod.rs @@ -0,0 +1 @@ +pub mod clc; diff --git a/src/gallium/frontends/rusticl/mesa/lib.rs b/src/gallium/frontends/rusticl/mesa/lib.rs new file mode 100644 index 0000000..f156310 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/lib.rs @@ -0,0 +1,2 @@ +pub mod compiler; +pub mod pipe; diff --git a/src/gallium/frontends/rusticl/mesa/pipe/context.rs b/src/gallium/frontends/rusticl/mesa/pipe/context.rs new file mode 100644 index 0000000..6f64c3a --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/context.rs @@ -0,0 +1,116 @@ +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, +} + +unsafe impl Send for PipeContext {} +unsafe impl Sync for PipeContext {} + +impl PipeContext { + pub(super) fn new(context: *mut pipe_context) -> Option> { + 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() +} diff --git a/src/gallium/frontends/rusticl/mesa/pipe/device.rs b/src/gallium/frontends/rusticl/mesa/pipe/device.rs new file mode 100644 index 0000000..0c1754d --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/device.rs @@ -0,0 +1,55 @@ +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 { + if ldev.is_null() { + return None; + } + Some(Self { ldev }) + } + + fn load_screen(self) -> Option> { + 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 { + 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> { + load_devs() + .into_iter() + .filter_map(PipeLoaderDevice::load_screen) + .collect() +} diff --git a/src/gallium/frontends/rusticl/mesa/pipe/mod.rs b/src/gallium/frontends/rusticl/mesa/pipe/mod.rs new file mode 100644 index 0000000..e9b8d48 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/mod.rs @@ -0,0 +1,5 @@ +pub mod context; +pub mod device; +pub mod resource; +pub mod screen; +pub mod transfer; diff --git a/src/gallium/frontends/rusticl/mesa/pipe/resource.rs b/src/gallium/frontends/rusticl/mesa/pipe/resource.rs new file mode 100644 index 0000000..febda57 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/resource.rs @@ -0,0 +1,29 @@ +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 { + 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()) } + } +} diff --git a/src/gallium/frontends/rusticl/mesa/pipe/screen.rs b/src/gallium/frontends/rusticl/mesa/pipe/screen.rs new file mode 100644 index 0000000..3da6c1e --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/screen.rs @@ -0,0 +1,175 @@ +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 { + 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> for PipeScreen { + fn compute_param(&self, cap: pipe_compute_cap) -> Vec { + let size = self.compute_param_wrapped(cap, ptr::null_mut()); + let elems = (size / 8) as usize; + + let mut res: Vec = Vec::new(); + let mut d: Vec = 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> { + if screen.is_null() || !has_required_cbs(screen) { + return None; + } + + Some(Arc::new(Self { ldev, screen })) + } + + pub fn create_context(self: &Arc) -> Option> { + 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 { + 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 { + 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() +} diff --git a/src/gallium/frontends/rusticl/mesa/pipe/transfer.rs b/src/gallium/frontends/rusticl/mesa/pipe/transfer.rs new file mode 100644 index 0000000..258a165 --- /dev/null +++ b/src/gallium/frontends/rusticl/mesa/pipe/transfer.rs @@ -0,0 +1,66 @@ +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); + } +} diff --git a/src/gallium/frontends/rusticl/meson.build b/src/gallium/frontends/rusticl/meson.build new file mode 100644 index 0000000..1745700 --- /dev/null +++ b/src/gallium/frontends/rusticl/meson.build @@ -0,0 +1,250 @@ +# 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, + ], +) diff --git a/src/gallium/frontends/rusticl/rusticl_mesa_bindings.h b/src/gallium/frontends/rusticl/rusticl_mesa_bindings.h new file mode 100644 index 0000000..d8c0d62 --- /dev/null +++ b/src/gallium/frontends/rusticl/rusticl_mesa_bindings.h @@ -0,0 +1,9 @@ +#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" diff --git a/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.c b/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.c new file mode 100644 index 0000000..80333f5 --- /dev/null +++ b/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.c @@ -0,0 +1,7 @@ +#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); +} diff --git a/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.h b/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.h new file mode 100644 index 0000000..36f5f20 --- /dev/null +++ b/src/gallium/frontends/rusticl/rusticl_mesa_inline_bindings_wrapper.h @@ -0,0 +1,5 @@ +#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); diff --git a/src/gallium/frontends/rusticl/rusticl_opencl_bindings.h b/src/gallium/frontends/rusticl/rusticl_opencl_bindings.h new file mode 100644 index 0000000..27e0498 --- /dev/null +++ b/src/gallium/frontends/rusticl/rusticl_opencl_bindings.h @@ -0,0 +1,13 @@ +#include + +#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 diff --git a/src/gallium/frontends/rusticl/util/lib.rs b/src/gallium/frontends/rusticl/util/lib.rs new file mode 100644 index 0000000..fe9cc89 --- /dev/null +++ b/src/gallium/frontends/rusticl/util/lib.rs @@ -0,0 +1,3 @@ +pub mod properties; +pub mod ptr; +pub mod string; diff --git a/src/gallium/frontends/rusticl/util/properties.rs b/src/gallium/frontends/rusticl/util/properties.rs new file mode 100644 index 0000000..aad44e5 --- /dev/null +++ b/src/gallium/frontends/rusticl/util/properties.rs @@ -0,0 +1,54 @@ +pub struct Properties { + pub props: Vec<(T, T)>, +} + +impl Properties { + #[allow(clippy::not_unsafe_ptr_arg_deref)] + pub fn from_ptr_raw(mut p: *const T) -> Vec { + let mut res: Vec = 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 { + let mut res = Self::default(); + + if !p.is_null() { + let mut k: Vec = Vec::new(); + let mut v: Vec = 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 Default for Properties { + fn default() -> Self { + Self { props: Vec::new() } + } +} diff --git a/src/gallium/frontends/rusticl/util/ptr.rs b/src/gallium/frontends/rusticl/util/ptr.rs new file mode 100644 index 0000000..9167be7f --- /dev/null +++ b/src/gallium/frontends/rusticl/util/ptr.rs @@ -0,0 +1,26 @@ +use std::ptr; + +pub trait CheckedPtr { + /// # 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 CheckedPtr 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; + } + } + } +} diff --git a/src/gallium/frontends/rusticl/util/string.rs b/src/gallium/frontends/rusticl/util/string.rs new file mode 100644 index 0000000..eea8e4f --- /dev/null +++ b/src/gallium/frontends/rusticl/util/string.rs @@ -0,0 +1,13 @@ +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("")) +} diff --git a/src/gallium/meson.build b/src/gallium/meson.build index fdb43f8..9268e40 100644 --- a/src/gallium/meson.build +++ b/src/gallium/meson.build @@ -182,6 +182,10 @@ if with_gallium_opencl 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') diff --git a/src/gallium/targets/rusticl/meson.build b/src/gallium/targets/rusticl/meson.build new file mode 100644 index 0000000..0c21cf9 --- /dev/null +++ b/src/gallium/targets/rusticl/meson.build @@ -0,0 +1,62 @@ +# 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'), +) diff --git a/src/gallium/targets/rusticl/rusticl.icd.in b/src/gallium/targets/rusticl/rusticl.icd.in new file mode 100644 index 0000000..4bf89f1 --- /dev/null +++ b/src/gallium/targets/rusticl/rusticl.icd.in @@ -0,0 +1 @@ +libRusticlOpenCL.so.@OPENCL_VERSION@ diff --git a/src/gallium/targets/rusticl/target.c b/src/gallium/targets/rusticl/target.c new file mode 100644 index 0000000..308e23b --- /dev/null +++ b/src/gallium/targets/rusticl/target.c @@ -0,0 +1,2 @@ +#include "target-helpers/drm_helper.h" +#include "target-helpers/sw_helper.h" -- 2.7.4