From 2273b6c46a5eee5cd0131d98c85a4efa8fe52898 Mon Sep 17 00:00:00 2001 From: Dave Airlie Date: Tue, 15 Jul 2025 06:26:52 +1000 Subject: [PATCH] nak: add divergent attribute and wrapper for nir_load_sysval_nv This wraps the sysval load in a builder where we can add proper divergence for ctaid later. Part-of: --- src/compiler/nir/nir_divergence_analysis.c | 2 +- src/compiler/nir/nir_intrinsics.py | 2 +- src/nouveau/compiler/nak_nir.c | 60 +++++++++++---------- src/nouveau/compiler/nak_nir_lower_vtg_io.c | 13 +++-- src/nouveau/compiler/nak_private.h | 3 ++ 5 files changed, 43 insertions(+), 37 deletions(-) diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c index cecf84c23be..e2d7a6f8648 100644 --- a/src/compiler/nir/nir_divergence_analysis.c +++ b/src/compiler/nir/nir_divergence_analysis.c @@ -377,6 +377,7 @@ visit_intrinsic(nir_intrinsic_instr *instr, struct divergence_state *state) break; case nir_intrinsic_decl_reg: + case nir_intrinsic_load_sysval_nv: is_divergent = nir_intrinsic_divergent(instr); break; @@ -917,7 +918,6 @@ visit_intrinsic(nir_intrinsic_instr *instr, struct divergence_state *state) case nir_intrinsic_load_ray_hit_kind: case nir_intrinsic_load_ray_flags: case nir_intrinsic_load_cull_mask: - case nir_intrinsic_load_sysval_nv: case nir_intrinsic_emit_vertex_nv: case nir_intrinsic_end_primitive_nv: case nir_intrinsic_report_ray_intersection: diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py index a640cf30ed1..dd49204045a 100644 --- a/src/compiler/nir/nir_intrinsics.py +++ b/src/compiler/nir/nir_intrinsics.py @@ -2511,7 +2511,7 @@ intrinsic("ldcx_nv", dest_comp=0, src_comp=[1, 1], indices=[ACCESS, ALIGN_MUL, ALIGN_OFFSET], flags=[CAN_ELIMINATE, CAN_REORDER]) intrinsic("load_sysval_nv", dest_comp=1, src_comp=[], bit_sizes=[32, 64], - indices=[ACCESS, BASE], flags=[CAN_ELIMINATE]) + indices=[ACCESS, BASE, DIVERGENT], flags=[CAN_ELIMINATE]) intrinsic("isberd_nv", dest_comp=1, src_comp=[1], bit_sizes=[32], flags=[CAN_ELIMINATE, CAN_REORDER]) intrinsic("vild_nv", dest_comp=1, src_comp=[1], bit_sizes=[32], diff --git a/src/nouveau/compiler/nak_nir.c b/src/nouveau/compiler/nak_nir.c index 49f15eed098..750981b70cb 100644 --- a/src/nouveau/compiler/nak_nir.c +++ b/src/nouveau/compiler/nak_nir.c @@ -20,6 +20,15 @@ #define OPT_V(nir, pass, ...) NIR_PASS(_, nir, pass, ##__VA_ARGS__) +nir_def * +nak_nir_load_sysval(nir_builder *b, enum nak_sv idx, + enum gl_access_qualifier access) +{ + return nir_load_sysval_nv(b, 32, .base = idx, + .access = access, + .divergent = true); +} + bool nak_nir_workgroup_has_one_subgroup(const nir_shader *nir) { @@ -458,15 +467,13 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, } case nir_intrinsic_load_patch_vertices_in: { - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_PRIM_TYPE, - .access = ACCESS_CAN_REORDER); + val = nak_nir_load_sysval(b, NAK_SV_PRIM_TYPE, ACCESS_CAN_REORDER); val = nir_extract_u8(b, val, nir_imm_int(b, 1)); break; } case nir_intrinsic_load_frag_shading_rate: { - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VARIABLE_RATE, - .access = ACCESS_CAN_REORDER); + val = nak_nir_load_sysval(b, NAK_SV_VARIABLE_RATE, ACCESS_CAN_REORDER); /* X is in bits 8..16 and Y is in bits 16..24. However, we actually * want the log2 of X and Y and, since we only support 1, 2, and 4, a @@ -492,8 +499,7 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, const gl_system_value sysval = nir_system_value_from_intrinsic(intrin->intrinsic); const uint32_t idx = nak_sysval_sysval_idx(sysval); - val = nir_load_sysval_nv(b, 32, .base = idx, - .access = ACCESS_CAN_REORDER); + val = nak_nir_load_sysval(b, idx, ACCESS_CAN_REORDER); /* Pad with 0 because all invocations above 31 are off */ if (intrin->def.bit_size == 64) { @@ -515,21 +521,19 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, nir_def *comps[3]; assert(intrin->def.num_components <= 3); for (unsigned c = 0; c < intrin->def.num_components; c++) { - comps[c] = nir_load_sysval_nv(b, 32, .base = idx + c, - .access = ACCESS_CAN_REORDER); + comps[c] = nak_nir_load_sysval(b, idx + c, ACCESS_CAN_REORDER); } val = nir_vec(b, comps, intrin->def.num_components); break; } case nir_intrinsic_load_local_invocation_id: { - nir_def *x = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_X, - .access = ACCESS_CAN_REORDER); - nir_def *y = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_Y, - .access = ACCESS_CAN_REORDER); - nir_def *z = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_Z, - .access = ACCESS_CAN_REORDER); - + nir_def *x = nak_nir_load_sysval(b, NAK_SV_TID_X, + ACCESS_CAN_REORDER); + nir_def *y = nak_nir_load_sysval(b, NAK_SV_TID_Y, + ACCESS_CAN_REORDER); + nir_def *z = nak_nir_load_sysval(b, NAK_SV_TID_Z, + ACCESS_CAN_REORDER); if (b->shader->info.derivative_group == DERIVATIVE_GROUP_QUADS) { nir_def *x_lo = nir_iand_imm(b, x, 0x1); nir_def *y_lo = nir_ushr_imm(b, nir_iand_imm(b, x, 0x2), 1); @@ -558,12 +562,12 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, val = nir_imm_int(b, 0); } else { assert(!b->shader->info.workgroup_size_variable); - nir_def *tid_x = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_X, - .access = ACCESS_CAN_REORDER); - nir_def *tid_y = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_Y, - .access = ACCESS_CAN_REORDER); - nir_def *tid_z = nir_load_sysval_nv(b, 32, .base = NAK_SV_TID_Z, - .access = ACCESS_CAN_REORDER); + nir_def *tid_x = nak_nir_load_sysval(b, NAK_SV_TID_X, + ACCESS_CAN_REORDER); + nir_def *tid_y = nak_nir_load_sysval(b, NAK_SV_TID_Y, + ACCESS_CAN_REORDER); + nir_def *tid_z = nak_nir_load_sysval(b, NAK_SV_TID_Z, + ACCESS_CAN_REORDER); const uint16_t *wg_size = b->shader->info.workgroup_size; nir_def *tid = @@ -577,7 +581,7 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, case nir_intrinsic_is_helper_invocation: { /* Unlike load_helper_invocation, this one isn't re-orderable */ - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_THREAD_KILL); + val = nak_nir_load_sysval(b, NAK_SV_THREAD_KILL, 0); break; } @@ -605,7 +609,7 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, nir_variable *clock = nir_local_variable_create(b->impl, glsl_uvec2_type(), NULL); - nir_def *clock_hi = nir_load_sysval_nv(b, 32, .base = NAK_SV_CLOCK_HI); + nir_def *clock_hi = nak_nir_load_sysval(b, NAK_SV_CLOCK_HI, 0); nir_ssa_bar_nv(b, clock_hi); nir_store_var(b, clock, nir_vec2(b, nir_imm_int(b, 0), clock_hi), 0x3); @@ -614,10 +618,10 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, { nir_def *last_clock = nir_load_var(b, clock); - nir_def *clock_lo = nir_load_sysval_nv(b, 32, .base = NAK_SV_CLOCK_LO); + nir_def *clock_lo = nak_nir_load_sysval(b, NAK_SV_CLOCK_LO, 0); nir_ssa_bar_nv(b, clock_lo); - clock_hi = nir_load_sysval_nv(b, 32, .base = NAK_SV_CLOCK + 1); + clock_hi = nak_nir_load_sysval(b, NAK_SV_CLOCK + 1, 0); nir_ssa_bar_nv(b, clock_hi); nir_store_var(b, clock, nir_vec2(b, clock_lo, clock_hi), 0x3); @@ -637,17 +641,17 @@ nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin, break; case nir_intrinsic_load_sm_count_nv: - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VIRTCFG); + val = nak_nir_load_sysval(b, NAK_SV_VIRTCFG, 0); val = nir_ubitfield_extract_imm(b, val, 20, 9); break; case nir_intrinsic_load_warp_id_nv: - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VIRTID); + val = nak_nir_load_sysval(b, NAK_SV_VIRTID, 0); val = nir_ubitfield_extract_imm(b, val, 8, 7); break; case nir_intrinsic_load_sm_id_nv: - val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VIRTID); + val = nak_nir_load_sysval(b, NAK_SV_VIRTID, 0); val = nir_ubitfield_extract_imm(b, val, 20, 9); break; diff --git a/src/nouveau/compiler/nak_nir_lower_vtg_io.c b/src/nouveau/compiler/nak_nir_lower_vtg_io.c index 4c528869c36..6ccd78260ed 100644 --- a/src/nouveau/compiler/nak_nir_lower_vtg_io.c +++ b/src/nouveau/compiler/nak_nir_lower_vtg_io.c @@ -40,10 +40,10 @@ tess_ctrl_output_vtx(nir_builder *b, nir_def *vtx) * range of lanes. We have to compute the lane index of the requested * invocation from the invocation index. */ - nir_def *lane = nir_load_sysval_nv(b, 32, .base = NAK_SV_LANE_ID, - .access = ACCESS_CAN_REORDER); - nir_def *invoc = nir_load_sysval_nv(b, 32, .base = NAK_SV_INVOCATION_ID, - .access = ACCESS_CAN_REORDER); + nir_def *lane = nak_nir_load_sysval(b, NAK_SV_LANE_ID, + ACCESS_CAN_REORDER); + nir_def *invoc = nak_nir_load_sysval(b, NAK_SV_INVOCATION_ID, + ACCESS_CAN_REORDER); return nir_iadd(b, lane, nir_iadd(b, vtx, nir_ineg(b, invoc))); } @@ -139,9 +139,8 @@ lower_vtg_io_intrin(nir_builder *b, if (vtx != NULL && !is_output) { if (nak->sm >= 50) { - nir_def *info = nir_load_sysval_nv(b, 32, - .base = NAK_SV_INVOCATION_INFO, - .access = ACCESS_CAN_REORDER); + nir_def *info = nak_nir_load_sysval(b, NAK_SV_INVOCATION_INFO, + ACCESS_CAN_REORDER); nir_def *lo = nir_extract_u8_imm(b, info, 0); nir_def *hi = nir_extract_u8_imm(b, info, 2); nir_def *idx = nir_iadd(b, nir_imul(b, lo, hi), vtx); diff --git a/src/nouveau/compiler/nak_private.h b/src/nouveau/compiler/nak_private.h index 1087823000c..969f97b2bc8 100644 --- a/src/nouveau/compiler/nak_private.h +++ b/src/nouveau/compiler/nak_private.h @@ -132,6 +132,9 @@ enum ENUM_PACKED nak_sv { bool nak_nir_workgroup_has_one_subgroup(const nir_shader *nir); +nir_def *nak_nir_load_sysval(nir_builder *b, enum nak_sv idx, + enum gl_access_qualifier access); + struct nak_xfb_info nak_xfb_from_nir(const struct nak_compiler *nak, const struct nir_xfb_info *nir_xfb);