rusticl/program: move attribute_str to the spirv module

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30351>

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

@@ -17,6 +17,7 @@ use mesa_rust_gen::*;
 use mesa_rust_util::math::*;
 use mesa_rust_util::serialize::*;
 use rusticl_opencl_gen::*;
+use spirv::SpirvKernelInfo;
 
 use std::cmp;
 use std::collections::HashMap;
@@ -886,7 +887,7 @@ pub(super) fn convert_spirv_to_nir(
         (nir, args, internal_args)
     };
 
-    let attributes_string = build.attribute_str(name, dev);
+    let attributes_string = build.spirv_info(name, dev).unwrap().attribute_str();
     let wgs = nir.workgroup_size();
     let kernel_info = KernelInfo {
         args: args,

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

@@ -98,21 +98,8 @@ pub struct ProgramBuild {
 }
 
 impl ProgramBuild {
-    pub fn attribute_str(&self, kernel: &str, d: &Device) -> String {
-        let info = self.dev_build(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 spirv_info(&self, kernel: &str, d: &Device) -> Option<&clc_kernel_info> {
+        self.dev_build(d).spirv.as_ref()?.kernel_info(kernel)
     }
 
     fn args(&self, dev: &Device, kernel: &str) -> Vec<spirv::SPIRVKernelArg> {

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

@@ -229,7 +229,7 @@ impl SPIRVBin {
         }
     }
 
-    fn kernel_info(&self, name: &str) -> Option<&clc_kernel_info> {
+    pub fn kernel_info(&self, name: &str) -> Option<&clc_kernel_info> {
         self.kernel_infos()
             .iter()
             .find(|i| c_string_to_string(i.name) == name)
@@ -243,46 +243,6 @@ 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) {
             Some(info) if info.num_args > 0 => {
@@ -545,3 +505,59 @@ impl CLCSpecConstantType for clc_spec_constant_type {
         }
     }
 }
+
+pub trait SpirvKernelInfo {
+    fn vec_type_hint(&self) -> Option<String>;
+    fn local_size(&self) -> Option<String>;
+    fn local_size_hint(&self) -> Option<String>;
+
+    fn attribute_str(&self) -> String {
+        let attributes_strings = [
+            self.vec_type_hint(),
+            self.local_size(),
+            self.local_size_hint(),
+        ];
+
+        let attributes_strings: Vec<_> = attributes_strings.into_iter().flatten().collect();
+        attributes_strings.join(",")
+    }
+}
+
+impl SpirvKernelInfo for clc_kernel_info {
+    fn vec_type_hint(&self) -> Option<String> {
+        if ![1, 2, 3, 4, 8, 16].contains(&self.vec_hint_size) {
+            return None;
+        }
+        let cltype = match self.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",
+        };
+
+        Some(format!("vec_type_hint({}{})", cltype, self.vec_hint_size))
+    }
+
+    fn local_size(&self) -> Option<String> {
+        if self.local_size == [0; 3] {
+            return None;
+        }
+        Some(format!(
+            "reqd_work_group_size({},{},{})",
+            self.local_size[0], self.local_size[1], self.local_size[2]
+        ))
+    }
+
+    fn local_size_hint(&self) -> Option<String> {
+        if self.local_size_hint == [0; 3] {
+            return None;
+        }
+        Some(format!(
+            "work_group_size_hint({},{},{})",
+            self.local_size_hint[0], self.local_size_hint[1], self.local_size_hint[2]
+        ))
+    }
+}