From e8de580998fb48bda58e35150614996b8573ba08 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Thu, 10 Mar 2022 19:32:35 +0100 Subject: [PATCH] rusticl/kernel: basic implementation still not able to run kernels, but most of the boilerplate code is there now Signed-off-by: Karol Herbst Acked-by: Alyssa Rosenzweig Part-of: --- src/gallium/frontends/rusticl/api/icd.rs | 101 ++++++++++++++-- src/gallium/frontends/rusticl/api/kernel.rs | 134 ++++++++++++++++++++- src/gallium/frontends/rusticl/api/program.rs | 1 + src/gallium/frontends/rusticl/core/kernel.rs | 74 +++++++++++- src/gallium/frontends/rusticl/core/program.rs | 8 ++ .../frontends/rusticl/mesa/compiler/clc/spirv.rs | 51 ++++++-- src/gallium/frontends/rusticl/meson.build | 5 + 7 files changed, 350 insertions(+), 24 deletions(-) diff --git a/src/gallium/frontends/rusticl/api/icd.rs b/src/gallium/frontends/rusticl/api/icd.rs index bbacf3a..e7febc9 100644 --- a/src/gallium/frontends/rusticl/api/icd.rs +++ b/src/gallium/frontends/rusticl/api/icd.rs @@ -58,14 +58,14 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch { 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, + clRetainKernel: Some(cl_retain_kernel), + clReleaseKernel: Some(cl_release_kernel), + clSetKernelArg: Some(cl_set_kernel_arg), + clGetKernelInfo: Some(cl_get_kernel_info), + clGetKernelWorkGroupInfo: Some(cl_get_kernel_work_group_info), clWaitForEvents: Some(cl_wait_for_events), clGetEventInfo: Some(cl_get_event_info), - clRetainEvent: None, + clRetainEvent: Some(cl_retain_event), clReleaseEvent: Some(cl_release_event), clGetEventProfilingInfo: Some(cl_get_event_profiling_info), clFlush: Some(cl_flush), @@ -81,7 +81,7 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch { clEnqueueMapBuffer: Some(cl_enqueue_map_buffer), clEnqueueMapImage: Some(cl_enqueue_map_image), clEnqueueUnmapMemObject: Some(cl_enqueue_unmap_mem_object), - clEnqueueNDRangeKernel: None, + clEnqueueNDRangeKernel: Some(cl_enqueue_ndrange_kernel), clEnqueueTask: None, clEnqueueNativeKernel: None, clEnqueueMarker: None, @@ -123,7 +123,7 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch { clCompileProgram: Some(cl_compile_program), clLinkProgram: Some(cl_link_program), clUnloadPlatformCompiler: Some(cl_unload_platform_compiler), - clGetKernelArgInfo: None, + clGetKernelArgInfo: Some(cl_get_kernel_arg_info), clEnqueueFillBuffer: None, clEnqueueFillImage: Some(cl_enqueue_fill_image), clEnqueueMigrateMemObjects: None, @@ -775,6 +775,55 @@ extern "C" fn cl_create_kernels_in_program( CL_OUT_OF_HOST_MEMORY } +extern "C" fn cl_retain_kernel(kernel: cl_kernel) -> cl_int { + match_err!(kernel.retain()) +} + +extern "C" fn cl_release_kernel(kernel: cl_kernel) -> cl_int { + match_err!(kernel.release()) +} + +extern "C" fn cl_set_kernel_arg( + kernel: cl_kernel, + arg_index: cl_uint, + arg_size: usize, + arg_value: *const ::std::os::raw::c_void, +) -> cl_int { + match_err!(set_kernel_arg(kernel, arg_index, arg_size, arg_value)) +} + +extern "C" fn cl_get_kernel_info( + kernel: cl_kernel, + param_name: cl_kernel_info, + param_value_size: usize, + param_value: *mut ::std::os::raw::c_void, + param_value_size_ret: *mut usize, +) -> cl_int { + match_err!(kernel.get_info( + param_name, + param_value_size, + param_value, + param_value_size_ret, + )) +} + +extern "C" fn cl_get_kernel_work_group_info( + kernel: cl_kernel, + device: cl_device_id, + param_name: cl_kernel_work_group_info, + param_value_size: usize, + param_value: *mut ::std::os::raw::c_void, + param_value_size_ret: *mut usize, +) -> cl_int { + match_err!(kernel.get_info_obj( + device, + param_name, + param_value_size, + param_value, + param_value_size_ret, + )) +} + 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 @@ -795,6 +844,10 @@ extern "C" fn cl_get_event_info( )) } +extern "C" fn cl_retain_event(event: cl_event) -> cl_int { + match_err!(event.retain()) +} + extern "C" fn cl_release_event(event: cl_event) -> cl_int { match_err!(event.release()) } @@ -1025,6 +1078,21 @@ extern "C" fn cl_enqueue_unmap_mem_object( )) } +extern "C" fn cl_enqueue_ndrange_kernel( + _command_queue: cl_command_queue, + _kernel: cl_kernel, + _work_dim: cl_uint, + _global_work_offset: *const usize, + _global_work_size: *const usize, + _local_work_size: *const usize, + _num_events_in_wait_list: cl_uint, + _event_wait_list: *const cl_event, + _event: *mut cl_event, +) -> cl_int { + println!("cl_enqueue_ndrange_kernel not implemented"); + CL_OUT_OF_HOST_MEMORY +} + extern "C" fn cl_get_extension_function_address( function_name: *const ::std::os::raw::c_char, ) -> *mut ::std::ffi::c_void { @@ -1249,6 +1317,23 @@ extern "C" fn cl_unload_platform_compiler(_platform: cl_platform_id) -> cl_int { CL_OUT_OF_HOST_MEMORY } +extern "C" fn cl_get_kernel_arg_info( + kernel: cl_kernel, + arg_indx: cl_uint, + param_name: cl_kernel_arg_info, + param_value_size: usize, + param_value: *mut ::std::os::raw::c_void, + param_value_size_ret: *mut usize, +) -> cl_int { + match_err!(kernel.get_info_obj( + arg_indx, + param_name, + param_value_size, + param_value, + param_value_size_ret, + )) +} + extern "C" fn cl_enqueue_fill_image( _command_queue: cl_command_queue, _image: cl_mem, diff --git a/src/gallium/frontends/rusticl/api/kernel.rs b/src/gallium/frontends/rusticl/api/kernel.rs index 6575474..0a67f8d 100644 --- a/src/gallium/frontends/rusticl/api/kernel.rs +++ b/src/gallium/frontends/rusticl/api/kernel.rs @@ -1,24 +1,146 @@ +extern crate mesa_rust_util; extern crate rusticl_opencl_gen; use crate::api::icd::*; +use crate::api::util::*; +use crate::core::kernel::*; +use self::mesa_rust_util::string::*; use self::rusticl_opencl_gen::*; +use std::collections::HashSet; +use std::sync::Arc; + +impl CLInfo for cl_kernel { + fn query(&self, q: cl_kernel_info) -> CLResult> { + let kernel = self.get_ref()?; + Ok(match q { + CL_KERNEL_CONTEXT => { + let ptr = Arc::as_ptr(&kernel.prog.context); + cl_prop::(cl_context::from_ptr(ptr)) + } + CL_KERNEL_FUNCTION_NAME => cl_prop::<&str>(&kernel.name), + CL_KERNEL_NUM_ARGS => cl_prop::(kernel.args.len() as cl_uint), + CL_KERNEL_PROGRAM => { + let ptr = Arc::as_ptr(&kernel.prog); + cl_prop::(cl_program::from_ptr(ptr)) + } + CL_KERNEL_REFERENCE_COUNT => cl_prop::(self.refcnt()?), + // CL_INVALID_VALUE if param_name is not one of the supported values + _ => return Err(CL_INVALID_VALUE), + }) + } +} + +impl CLInfoObj for cl_kernel { + fn query(&self, idx: cl_uint, q: cl_kernel_arg_info) -> CLResult> { + let kernel = self.get_ref()?; + + // CL_INVALID_ARG_INDEX if arg_index is not a valid argument index. + if idx as usize >= kernel.args.len() { + return Err(CL_INVALID_ARG_INDEX); + } + + Ok(match *q { + CL_KERNEL_ARG_ACCESS_QUALIFIER => { + cl_prop::(kernel.access_qualifier(idx)) + } + CL_KERNEL_ARG_ADDRESS_QUALIFIER => { + cl_prop::(kernel.address_qualifier(idx)) + } + CL_KERNEL_ARG_NAME => cl_prop::<&str>(kernel.arg_name(idx)), + CL_KERNEL_ARG_TYPE_NAME => cl_prop::<&str>(kernel.arg_type_name(idx)), + CL_KERNEL_ARG_TYPE_QUALIFIER => { + cl_prop::(kernel.type_qualifier(idx)) + } + // CL_INVALID_VALUE if param_name is not one of the supported values + _ => return Err(CL_INVALID_VALUE), + }) + } +} + +impl CLInfoObj for cl_kernel { + fn query(&self, dev: cl_device_id, q: cl_kernel_work_group_info) -> CLResult> { + let _kernel = self.get_ref()?; + let _dev = dev.get_ref()?; + Ok(match *q { + CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => cl_prop::(1), + // TODO + CL_KERNEL_WORK_GROUP_SIZE => cl_prop::(1), + // CL_INVALID_VALUE if param_name is not one of the supported values + _ => return Err(CL_INVALID_VALUE), + }) + } +} + pub fn create_kernel( program: cl_program, kernel_name: *const ::std::os::raw::c_char, ) -> CLResult { - let _p = program.get_ref()?; + let p = program.get_arc()?; + let name = c_string_to_string(kernel_name); // 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. + if p.kernels().is_empty() { + return Err(CL_INVALID_PROGRAM_EXECUTABLE); + } - //• 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. + // CL_INVALID_KERNEL_NAME if kernel_name is not found in program. + if !p.kernels().contains(&name) { + return Err(CL_INVALID_KERNEL_NAME); + } + + // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program. + let devs: Vec<_> = p + .devs + .iter() + .filter(|d| p.status(d) == CL_BUILD_SUCCESS as cl_build_status) + .collect(); + if devs.is_empty() { + return Err(CL_INVALID_PROGRAM_EXECUTABLE); + } + + // 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. + let kernel_args: HashSet<_> = devs.iter().map(|d| p.args(d, &name)).collect(); + if kernel_args.len() != 1 { + return Err(CL_INVALID_KERNEL_DEFINITION); + } + + Ok(cl_kernel::from_arc(Kernel::new( + name, + p, + kernel_args.into_iter().next().unwrap(), + ))) +} + +pub fn set_kernel_arg( + kernel: cl_kernel, + arg_index: cl_uint, + _arg_size: usize, + _arg_value: *const ::std::os::raw::c_void, +) -> CLResult<()> { + let k = kernel.get_arc()?; + + // CL_INVALID_ARG_INDEX if arg_index is not a valid argument index. + if arg_index as usize >= k.args.len() { + return Err(CL_INVALID_ARG_INDEX); + } + + //• CL_INVALID_ARG_VALUE if arg_value specified is not a valid value. + //• CL_INVALID_MEM_OBJECT for an argument declared to be a memory object when the specified arg_value is not a valid memory object. + //• CL_INVALID_SAMPLER for an argument declared to be of type sampler_t when the specified arg_value is not a valid sampler object. + //• CL_INVALID_DEVICE_QUEUE for an argument declared to be of type queue_t when the specified arg_value is not a valid device queue object. This error code is missing before version 2.0. + //• CL_INVALID_ARG_SIZE if arg_size does not match the size of the data type for an argument that is not a memory object or if the argument is a memory object and arg_size != sizeof(cl_mem) or if arg_size is zero and the argument is declared with the local qualifier or if the argument is a sampler and arg_size != sizeof(cl_sampler). + //• CL_MAX_SIZE_RESTRICTION_EXCEEDED if the size in bytes of the memory object (if the argument is a memory object) or arg_size (if the argument is declared with local qualifier) exceeds a language- specified maximum size restriction for this argument, such as the MaxByteOffset SPIR-V decoration. This error code is missing before version 2.2. + //• CL_INVALID_ARG_VALUE if the argument is an image declared with the read_only qualifier and arg_value refers to an image object created with cl_mem_flags of CL_MEM_WRITE_ONLY or if the image argument is declared with the write_only qualifier and arg_value refers to an image object created with cl_mem_flags of CL_MEM_READ_ONLY. + + println!("set_kernel_arg not implemented"); + Err(CL_OUT_OF_HOST_MEMORY) } diff --git a/src/gallium/frontends/rusticl/api/program.rs b/src/gallium/frontends/rusticl/api/program.rs index 046762f..82376fe 100644 --- a/src/gallium/frontends/rusticl/api/program.rs +++ b/src/gallium/frontends/rusticl/api/program.rs @@ -40,6 +40,7 @@ impl CLInfo for cl_program { .collect(), ) } + CL_PROGRAM_KERNEL_NAMES => cl_prop::(prog.kernels().join(";")), 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()?), diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs index 93fd9f0..8b48755 100644 --- a/src/gallium/frontends/rusticl/core/kernel.rs +++ b/src/gallium/frontends/rusticl/core/kernel.rs @@ -1,9 +1,13 @@ extern crate mesa_rust; +extern crate mesa_rust_gen; extern crate rusticl_opencl_gen; use crate::api::icd::*; +use crate::core::program::*; use crate::impl_cl_type_trait; +use self::mesa_rust::compiler::clc::*; +use self::mesa_rust_gen::*; use self::rusticl_opencl_gen::*; use std::sync::Arc; @@ -11,14 +15,82 @@ use std::sync::Arc; #[repr(C)] pub struct Kernel { pub base: CLObjectBase, + pub prog: Arc, + pub name: String, + pub args: Vec, } impl_cl_type_trait!(cl_kernel, Kernel, CL_INVALID_KERNEL); impl Kernel { - pub fn new() -> Arc { + pub fn new(name: String, prog: Arc, args: Vec) -> Arc { Arc::new(Self { base: CLObjectBase::new(), + prog: prog, + name: name, + args: args, }) } + + pub fn access_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_access_qualifier { + let aq = self.args[idx as usize].access_qualifier; + + if aq + == clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_READ + | clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_WRITE + { + CL_KERNEL_ARG_ACCESS_READ_WRITE + } else if aq == clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_READ { + CL_KERNEL_ARG_ACCESS_READ_ONLY + } else if aq == clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_WRITE { + CL_KERNEL_ARG_ACCESS_WRITE_ONLY + } else { + CL_KERNEL_ARG_ACCESS_NONE + } + } + + pub fn address_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_address_qualifier { + match self.args[idx as usize].address_qualifier { + clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_PRIVATE => { + CL_KERNEL_ARG_ADDRESS_PRIVATE + } + clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_CONSTANT => { + CL_KERNEL_ARG_ADDRESS_CONSTANT + } + clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_LOCAL => { + CL_KERNEL_ARG_ADDRESS_LOCAL + } + clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_GLOBAL => { + CL_KERNEL_ARG_ADDRESS_GLOBAL + } + } + } + + pub fn type_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_type_qualifier { + let tq = self.args[idx as usize].type_qualifier; + let zero = clc_kernel_arg_type_qualifier(0); + let mut res = CL_KERNEL_ARG_TYPE_NONE; + + if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_CONST != zero { + res |= CL_KERNEL_ARG_TYPE_CONST; + } + + if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_RESTRICT != zero { + res |= CL_KERNEL_ARG_TYPE_RESTRICT; + } + + if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_VOLATILE != zero { + res |= CL_KERNEL_ARG_TYPE_VOLATILE; + } + + res.into() + } + + pub fn arg_name(&self, idx: cl_uint) -> &String { + &self.args[idx as usize].name + } + + pub fn arg_type_name(&self, idx: cl_uint) -> &String { + &self.args[idx as usize].type_name + } } diff --git a/src/gallium/frontends/rusticl/core/program.rs b/src/gallium/frontends/rusticl/core/program.rs index a5826bd..3c96a0d 100644 --- a/src/gallium/frontends/rusticl/core/program.rs +++ b/src/gallium/frontends/rusticl/core/program.rs @@ -107,6 +107,14 @@ impl Program { .clone() } + pub fn args(&self, dev: &Arc, kernel: &str) -> Vec { + Self::dev_build_info(&mut self.build_info(), dev) + .spirv + .as_ref() + .unwrap() + .args(kernel) + } + pub fn kernels(&self) -> Vec { self.build_info().kernels.clone() } diff --git a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs index ee36e55..89f7655 100644 --- a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs +++ b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs @@ -17,6 +17,15 @@ pub struct SPIRVBin { info: Option, } +#[derive(PartialEq, Eq, Hash)] +pub struct SPIRVKernelArg { + pub name: String, + pub type_name: String, + pub access_qualifier: clc_kernel_arg_access_qualifier, + pub address_qualifier: clc_kernel_arg_address_qualifier, + pub type_qualifier: clc_kernel_arg_type_qualifier, +} + pub struct CLCHeader<'a> { pub name: CString, pub source: &'a CString, @@ -122,16 +131,40 @@ impl SPIRVBin { (res, msgs.join("\n")) } + fn kernel_infos(&self) -> &[clc_kernel_info] { + match self.info { + None => &[], + Some(info) => unsafe { slice::from_raw_parts(info.kernels, info.num_kernels as usize) }, + } + } + + fn kernel_info(&self, name: &str) -> Option<&clc_kernel_info> { + self.kernel_infos() + .iter() + .find(|i| c_string_to_string(i.name) == name) + } + 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(), - } + self.kernel_infos() + .iter() + .map(|i| i.name) + .map(c_string_to_string) + .collect() + } + + pub fn args(&self, name: &str) -> Vec { + match self.kernel_info(name) { + None => Vec::new(), + Some(info) => unsafe { slice::from_raw_parts(info.args, info.num_args) } + .iter() + .map(|a| SPIRVKernelArg { + name: c_string_to_string(a.name), + type_name: c_string_to_string(a.type_name), + access_qualifier: clc_kernel_arg_access_qualifier(a.access_qualifier), + address_qualifier: a.address_qualifier, + type_qualifier: clc_kernel_arg_type_qualifier(a.type_qualifier), + }) + .collect(), } } } diff --git a/src/gallium/frontends/rusticl/meson.build b/src/gallium/frontends/rusticl/meson.build index ab0f519..b55ba94 100644 --- a/src/gallium/frontends/rusticl/meson.build +++ b/src/gallium/frontends/rusticl/meson.build @@ -127,6 +127,7 @@ rusticl_opencl_bindings_rs = rust.bindgen( '--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', + '--new-type-alias-deref', 'cl_kernel_(arg|work_group)_info', ], ) @@ -177,6 +178,10 @@ rusticl_mesa_bindings_rs = rust.bindgen( '--whitelist-function', 'clc_.*', '--whitelist-function', 'pipe_.*', '--whitelist-type', 'pipe_endian', + '--whitelist-type', 'clc_kernel_arg_access_qualifier', + '--bitfield-enum', 'clc_kernel_arg_access_qualifier', + '--whitelist-type', 'clc_kernel_arg_type_qualifier', + '--bitfield-enum', 'clc_kernel_arg_type_qualifier', '--whitelist-var', 'PIPE_.*', '--bitfield-enum', 'pipe_map_flags', ], -- 2.7.4