From 56eb83115538a05f656503f8d2a71a837490af6c Mon Sep 17 00:00:00 2001
From: Rhys Perry <pendingchaos02@gmail.com>
Date: Wed, 7 Jun 2023 16:37:51 +0100
Subject: [PATCH] aco: run nir_lower_int64 after nir_opt_uniform_atomics

nir_opt_uniform_atomics can create 64-bit ALU instructions which need to
be lowered.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Qiang Yu <yuq825@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23502>
---
 src/amd/compiler/aco_instruction_selection_setup.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp
index 4d4cf5b..2c8ca8c 100644
--- a/src/amd/compiler/aco_instruction_selection_setup.cpp
+++ b/src/amd/compiler/aco_instruction_selection_setup.cpp
@@ -308,7 +308,8 @@ init_context(isel_context* ctx, nir_shader* shader)
    ctx->ub_config.max_workgroup_size[2] = 2048;
 
    nir_divergence_analysis(shader);
-   nir_opt_uniform_atomics(shader);
+   if (nir_opt_uniform_atomics(shader) && nir_lower_int64(shader))
+      nir_divergence_analysis(shader);
 
    apply_nuw_to_offsets(ctx, impl);
 
-- 
2.7.4