fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-21 06:48:09 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

rusticl/kernel: move things around in lower_and_optimize_nir

Browse source 
Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: @LingMan <18294-LingMan@users.noreply.gitlab.freedesktop.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24470>
This commit is contained in:
Karol Herbst 2023-08-03 14:09:37 +02:00 • committed by Marge Bot
parent f06f59ea3b
commit b9e47cf5fd
1 changed files with 24 additions and 29 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

53
src/gallium/frontends/rusticl/core/kernel.rs
View file

@ -392,6 +392,30 @@ fn lower_and_optimize_nir(
args: &[spirv::SPIRVKernelArg],
lib_clc: &NirShader,
) -> (Vec<KernelArg>, Vec<InternalKernelArg>) {
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,