Google Git
Sign in
fuchsia / third_party / mesa / refs/heads/sandbox/turnip / . / src / nouveau / compiler / nak_nir.c
blob: 4835de038f16607d580d38ec4b8892997d6d6cd9 [file] [log] [blame] [edit]
/*
* Copyright © 2022 Collabora, Ltd.
* SPDX-License-Identifier: MIT
*/
#include "nak_private.h"
#include "nir_builder.h"
#include "nir_control_flow.h"
#include "nir_xfb_info.h"
#include "util/u_math.h"
#define OPT(nir, pass, ...) ({ \
bool this_progress = false; \
NIR_PASS(this_progress, nir, pass, ##__VA_ARGS__); \
if (this_progress) \
progress = true; \
this_progress; \
})
#define OPT_V(nir, pass, ...) NIR_PASS_V(nir, pass, ##__VA_ARGS__)
bool
nak_nir_workgroup_has_one_subgroup(const nir_shader *nir)
{
switch (nir->info.stage) {
case MESA_SHADER_VERTEX:
case MESA_SHADER_TESS_EVAL:
case MESA_SHADER_GEOMETRY:
case MESA_SHADER_FRAGMENT:
unreachable("Shader stage does not have workgroups");
break;
case MESA_SHADER_TESS_CTRL:
/* Tessellation only ever has one subgroup per workgroup. The Vulkan
* limit on the number of tessellation invocations is 32 to allow for
* this.
*/
return true;
case MESA_SHADER_COMPUTE:
case MESA_SHADER_KERNEL: {
if (nir->info.workgroup_size_variable)
return false;
uint16_t wg_sz = nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
nir->info.workgroup_size[2];
return wg_sz <= NAK_SUBGROUP_SIZE;
}
default:
unreachable("Unknown shader stage");
}
}
static uint8_t
vectorize_filter_cb(const nir_instr *instr, const void *_data)
{
if (instr->type != nir_instr_type_alu)
return 0;
const nir_alu_instr *alu = nir_instr_as_alu(instr);
const unsigned bit_size = nir_alu_instr_is_comparison(alu)
? alu->src[0].src.ssa->bit_size
: alu->def.bit_size;
switch (alu->op) {
case nir_op_fadd:
case nir_op_fsub:
case nir_op_fabs:
case nir_op_fneg:
case nir_op_feq:
case nir_op_fge:
case nir_op_flt:
case nir_op_fneu:
case nir_op_fmul:
case nir_op_ffma:
case nir_op_fsign:
case nir_op_fsat:
case nir_op_fmax:
case nir_op_fmin:
return bit_size == 16 ? 2 : 1;
default:
return 1;
}
}
static void
optimize_nir(nir_shader *nir, const struct nak_compiler *nak, bool allow_copies)
{
bool progress;
unsigned lower_flrp =
(nir->options->lower_flrp16 ? 16 : 0) |
(nir->options->lower_flrp32 ? 32 : 0) |
(nir->options->lower_flrp64 ? 64 : 0);
do {
progress = false;
/* This pass is causing problems with types used by OpenCL :
* https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13955
*
* Running with it disabled made no difference in the resulting assembly
* code.
*/
if (nir->info.stage != MESA_SHADER_KERNEL)
OPT(nir, nir_split_array_vars, nir_var_function_temp);
OPT(nir, nir_shrink_vec_array_vars, nir_var_function_temp);
OPT(nir, nir_opt_deref);
if (OPT(nir, nir_opt_memcpy))
OPT(nir, nir_split_var_copies);
OPT(nir, nir_lower_vars_to_ssa);
if (allow_copies) {
/* Only run this pass in the first call to brw_nir_optimize. Later
* calls assume that we've lowered away any copy_deref instructions
* and we don't want to introduce any more.
*/
OPT(nir, nir_opt_find_array_copies);
}
OPT(nir, nir_opt_copy_prop_vars);
OPT(nir, nir_opt_dead_write_vars);
OPT(nir, nir_opt_combine_stores, nir_var_all);
OPT(nir, nir_lower_alu_width, vectorize_filter_cb, NULL);
OPT(nir, nir_opt_vectorize, vectorize_filter_cb, NULL);
OPT(nir, nir_lower_phis_to_scalar, false);
OPT(nir, nir_lower_frexp);
OPT(nir, nir_copy_prop);
OPT(nir, nir_opt_dce);
OPT(nir, nir_opt_cse);
OPT(nir, nir_opt_peephole_select, 0, false, false);
OPT(nir, nir_opt_intrinsics);
OPT(nir, nir_opt_idiv_const, 32);
OPT(nir, nir_opt_algebraic);
OPT(nir, nir_lower_constant_convert_alu_types);
OPT(nir, nir_opt_constant_folding);
if (lower_flrp != 0) {
if (OPT(nir, nir_lower_flrp, lower_flrp, false /* always_precise */))
OPT(nir, nir_opt_constant_folding);
/* Nothing should rematerialize any flrps */
lower_flrp = 0;
}
OPT(nir, nir_opt_dead_cf);
if (OPT(nir, nir_opt_loop)) {
/* If nir_opt_loop makes progress, then we need to clean things up
* if we want any hope of nir_opt_if or nir_opt_loop_unroll to make
* progress.
*/
OPT(nir, nir_copy_prop);
OPT(nir, nir_opt_dce);
}
OPT(nir, nir_opt_if, nir_opt_if_optimize_phi_true_false);
OPT(nir, nir_opt_conditional_discard);
if (nir->options->max_unroll_iterations != 0) {
OPT(nir, nir_opt_loop_unroll);
}
OPT(nir, nir_opt_remove_phis);
OPT(nir, nir_opt_gcm, false);
OPT(nir, nir_opt_undef);
OPT(nir, nir_lower_pack);
} while (progress);
OPT(nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
}
void
nak_optimize_nir(nir_shader *nir, const struct nak_compiler *nak)
{
optimize_nir(nir, nak, false);
}
static unsigned
lower_bit_size_cb(const nir_instr *instr, void *data)
{
const struct nak_compiler *nak = data;
switch (instr->type) {
case nir_instr_type_alu: {
nir_alu_instr *alu = nir_instr_as_alu(instr);
if (nir_op_infos[alu->op].is_conversion)
return 0;
const unsigned bit_size = nir_alu_instr_is_comparison(alu)
? alu->src[0].src.ssa->bit_size
: alu->def.bit_size;
switch (alu->op) {
case nir_op_bit_count:
case nir_op_ufind_msb:
case nir_op_ifind_msb:
case nir_op_find_lsb:
/* These are handled specially because the destination is always
* 32-bit and so the bit size of the instruction is given by the
* source.
*/
return alu->src[0].src.ssa->bit_size == 32 ? 0 : 32;
case nir_op_fabs:
case nir_op_fadd:
case nir_op_fneg:
case nir_op_feq:
case nir_op_fge:
case nir_op_flt:
case nir_op_fneu:
case nir_op_fmul:
case nir_op_ffma:
case nir_op_ffmaz:
case nir_op_fsign:
case nir_op_fsat:
case nir_op_fceil:
case nir_op_ffloor:
case nir_op_fround_even:
case nir_op_ftrunc:
if (bit_size == 16 && nak->sm >= 70)
return 0;
break;
case nir_op_fmax:
case nir_op_fmin:
if (bit_size == 16 && nak->sm >= 80)
return 0;
break;
default:
break;
}
if (bit_size >= 32)
return 0;
if (bit_size & (8 | 16))
return 32;
return 0;
}
case nir_instr_type_intrinsic: {
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_vote_ieq:
if (intrin->src[0].ssa->bit_size != 1 &&
intrin->src[0].ssa->bit_size < 32)
return 32;
return 0;
case nir_intrinsic_vote_feq:
case nir_intrinsic_read_invocation:
case nir_intrinsic_read_first_invocation:
case nir_intrinsic_shuffle:
case nir_intrinsic_shuffle_xor:
case nir_intrinsic_shuffle_up:
case nir_intrinsic_shuffle_down:
case nir_intrinsic_quad_broadcast:
case nir_intrinsic_quad_swap_horizontal:
case nir_intrinsic_quad_swap_vertical:
case nir_intrinsic_quad_swap_diagonal:
case nir_intrinsic_reduce:
case nir_intrinsic_inclusive_scan:
case nir_intrinsic_exclusive_scan:
if (intrin->src[0].ssa->bit_size < 32)
return 32;
return 0;
default:
return 0;
}
}
case nir_instr_type_phi: {
nir_phi_instr *phi = nir_instr_as_phi(instr);
if (phi->def.bit_size < 32 && phi->def.bit_size != 1)
return 32;
return 0;
}
default:
return 0;
}
}
void
nak_preprocess_nir(nir_shader *nir, const struct nak_compiler *nak)
{
UNUSED bool progress = false;
nir_validate_ssa_dominance(nir, "before nak_preprocess_nir");
if (nir->info.stage == MESA_SHADER_FRAGMENT) {
nir_lower_io_to_temporaries(nir, nir_shader_get_entrypoint(nir),
true /* outputs */, false /* inputs */);
}
const nir_lower_tex_options tex_options = {
.lower_txd_3d = true,
.lower_txd_cube_map = true,
.lower_txd_clamp = true,
.lower_txd_shadow = true,
.lower_txp = ~0,
/* TODO: More lowering */
};
OPT(nir, nir_lower_tex, &tex_options);
OPT(nir, nir_normalize_cubemap_coords);
nir_lower_image_options image_options = {
.lower_cube_size = true,
};
OPT(nir, nir_lower_image, &image_options);
OPT(nir, nir_lower_global_vars_to_local);
OPT(nir, nir_split_var_copies);
OPT(nir, nir_split_struct_vars, nir_var_function_temp);
/* Optimize but allow copies because we haven't lowered them yet */
optimize_nir(nir, nak, true /* allow_copies */);
OPT(nir, nir_lower_load_const_to_scalar);
OPT(nir, nir_lower_var_copies);
OPT(nir, nir_lower_system_values);
OPT(nir, nir_lower_compute_system_values, NULL);
if (nir->info.stage == MESA_SHADER_FRAGMENT)
OPT(nir, nir_lower_terminate_to_demote);
}
uint16_t
nak_varying_attr_addr(gl_varying_slot slot)
{
if (slot >= VARYING_SLOT_PATCH0) {
return NAK_ATTR_PATCH_START + (slot - VARYING_SLOT_PATCH0) * 0x10;
} else if (slot >= VARYING_SLOT_VAR0) {
return NAK_ATTR_GENERIC_START + (slot - VARYING_SLOT_VAR0) * 0x10;
} else {
switch (slot) {
case VARYING_SLOT_TESS_LEVEL_OUTER: return NAK_ATTR_TESS_LOD;
case VARYING_SLOT_TESS_LEVEL_INNER: return NAK_ATTR_TESS_INTERRIOR;
case VARYING_SLOT_PRIMITIVE_ID: return NAK_ATTR_PRIMITIVE_ID;
case VARYING_SLOT_LAYER: return NAK_ATTR_RT_ARRAY_INDEX;
case VARYING_SLOT_VIEWPORT: return NAK_ATTR_VIEWPORT_INDEX;
case VARYING_SLOT_PSIZ: return NAK_ATTR_POINT_SIZE;
case VARYING_SLOT_POS: return NAK_ATTR_POSITION;
case VARYING_SLOT_CLIP_DIST0: return NAK_ATTR_CLIP_CULL_DIST_0;
case VARYING_SLOT_CLIP_DIST1: return NAK_ATTR_CLIP_CULL_DIST_4;
default: unreachable("Invalid varying slot");
}
}
}
static uint16_t
nak_fs_out_addr(gl_frag_result slot, uint32_t blend_idx)
{
switch (slot) {
case FRAG_RESULT_DEPTH:
assert(blend_idx == 0);
return NAK_FS_OUT_DEPTH;
case FRAG_RESULT_STENCIL:
unreachable("EXT_shader_stencil_export not supported");
case FRAG_RESULT_COLOR:
unreachable("Vulkan alway uses explicit locations");
case FRAG_RESULT_SAMPLE_MASK:
assert(blend_idx == 0);
return NAK_FS_OUT_SAMPLE_MASK;
default:
assert(blend_idx < 2);
return NAK_FS_OUT_COLOR((slot - FRAG_RESULT_DATA0) + blend_idx);
}
}
uint16_t
nak_sysval_attr_addr(gl_system_value sysval)
{
switch (sysval) {
case SYSTEM_VALUE_PRIMITIVE_ID: return NAK_ATTR_PRIMITIVE_ID;
case SYSTEM_VALUE_FRAG_COORD: return NAK_ATTR_POSITION;
case SYSTEM_VALUE_POINT_COORD: return NAK_ATTR_POINT_SPRITE;
case SYSTEM_VALUE_TESS_COORD: return NAK_ATTR_TESS_COORD;
case SYSTEM_VALUE_INSTANCE_ID: return NAK_ATTR_INSTANCE_ID;
case SYSTEM_VALUE_VERTEX_ID: return NAK_ATTR_VERTEX_ID;
case SYSTEM_VALUE_FRONT_FACE: return NAK_ATTR_FRONT_FACE;
case SYSTEM_VALUE_LAYER_ID: return NAK_ATTR_RT_ARRAY_INDEX;
default: unreachable("Invalid system value");
}
}
static uint8_t
nak_sysval_sysval_idx(gl_system_value sysval)
{
switch (sysval) {
case SYSTEM_VALUE_SUBGROUP_INVOCATION: return NAK_SV_LANE_ID;
case SYSTEM_VALUE_VERTICES_IN: return NAK_SV_VERTEX_COUNT;
case SYSTEM_VALUE_INVOCATION_ID: return NAK_SV_INVOCATION_ID;
case SYSTEM_VALUE_HELPER_INVOCATION: return NAK_SV_THREAD_KILL;
case SYSTEM_VALUE_LOCAL_INVOCATION_ID: return NAK_SV_TID;
case SYSTEM_VALUE_WORKGROUP_ID: return NAK_SV_CTAID;
case SYSTEM_VALUE_SUBGROUP_EQ_MASK: return NAK_SV_LANEMASK_EQ;
case SYSTEM_VALUE_SUBGROUP_LT_MASK: return NAK_SV_LANEMASK_LT;
case SYSTEM_VALUE_SUBGROUP_LE_MASK: return NAK_SV_LANEMASK_LE;
case SYSTEM_VALUE_SUBGROUP_GT_MASK: return NAK_SV_LANEMASK_GT;
case SYSTEM_VALUE_SUBGROUP_GE_MASK: return NAK_SV_LANEMASK_GE;
default: unreachable("Invalid system value");
}
}
static bool
nak_nir_lower_system_value_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
void *data)
{
const struct nak_compiler *nak = data;
b->cursor = nir_before_instr(&intrin->instr);
nir_def *val;
switch (intrin->intrinsic) {
case nir_intrinsic_load_primitive_id:
case nir_intrinsic_load_instance_id:
case nir_intrinsic_load_vertex_id: {
assert(b->shader->info.stage != MESA_SHADER_VERTEX ||
b->shader->info.stage != MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL ||
b->shader->info.stage == MESA_SHADER_GEOMETRY);
const gl_system_value sysval =
nir_system_value_from_intrinsic(intrin->intrinsic);
const uint32_t addr = nak_sysval_attr_addr(sysval);
val = nir_ald_nv(b, 1, nir_imm_int(b, 0), nir_imm_int(b, 0),
.base = addr, .flags = 0,
.range_base = addr, .range = 4,
.access = ACCESS_CAN_REORDER);
break;
}
case nir_intrinsic_load_patch_vertices_in: {
val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VERTEX_COUNT,
.access = ACCESS_CAN_REORDER);
val = nir_extract_u8(b, val, nir_imm_int(b, 1));
break;
}
case nir_intrinsic_load_subgroup_eq_mask:
case nir_intrinsic_load_subgroup_lt_mask:
case nir_intrinsic_load_subgroup_le_mask:
case nir_intrinsic_load_subgroup_gt_mask:
case nir_intrinsic_load_subgroup_ge_mask: {
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);
/* Pad with 0 because all invocations above 31 are off */
if (intrin->def.bit_size == 64) {
val = nir_u2u32(b, val);
} else {
assert(intrin->def.bit_size == 32);
val = nir_pad_vector_imm_int(b, val, 0, intrin->def.num_components);
}
break;
}
case nir_intrinsic_load_subgroup_invocation:
case nir_intrinsic_load_helper_invocation:
case nir_intrinsic_load_invocation_id:
case nir_intrinsic_load_workgroup_id: {
const gl_system_value sysval =
nir_system_value_from_intrinsic(intrin->intrinsic);
const uint32_t idx = nak_sysval_sysval_idx(sysval);
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);
}
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);
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);
nir_def *x_hi = nir_ushr_imm(b, nir_iand_imm(b, x, ~0x3), 1);
nir_def *y_hi = nir_ishl_imm(b, y, 1);
x = nir_ior(b, x_lo, x_hi);
y = nir_ior(b, y_lo, y_hi);
}
val = nir_vec3(b, x, y, z);
break;
}
case nir_intrinsic_load_num_subgroups: {
assert(!b->shader->info.workgroup_size_variable);
uint16_t wg_size = b->shader->info.workgroup_size[0] *
b->shader->info.workgroup_size[1] *
b->shader->info.workgroup_size[2];
val = nir_imm_int(b, DIV_ROUND_UP(wg_size, 32));
break;
}
case nir_intrinsic_load_subgroup_id:
if (nak_nir_workgroup_has_one_subgroup(b->shader)) {
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);
const uint16_t *wg_size = b->shader->info.workgroup_size;
nir_def *tid =
nir_iadd(b, tid_x,
nir_iadd(b, nir_imul_imm(b, tid_y, wg_size[0]),
nir_imul_imm(b, tid_z, wg_size[0] * wg_size[1])));
val = nir_udiv_imm(b, tid, 32);
}
break;
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);
break;
}
case nir_intrinsic_shader_clock: {
/* The CS2R opcode can load 64 bits worth of sysval data at a time but
* it's not actually atomic. In order to get correct shader clocks, we
* need to do a loop where we do
*
* CS2R SV_CLOCK_HI
* CS2R SV_CLOCK_LO
* CS2R SV_CLOCK_HI
* CS2R SV_CLOCK_LO
* CS2R SV_CLOCK_HI
* ...
*
* The moment two high values are the same, we take the low value
* between them and that gives us our clock.
*
* In order to make sure we don't run into any weird races, we also need
* to insert a barrier after every load to ensure the one load completes
* before we kick off the next load. Otherwise, if one load happens to
* be faster than the other (they are variable latency, after all) we're
* still guaranteed that the loads happen in the order we want.
*/
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_ssa_bar_nv(b, clock_hi);
nir_store_var(b, clock, nir_vec2(b, nir_imm_int(b, 0), clock_hi), 0x3);
nir_push_loop(b);
{
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_ssa_bar_nv(b, clock_lo);
clock_hi = nir_load_sysval_nv(b, 32, .base = NAK_SV_CLOCK + 1);
nir_ssa_bar_nv(b, clock_hi);
nir_store_var(b, clock, nir_vec2(b, clock_lo, clock_hi), 0x3);
nir_break_if(b, nir_ieq(b, clock_hi, nir_channel(b, last_clock, 1)));
}
nir_pop_loop(b, NULL);
val = nir_load_var(b, clock);
if (intrin->def.bit_size == 64)
val = nir_pack_64_2x32(b, val);
break;
}
case nir_intrinsic_load_warps_per_sm_nv:
val = nir_imm_int(b, nak->warps_per_sm);
break;
case nir_intrinsic_load_sm_count_nv:
val = nir_load_sysval_nv(b, 32, .base = NAK_SV_VIRTCFG);
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 = 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 = nir_ubitfield_extract_imm(b, val, 20, 9);
break;
default:
return false;
}
if (intrin->def.bit_size == 1)
val = nir_i2b(b, val);
nir_def_rewrite_uses(&intrin->def, val);
return true;
}
static bool
nak_nir_lower_system_values(nir_shader *nir, const struct nak_compiler *nak)
{
return nir_shader_intrinsics_pass(nir, nak_nir_lower_system_value_intrin,
nir_metadata_none,
(void *)nak);
}
struct nak_xfb_info
nak_xfb_from_nir(const struct nir_xfb_info *nir_xfb)
{
if (nir_xfb == NULL)
return (struct nak_xfb_info) { };
struct nak_xfb_info nak_xfb = { };
u_foreach_bit(b, nir_xfb->buffers_written) {
nak_xfb.stride[b] = nir_xfb->buffers[b].stride;
nak_xfb.stream[b] = nir_xfb->buffer_to_stream[b];
}
memset(nak_xfb.attr_index, 0xff, sizeof(nak_xfb.attr_index)); /* = skip */
for (unsigned o = 0; o < nir_xfb->output_count; o++) {
const nir_xfb_output_info *out = &nir_xfb->outputs[o];
const uint8_t b = out->buffer;
assert(nir_xfb->buffers_written & BITFIELD_BIT(b));
const uint16_t attr_addr = nak_varying_attr_addr(out->location);
assert(attr_addr % 4 == 0);
const uint16_t attr_idx = attr_addr / 4;
assert(out->offset % 4 == 0);
uint8_t out_idx = out->offset / 4;
u_foreach_bit(c, out->component_mask)
nak_xfb.attr_index[b][out_idx++] = attr_idx + c;
nak_xfb.attr_count[b] = MAX2(nak_xfb.attr_count[b], out_idx);
}
return nak_xfb;
}
static bool
lower_fs_output_intrin(nir_builder *b, nir_intrinsic_instr *intrin, void *_data)
{
if (intrin->intrinsic != nir_intrinsic_store_output)
return false;
b->cursor = nir_before_instr(&intrin->instr);
const nir_io_semantics sem = nir_intrinsic_io_semantics(intrin);
uint16_t addr = nak_fs_out_addr(sem.location, sem.dual_source_blend_index) +
nir_src_as_uint(intrin->src[1]) * 16 +
nir_intrinsic_component(intrin) * 4;
nir_def *data = intrin->src[0].ssa;
/* The fs_out_nv intrinsic is always scalar */
u_foreach_bit(c, nir_intrinsic_write_mask(intrin)) {
if (nir_scalar_is_undef(nir_scalar_resolved(data, c)))
continue;
nir_fs_out_nv(b, nir_channel(b, data, c), .base = addr + c * 4);
}
nir_instr_remove(&intrin->instr);
return true;
}
static bool
nak_nir_lower_fs_outputs(nir_shader *nir)
{
if (nir->info.outputs_written == 0)
return false;
bool progress = nir_shader_intrinsics_pass(nir, lower_fs_output_intrin,
nir_metadata_control_flow,
NULL);
if (progress) {
/* We need a copy_fs_outputs_nv intrinsic so NAK knows where to place
* the final copy. This needs to be in the last block, after all
* store_output intrinsics.
*/
nir_function_impl *impl = nir_shader_get_entrypoint(nir);
nir_builder b = nir_builder_at(nir_after_impl(impl));
nir_copy_fs_outputs_nv(&b);
}
return progress;
}
static bool
nak_nir_remove_barrier_intrin(nir_builder *b, nir_intrinsic_instr *barrier,
UNUSED void *_data)
{
if (barrier->intrinsic != nir_intrinsic_barrier)
return false;
mesa_scope exec_scope = nir_intrinsic_execution_scope(barrier);
assert(exec_scope <= SCOPE_WORKGROUP &&
"Control barrier with scope > WORKGROUP");
if (exec_scope == SCOPE_WORKGROUP &&
nak_nir_workgroup_has_one_subgroup(b->shader))
exec_scope = SCOPE_SUBGROUP;
/* Because we're guaranteeing maximal convergence via warp barriers,
* subgroup barriers do nothing.
*/
if (exec_scope <= SCOPE_SUBGROUP)
exec_scope = SCOPE_NONE;
const nir_variable_mode mem_modes = nir_intrinsic_memory_modes(barrier);
if (exec_scope == SCOPE_NONE && mem_modes == 0) {
nir_instr_remove(&barrier->instr);
return true;
}
/* In this case, we're leaving the barrier there */
b->shader->info.uses_control_barrier = true;
bool progress = false;
if (exec_scope != nir_intrinsic_execution_scope(barrier)) {
nir_intrinsic_set_execution_scope(barrier, exec_scope);
progress = true;
}
return progress;
}
static bool
nak_nir_remove_barriers(nir_shader *nir)
{
/* We'll set this back to true if we leave any barriers in place */
nir->info.uses_control_barrier = false;
return nir_shader_intrinsics_pass(nir, nak_nir_remove_barrier_intrin,
nir_metadata_control_flow,
NULL);
}
static bool
nak_mem_vectorize_cb(unsigned align_mul, unsigned align_offset,
unsigned bit_size, unsigned num_components,
nir_intrinsic_instr *low, nir_intrinsic_instr *high,
void *cb_data)
{
/*
* Since we legalize these later with nir_lower_mem_access_bit_sizes,
* we can optimistically combine anything that might be profitable
*/
assert(util_is_power_of_two_nonzero(align_mul));
unsigned max_bytes = 128u / 8u;
if (low->intrinsic == nir_intrinsic_ldc_nv ||
low->intrinsic == nir_intrinsic_ldcx_nv)
max_bytes = 64u / 8u;
align_mul = MIN2(align_mul, max_bytes);
align_offset = align_offset % align_mul;
return align_offset + num_components * (bit_size / 8) <= align_mul;
}
static nir_mem_access_size_align
nak_mem_access_size_align(nir_intrinsic_op intrin,
uint8_t bytes, uint8_t bit_size,
uint32_t align_mul, uint32_t align_offset,
bool offset_is_const, const void *cb_data)
{
const uint32_t align = nir_combined_align(align_mul, align_offset);
assert(util_is_power_of_two_nonzero(align));
unsigned bytes_pow2;
if (nir_intrinsic_infos[intrin].has_dest) {
/* Reads can over-fetch a bit if the alignment is okay. */
bytes_pow2 = util_next_power_of_two(bytes);
} else {
bytes_pow2 = 1 << (util_last_bit(bytes) - 1);
}
unsigned chunk_bytes = MIN3(bytes_pow2, align, 16);
assert(util_is_power_of_two_nonzero(chunk_bytes));
if (intrin == nir_intrinsic_ldc_nv ||
intrin == nir_intrinsic_ldcx_nv)
chunk_bytes = MIN2(chunk_bytes, 8);
if ((intrin == nir_intrinsic_ldc_nv ||
intrin == nir_intrinsic_ldcx_nv) && align < 4) {
/* CBufs require 4B alignment unless we're doing a ldc.u8 or ldc.i8.
* In particular, this applies to ldc.u16 which means we either have to
* fall back to two ldc.u8 or use ldc.u32 and shift stuff around to get
* the 16bit value out. Fortunately, nir_lower_mem_access_bit_sizes()
* can handle over-alignment for reads.
*/
if (align == 2 || offset_is_const) {
return (nir_mem_access_size_align) {
.bit_size = 32,
.num_components = 1,
.align = 4,
};
} else {
assert(align == 1);
return (nir_mem_access_size_align) {
.bit_size = 8,
.num_components = 1,
.align = 1,
};
}
} else if (chunk_bytes < 4) {
return (nir_mem_access_size_align) {
.bit_size = chunk_bytes * 8,
.num_components = 1,
.align = chunk_bytes,
};
} else {
return (nir_mem_access_size_align) {
.bit_size = 32,
.num_components = chunk_bytes / 4,
.align = chunk_bytes,
};
}
}
static bool
nir_shader_has_local_variables(const nir_shader *nir)
{
nir_foreach_function(func, nir) {
if (func->impl && !exec_list_is_empty(&func->impl->locals))
return true;
}
return false;
}
static int
type_size_vec4(const struct glsl_type *type, bool bindless)
{
return glsl_count_vec4_slots(type, false, bindless);
}
void
nak_postprocess_nir(nir_shader *nir,
const struct nak_compiler *nak,
nir_variable_mode robust2_modes,
const struct nak_fs_key *fs_key)
{
UNUSED bool progress = false;
nak_optimize_nir(nir, nak);
const nir_lower_subgroups_options subgroups_options = {
.subgroup_size = NAK_SUBGROUP_SIZE,
.ballot_bit_size = 32,
.ballot_components = 1,
.lower_to_scalar = true,
.lower_vote_eq = true,
.lower_first_invocation_to_ballot = true,
.lower_read_first_invocation = true,
.lower_elect = true,
.lower_inverse_ballot = true,
.lower_rotate_to_shuffle = true
};
OPT(nir, nir_lower_subgroups, &subgroups_options);
OPT(nir, nak_nir_lower_scan_reduce);
if (nir_shader_has_local_variables(nir)) {
OPT(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp,
glsl_get_natural_size_align_bytes);
OPT(nir, nir_lower_explicit_io, nir_var_function_temp,
nir_address_format_32bit_offset);
nak_optimize_nir(nir, nak);
}
OPT(nir, nir_opt_shrink_vectors, true);
nir_load_store_vectorize_options vectorize_opts = {};
vectorize_opts.modes = nir_var_mem_global |
nir_var_mem_ssbo |
nir_var_mem_shared |
nir_var_shader_temp;
vectorize_opts.callback = nak_mem_vectorize_cb;
vectorize_opts.robust_modes = robust2_modes;
OPT(nir, nir_opt_load_store_vectorize, &vectorize_opts);
nir_lower_mem_access_bit_sizes_options mem_bit_size_options = {
.modes = nir_var_mem_constant | nir_var_mem_ubo | nir_var_mem_generic,
.callback = nak_mem_access_size_align,
};
OPT(nir, nir_lower_mem_access_bit_sizes, &mem_bit_size_options);
OPT(nir, nir_lower_bit_size, lower_bit_size_cb, (void *)nak);
OPT(nir, nir_opt_combine_barriers, NULL, NULL);
nak_optimize_nir(nir, nak);
OPT(nir, nak_nir_lower_tex, nak);
OPT(nir, nir_lower_idiv, NULL);
nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir));
OPT(nir, nir_lower_indirect_derefs, 0, UINT32_MAX);
if (nir->info.stage == MESA_SHADER_TESS_EVAL) {
OPT(nir, nir_lower_tess_coord_z,
nir->info.tess._primitive_mode == TESS_PRIMITIVE_TRIANGLES);
}
/* We need to do this before nak_nir_lower_system_values() because it
* relies on the workgroup size being the actual HW workgroup size in
* nir_intrinsic_load_subgroup_id.
*/
if (gl_shader_stage_uses_workgroup(nir->info.stage) &&
nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) {
assert(nir->info.workgroup_size[0] % 2 == 0);
assert(nir->info.workgroup_size[1] % 2 == 0);
nir->info.workgroup_size[0] *= 2;
nir->info.workgroup_size[1] /= 2;
}
OPT(nir, nak_nir_lower_system_values, nak);
switch (nir->info.stage) {
case MESA_SHADER_VERTEX:
case MESA_SHADER_TESS_CTRL:
case MESA_SHADER_TESS_EVAL:
case MESA_SHADER_GEOMETRY:
OPT(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out,
type_size_vec4, nir_lower_io_lower_64bit_to_32_new);
OPT(nir, nir_opt_constant_folding);
OPT(nir, nak_nir_lower_vtg_io, nak);
if (nir->info.stage == MESA_SHADER_GEOMETRY)
OPT(nir, nak_nir_lower_gs_intrinsics);
break;
case MESA_SHADER_FRAGMENT:
OPT(nir, nir_lower_indirect_derefs,
nir_var_shader_in | nir_var_shader_out, UINT32_MAX);
OPT(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out,
type_size_vec4, nir_lower_io_lower_64bit_to_32_new);
OPT(nir, nir_opt_constant_folding);
OPT(nir, nak_nir_lower_fs_inputs, nak, fs_key);
OPT(nir, nak_nir_lower_fs_outputs);
break;
case MESA_SHADER_COMPUTE:
case MESA_SHADER_KERNEL:
break;
default:
unreachable("Unsupported shader stage");
}
OPT(nir, nir_lower_doubles, NULL, nak->nir_options.lower_doubles_options);
OPT(nir, nir_lower_int64);
nak_optimize_nir(nir, nak);
do {
progress = false;
OPT(nir, nir_opt_algebraic_late);
OPT(nir, nak_nir_lower_algebraic_late, nak);
/* If we're lowering fp64 sat but not min/max, the sat lowering may have
* been undone by nir_opt_algebraic. Lower sat again just to be sure.
*/
if ((nak->nir_options.lower_doubles_options & nir_lower_dsat) &&
!(nak->nir_options.lower_doubles_options & nir_lower_dminmax))
OPT(nir, nir_lower_doubles, NULL, nir_lower_dsat);
if (progress) {
OPT(nir, nir_opt_constant_folding);
OPT(nir, nir_copy_prop);
OPT(nir, nir_opt_dce);
OPT(nir, nir_opt_cse);
}
} while (progress);
if (nak->sm < 70)
OPT(nir, nak_nir_split_64bit_conversions);
nir_convert_to_lcssa(nir, true, true);
nir_divergence_analysis(nir);
if (nak->sm >= 75) {
if (OPT(nir, nak_nir_lower_non_uniform_ldcx)) {
OPT(nir, nir_copy_prop);
OPT(nir, nir_opt_dce);
nir_divergence_analysis(nir);
}
}
OPT(nir, nak_nir_remove_barriers);
if (nak->sm >= 70) {
if (nak_should_print_nir()) {
fprintf(stderr, "Structured NIR for %s shader:\n",
_mesa_shader_stage_to_string(nir->info.stage));
nir_print_shader(nir, stderr);
}
OPT(nir, nak_nir_lower_cf);
}
/* Re-index blocks and compact SSA defs because we'll use them to index
* arrays
*/
nir_foreach_function(func, nir) {
if (func->impl) {
nir_index_blocks(func->impl);
nir_index_ssa_defs(func->impl);
}
}
if (nak_should_print_nir()) {
fprintf(stderr, "NIR for %s shader:\n",
_mesa_shader_stage_to_string(nir->info.stage));
nir_print_shader(nir, stderr);
}
}
static bool
scalar_is_imm_int(nir_scalar x, unsigned bits)
{
if (!nir_scalar_is_const(x))
return false;
int64_t imm = nir_scalar_as_int(x);
return u_intN_min(bits) <= imm && imm <= u_intN_max(bits);
}
struct nak_io_addr_offset
nak_get_io_addr_offset(nir_def *addr, uint8_t imm_bits)
{
nir_scalar addr_s = {
.def = addr,
.comp = 0,
};
if (scalar_is_imm_int(addr_s, imm_bits)) {
/* Base is a dumb name for this. It should be offset */
return (struct nak_io_addr_offset) {
.offset = nir_scalar_as_int(addr_s),
};
}
addr_s = nir_scalar_chase_movs(addr_s);
if (!nir_scalar_is_alu(addr_s) ||
nir_scalar_alu_op(addr_s) != nir_op_iadd) {
return (struct nak_io_addr_offset) {
.base = addr_s,
};
}
for (unsigned i = 0; i < 2; i++) {
nir_scalar off_s = nir_scalar_chase_alu_src(addr_s, i);
off_s = nir_scalar_chase_movs(off_s);
if (scalar_is_imm_int(off_s, imm_bits)) {
return (struct nak_io_addr_offset) {
.base = nir_scalar_chase_alu_src(addr_s, 1 - i),
.offset = nir_scalar_as_int(off_s),
};
}
}
return (struct nak_io_addr_offset) {
.base = addr_s,
};
}