rusticl/kernel: implement CL_KERNEL_ATTRIBUTES

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/kernel.rs

@@ -20,6 +20,7 @@ impl CLInfo<cl_kernel_info> for cl_kernel {
     fn query(&self, q: cl_kernel_info, _: &[u8]) -> CLResult<Vec<u8>> {
         let kernel = self.get_ref()?;
         Ok(match q {
+            CL_KERNEL_ATTRIBUTES => cl_prop::<&str>(&kernel.attributes_string),
             CL_KERNEL_CONTEXT => {
                 let ptr = Arc::as_ptr(&kernel.prog.context);
                 cl_prop::<cl_context>(cl_context::from_ptr(ptr))

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

@@ -241,6 +241,7 @@ pub struct Kernel {
     pub args: Vec<KernelArg>,
     pub values: Vec<RefCell<Option<KernelArgValue>>>,
     pub work_group_size: [usize; 3],
+    pub attributes_string: String,
     internal_args: Vec<InternalKernelArg>,
     nirs: HashMap<Arc<Device>, NirShader>,
 }
@@ -584,10 +585,12 @@ fn convert_spirv_to_nir(
     HashMap<Arc<Device>, NirShader>,
     Vec<KernelArg>,
     Vec<InternalKernelArg>,
+    String,
 ) {
     let mut nirs = HashMap::new();
     let mut args_set = HashSet::new();
     let mut internal_args_set = HashSet::new();
+    let mut attributes_string_set = HashSet::new();
 
     // TODO: we could run this in parallel?
     for d in p.devs_with_build() {
@@ -639,15 +642,18 @@ fn convert_spirv_to_nir(
         args_set.insert(args);
         internal_args_set.insert(internal_args);
         nirs.insert(d.clone(), nir);
+        attributes_string_set.insert(p.attribute_str(name, d));
     }
 
     // we want the same (internal) args for every compiled kernel, for now
     assert!(args_set.len() == 1);
     assert!(internal_args_set.len() == 1);
+    assert!(attributes_string_set.len() == 1);
     let args = args_set.into_iter().next().unwrap();
     let internal_args = internal_args_set.into_iter().next().unwrap();
+    let attributes_string = attributes_string_set.into_iter().next().unwrap();
 
-    (nirs, args, internal_args)
+    (nirs, args, internal_args, attributes_string)
 }
 
 fn extract<'a, const S: usize>(buf: &'a mut &[u8]) -> &'a [u8; S] {
@@ -698,7 +704,8 @@ fn optimize_local_size(d: &Device, grid: &mut [u32; 3], block: &mut [u32; 3]) {
 
 impl Kernel {
     pub fn new(name: String, prog: Arc<Program>, args: Vec<spirv::SPIRVKernelArg>) -> Arc<Kernel> {
-        let (mut nirs, args, internal_args) = convert_spirv_to_nir(&prog, &name, args);
+        let (mut nirs, args, internal_args, attributes_string) =
+            convert_spirv_to_nir(&prog, &name, args);
 
         let nir = nirs.values_mut().next().unwrap();
         let wgs = nir.workgroup_size();
@@ -713,6 +720,7 @@ impl Kernel {
             name: name,
             args: args,
             work_group_size: work_group_size,
+            attributes_string: attributes_string,
             values: values,
             internal_args: internal_args,
             // caller has to verify all kernels have the same sig
@@ -1010,6 +1018,7 @@ impl Clone for Kernel {
             args: self.args.clone(),
             values: self.values.clone(),
             work_group_size: self.work_group_size,
+            attributes_string: self.attributes_string.clone(),
             internal_args: self.internal_args.clone(),
             nirs: self.nirs.clone(),
         }

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

@@ -462,6 +462,24 @@ impl Program {
             .collect()
     }
 
+    pub fn attribute_str(&self, kernel: &str, d: &Arc<Device>) -> String {
+        let mut lock = self.build_info();
+        let info = Self::dev_build_info(&mut lock, d);
+
+        let attributes_strings = [
+            info.spirv.as_ref().unwrap().vec_type_hint(kernel),
+            info.spirv.as_ref().unwrap().local_size(kernel),
+            info.spirv.as_ref().unwrap().local_size_hint(kernel),
+        ];
+
+        let attributes_strings: Vec<_> = attributes_strings
+            .iter()
+            .flatten()
+            .map(String::as_str)
+            .collect();
+        attributes_strings.join(",")
+    }
+
     pub fn to_nir(&self, kernel: &str, d: &Arc<Device>) -> NirShader {
         let mut lock = self.build_info();
         let info = Self::dev_build_info(&mut lock, d);

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

@@ -190,6 +190,46 @@ impl SPIRVBin {
             .collect()
     }
 
+    pub fn vec_type_hint(&self, name: &str) -> Option<String> {
+        self.kernel_info(name)
+            .filter(|info| [1, 2, 3, 4, 8, 16].contains(&info.vec_hint_size))
+            .map(|info| {
+                let cltype = match info.vec_hint_type {
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_CHAR => "uchar",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_SHORT => "ushort",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_INT => "uint",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_LONG => "ulong",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_HALF => "half",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_FLOAT => "float",
+                    clc_vec_hint_type::CLC_VEC_HINT_TYPE_DOUBLE => "double",
+                };
+
+                format!("vec_type_hint({}{})", cltype, info.vec_hint_size)
+            })
+    }
+
+    pub fn local_size(&self, name: &str) -> Option<String> {
+        self.kernel_info(name)
+            .filter(|info| info.local_size != [0; 3])
+            .map(|info| {
+                format!(
+                    "reqd_work_group_size({},{},{})",
+                    info.local_size[0], info.local_size[1], info.local_size[2]
+                )
+            })
+    }
+
+    pub fn local_size_hint(&self, name: &str) -> Option<String> {
+        self.kernel_info(name)
+            .filter(|info| info.local_size_hint != [0; 3])
+            .map(|info| {
+                format!(
+                    "work_group_size_hint({},{},{})",
+                    info.local_size_hint[0], info.local_size_hint[1], info.local_size_hint[2]
+                )
+            })
+    }
+
     pub fn args(&self, name: &str) -> Vec<SPIRVKernelArg> {
         match self.kernel_info(name) {
             None => Vec::new(),