rusticl/kernel: basic implementation

still not able to run kernels, but most of the boilerplate code is there now

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15439>

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,
+    clRetainKernel: Some(cl_retain_kernel),
-    clReleaseKernel: None,
+    clReleaseKernel: Some(cl_release_kernel),
-    clSetKernelArg: None,
+    clSetKernelArg: Some(cl_set_kernel_arg),
-    clGetKernelInfo: None,
+    clGetKernelInfo: Some(cl_get_kernel_info),
-    clGetKernelWorkGroupInfo: None,
+    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,

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<cl_kernel_info> for cl_kernel {
+    fn query(&self, q: cl_kernel_info) -> CLResult<Vec<u8>> {
+        let kernel = self.get_ref()?;
+        Ok(match q {
+            CL_KERNEL_CONTEXT => {
+                let ptr = Arc::as_ptr(&kernel.prog.context);
+                cl_prop::<cl_context>(cl_context::from_ptr(ptr))
+            }
+            CL_KERNEL_FUNCTION_NAME => cl_prop::<&str>(&kernel.name),
+            CL_KERNEL_NUM_ARGS => cl_prop::<cl_uint>(kernel.args.len() as cl_uint),
+            CL_KERNEL_PROGRAM => {
+                let ptr = Arc::as_ptr(&kernel.prog);
+                cl_prop::<cl_program>(cl_program::from_ptr(ptr))
+            }
+            CL_KERNEL_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
+            // CL_INVALID_VALUE if param_name is not one of the supported values
+            _ => return Err(CL_INVALID_VALUE),
+        })
+    }
+}
+
+impl CLInfoObj<cl_kernel_arg_info, cl_uint> for cl_kernel {
+    fn query(&self, idx: cl_uint, q: cl_kernel_arg_info) -> CLResult<Vec<u8>> {
+        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::<cl_kernel_arg_access_qualifier>(kernel.access_qualifier(idx))
+            }
+            CL_KERNEL_ARG_ADDRESS_QUALIFIER => {
+                cl_prop::<cl_kernel_arg_address_qualifier>(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::<cl_kernel_arg_type_qualifier>(kernel.type_qualifier(idx))
+            }
+            // CL_INVALID_VALUE if param_name is not one of the supported values
+            _ => return Err(CL_INVALID_VALUE),
+        })
+    }
+}
+
+impl CLInfoObj<cl_kernel_work_group_info, cl_device_id> for cl_kernel {
+    fn query(&self, dev: cl_device_id, q: cl_kernel_work_group_info) -> CLResult<Vec<u8>> {
+        let _kernel = self.get_ref()?;
+        let _dev = dev.get_ref()?;
+        Ok(match *q {
+            CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => cl_prop::<usize>(1),
+            // TODO
+            CL_KERNEL_WORK_GROUP_SIZE => cl_prop::<usize>(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<cl_kernel> {
-    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");
+    // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
-    Err(CL_OUT_OF_HOST_MEMORY)
+    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_NAME if kernel_name is not found in program.
+    if !p.kernels().contains(&name) {
-    //• 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.
+        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)
 }

src/gallium/frontends/rusticl/api/program.rs

@@ -40,6 +40,7 @@ impl CLInfo<cl_program_info> for cl_program {
                         .collect(),
                 )
             }
+            CL_PROGRAM_KERNEL_NAMES => cl_prop::<String>(prog.kernels().join(";")),
             CL_PROGRAM_NUM_DEVICES => cl_prop::<cl_uint>(prog.devs.len() as cl_uint),
             CL_PROGRAM_NUM_KERNELS => cl_prop::<usize>(prog.kernels().len()),
             CL_PROGRAM_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),

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<CL_INVALID_KERNEL>,
+    pub prog: Arc<Program>,
+    pub name: String,
+    pub args: Vec<spirv::SPIRVKernelArg>,
 }
 
 impl_cl_type_trait!(cl_kernel, Kernel, CL_INVALID_KERNEL);
 
 impl Kernel {
-    pub fn new() -> Arc<Kernel> {
+    pub fn new(name: String, prog: Arc<Program>, args: Vec<spirv::SPIRVKernelArg>) -> Arc<Kernel> {
         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
+    }
 }

src/gallium/frontends/rusticl/core/program.rs

@@ -107,6 +107,14 @@ impl Program {
             .clone()
     }
 
+    pub fn args(&self, dev: &Arc<Device>, kernel: &str) -> Vec<spirv::SPIRVKernelArg> {
+        Self::dev_build_info(&mut self.build_info(), dev)
+            .spirv
+            .as_ref()
+            .unwrap()
+            .args(kernel)
+    }
+
     pub fn kernels(&self) -> Vec<String> {
         self.build_info().kernels.clone()
     }

src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs

@@ -17,6 +17,15 @@ pub struct SPIRVBin {
     info: Option<clc_parsed_spirv>,
 }
 
+#[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<String> {
-        unsafe {
+        self.kernel_infos()
-            match self.info {
+            .iter()
-                None => Vec::new(),
+            .map(|i| i.name)
-                Some(info) => slice::from_raw_parts(info.kernels, info.num_kernels as usize)
+            .map(c_string_to_string)
-                    .iter()
+            .collect()
-                    .map(|i| i.name)
+    }
-                    .map(c_string_to_string)
+
-                    .collect(),
+    pub fn args(&self, name: &str) -> Vec<SPIRVKernelArg> {
-            }
+        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(),
         }
     }
 }

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',
   ],