From b9e47cf5fd8c2279df5efa2c1b4b7eab09622b55 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Thu, 3 Aug 2023 14:09:37 +0200 Subject: [PATCH] rusticl/kernel: move things around in lower_and_optimize_nir Signed-off-by: Karol Herbst Reviewed-by: @LingMan <18294-LingMan@users.noreply.gitlab.freedesktop.org> Part-of: --- src/gallium/frontends/rusticl/core/kernel.rs | 53 +++++++++++++--------------- 1 file changed, 24 insertions(+), 29 deletions(-) diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs index a6015c1..9c56742 100644 --- a/src/gallium/frontends/rusticl/core/kernel.rs +++ b/src/gallium/frontends/rusticl/core/kernel.rs @@ -392,6 +392,30 @@ fn lower_and_optimize_nir( args: &[spirv::SPIRVKernelArg], lib_clc: &NirShader, ) -> (Vec, Vec) { + let address_bits_base_type; + let address_bits_ptr_type; + let global_address_format; + let shared_address_format; + + if dev.address_bits() == 64 { + address_bits_base_type = glsl_base_type::GLSL_TYPE_UINT64; + address_bits_ptr_type = unsafe { glsl_uint64_t_type() }; + global_address_format = nir_address_format::nir_address_format_64bit_global; + shared_address_format = nir_address_format::nir_address_format_32bit_offset_as_64bit; + } else { + address_bits_base_type = glsl_base_type::GLSL_TYPE_UINT; + address_bits_ptr_type = unsafe { glsl_uint_type() }; + global_address_format = nir_address_format::nir_address_format_32bit_global; + shared_address_format = nir_address_format::nir_address_format_32bit_offset; + } + + let mut lower_state = rusticl_lower_state::default(); + let nir_options = unsafe { + &*dev + .screen + .nir_shader_compiler_options(pipe_shader_type::PIPE_SHADER_COMPUTE) + }; + nir_pass!(nir, nir_scale_fdiv); nir.set_workgroup_size_variable_if_zero(); nir.structurize(); @@ -431,26 +455,7 @@ fn lower_and_optimize_nir( opt_nir(nir, dev); let mut args = KernelArg::from_spirv_nir(args, nir); - - let address_bits_base_type; - let address_bits_ptr_type; - - if dev.address_bits() == 64 { - address_bits_base_type = glsl_base_type::GLSL_TYPE_UINT64; - address_bits_ptr_type = unsafe { glsl_uint64_t_type() }; - } else { - address_bits_base_type = glsl_base_type::GLSL_TYPE_UINT; - address_bits_ptr_type = unsafe { glsl_uint_type() }; - }; - let mut internal_args = Vec::new(); - let nir_options = unsafe { - &*dev - .screen - .nir_shader_compiler_options(pipe_shader_type::PIPE_SHADER_COMPUTE) - }; - let mut lower_state = rusticl_lower_state::default(); - nir_pass!(nir, nir_lower_memcpy); let dv_opts = nir_remove_dead_variables_options { @@ -615,16 +620,6 @@ fn lower_and_optimize_nir( Some(glsl_get_cl_type_size_align), ); - let global_address_format; - let shared_address_format; - if dev.address_bits() == 32 { - global_address_format = nir_address_format::nir_address_format_32bit_global; - shared_address_format = nir_address_format::nir_address_format_32bit_offset; - } else { - global_address_format = nir_address_format::nir_address_format_64bit_global; - shared_address_format = nir_address_format::nir_address_format_32bit_offset_as_64bit; - } - nir_pass!( nir, nir_lower_explicit_io, -- 2.7.4