From 1b00d4f22e3d4d96a1174c684b16b4b128d74b3f Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sun, 24 Apr 2022 13:21:32 +0200 Subject: [PATCH] rusticl/kernel: implement CL_KERNEL_ATTRIBUTES Signed-off-by: Karol Herbst Acked-by: Alyssa Rosenzweig Part-of: --- src/gallium/frontends/rusticl/api/kernel.rs | 1 + src/gallium/frontends/rusticl/core/kernel.rs | 13 +++++- src/gallium/frontends/rusticl/core/program.rs | 18 +++++++++ .../rusticl/mesa/compiler/clc/spirv.rs | 40 +++++++++++++++++++ 4 files changed, 70 insertions(+), 2 deletions(-) diff --git a/src/gallium/frontends/rusticl/api/kernel.rs b/src/gallium/frontends/rusticl/api/kernel.rs index 7258601f288..3b0bf0033e9 100644 --- a/src/gallium/frontends/rusticl/api/kernel.rs +++ b/src/gallium/frontends/rusticl/api/kernel.rs @@ -20,6 +20,7 @@ impl CLInfo for cl_kernel { fn query(&self, q: cl_kernel_info, _: &[u8]) -> CLResult> { 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::from_ptr(ptr)) diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs index 14d85e95624..f67c9d516eb 100644 --- a/src/gallium/frontends/rusticl/core/kernel.rs +++ b/src/gallium/frontends/rusticl/core/kernel.rs @@ -241,6 +241,7 @@ pub struct Kernel { pub args: Vec, pub values: Vec>>, pub work_group_size: [usize; 3], + pub attributes_string: String, internal_args: Vec, nirs: HashMap, NirShader>, } @@ -584,10 +585,12 @@ fn convert_spirv_to_nir( HashMap, NirShader>, Vec, Vec, + 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, args: Vec) -> Arc { - 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(), } diff --git a/src/gallium/frontends/rusticl/core/program.rs b/src/gallium/frontends/rusticl/core/program.rs index 23e7f65f3cb..b97a183b3b6 100644 --- a/src/gallium/frontends/rusticl/core/program.rs +++ b/src/gallium/frontends/rusticl/core/program.rs @@ -462,6 +462,24 @@ impl Program { .collect() } + pub fn attribute_str(&self, kernel: &str, d: &Arc) -> 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) -> NirShader { let mut lock = self.build_info(); let info = Self::dev_build_info(&mut lock, d); diff --git a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs b/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs index 01415c6d15b..a9bfca38c86 100644 --- a/src/gallium/frontends/rusticl/mesa/compiler/clc/spirv.rs +++ b/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 { + 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) { None => Vec::new(),