From dc896c31f91f61889a1dca50c520aacab994dd7f Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Wed, 24 Jul 2024 17:42:08 +0200 Subject: [PATCH] rusticl/program: move attribute_str to the spirv module Part-of: --- src/gallium/frontends/rusticl/core/kernel.rs | 3 +- src/gallium/frontends/rusticl/core/program.rs | 17 +--- .../rusticl/mesa/compiler/clc/spirv.rs | 98 +++++++++++-------- 3 files changed, 61 insertions(+), 57 deletions(-) diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs index ef298a893db..57bbaf18f90 100644 --- a/src/gallium/frontends/rusticl/core/kernel.rs +++ b/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, diff --git a/src/gallium/frontends/rusticl/core/program.rs b/src/gallium/frontends/rusticl/core/program.rs index 134d6742d2b..1653da8fa68 100644 --- a/src/gallium/frontends/rusticl/core/program.rs +++ b/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 { diff --git a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs index a58dcd7ae80..0fbd39c97b5 100644 --- a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs +++ b/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 { - 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 { - 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 { - 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 { 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; + fn local_size(&self) -> Option; + fn local_size_hint(&self) -> Option; + + 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 { + 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 { + 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 { + 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] + )) + } +}