OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nir/lower_system_values change num_workgroups to uint32_t

Browse Source 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24882>
This commit is contained in:
Rhys Perry
2023-09-08 15:52:21 +01:00
committed by Marge Bot
parent b3b58b91f4
commit 81d17246ec
2 changed files with 8 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/compiler/nir/nir.h
View File

@@ -5375,7 +5375,7 @@ typedef struct nir_lower_compute_system_values_options {
* and compute it quickly. Fall back to slow computation if not.
*/
bool shortcut_1d_workgroup_id : 1;
uint16_t num_workgroups[3]; /* Compile-time-known dispatch sizes, or 0 if unknown. */
uint32_t num_workgroups[3]; /* Compile-time-known dispatch sizes, or 0 if unknown. */
} nir_lower_compute_system_values_options;
bool nir_lower_compute_system_values(nir_shader *shader,

12
src/compiler/nir/nir_lower_system_values.c
View File

@@ -384,7 +384,7 @@ id_to_index_no_umod_slow(nir_builder *b, nir_def *index,
static nir_def *
lower_id_to_index_no_umod(nir_builder *b, nir_def *index,
nir_def *size, unsigned bit_size,
const uint16_t *size_imm,
const uint32_t *size_imm,
bool shortcut_1d)
{
nir_def *size_x, *size_y;
@@ -465,7 +465,7 @@ lower_compute_system_value_filter(const nir_instr *instr, const void *_state)
}
static nir_def *
try_lower_id_to_index_1d(nir_builder *b, nir_def *index, const uint16_t *size)
try_lower_id_to_index_1d(nir_builder *b, nir_def *index, const uint32_t *size)
{
/* size_x = 1, size_y = 1, therefore Z = local index */
if (size[0] == 1 && size[1] == 1)
@@ -512,8 +512,10 @@ lower_compute_system_value_instr(nir_builder *b,
* this way we don't leave behind extra ALU instrs.
*/
nir_def *val = try_lower_id_to_index_1d(b, local_index,
b->shader->info.workgroup_size);
uint32_t wg_size[3] = {b->shader->info.workgroup_size[0],
b->shader->info.workgroup_size[1],
b->shader->info.workgroup_size[2]};
nir_def *val = try_lower_id_to_index_1d(b, local_index, wg_size);
if (val)
return val;
}
@@ -734,7 +736,7 @@ lower_compute_system_value_instr(nir_builder *b,
if (!options)
return NULL;
const uint16_t *num_wgs_imm = options->num_workgroups;
const uint32_t *num_wgs_imm = options->num_workgroups;
/* Exit early when none of the num workgroups components are known at
* compile time.