AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nak: add divergent attribute and wrapper for nir_load_sysval_nv

Browse Source 
This wraps the sysval load in a builder where we can add
proper divergence for ctaid later.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36105>
This commit is contained in:
Dave Airlie
2025-07-15 06:26:52 +10:00
committed by Marge Bot
parent d12bc87dda
commit 2273b6c46a
5 changed files with 43 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/compiler/nir/nir_divergence_analysis.c
+1 -1
View File
@@ -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:
src/compiler/nir/nir_intrinsics.py
+1 -1
View File
@@ -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],
src/nouveau/compiler/nak_nir.c
+32 -28
View File
@@ -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;
src/nouveau/compiler/nak_nir_lower_vtg_io.c
+6 -7
View File
@@ -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);
src/nouveau/compiler/nak_private.h
+3
View File
@@ -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);