AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
feec9166cdb2d562e741a2775b3fa87fc0876707
mesa/src/intel/compiler/brw_mesh.cpp
T
Caio Oliveira 1cdc4be14b intel/compiler: Don't allocate memory for SIMD select error handling 
The position in the error array already indicate the SIMD in question,
so take off all the formatted printing from the errors -- which in some
cases were just not needed.  We lose a little bit of extra context but
it is all easily derivable from the message and the SIMD.

This also will remove the overhead when SIMD selection is being used to
just to find the selected dispatch width -- at a point where the shaders
were already compiled -- and the errors are not used at all.

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9849
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25336>
2023-09-22 16:23:02 +00:00

2092 lines
72 KiB
C++
Raw Blame History

/*
* Copyright © 2021 Intel Corporation
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and associated documentation files (the "Software"),
* to deal in the Software without restriction, including without limitation
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
* and/or sell copies of the Software, and to permit persons to whom the
* Software is furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice (including the next
* paragraph) shall be included in all copies or substantial portions of the
* Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*/
#include <list>
#include <vector>
#include "brw_compiler.h"
#include "brw_fs.h"
#include "brw_nir.h"
#include "brw_private.h"
#include "compiler/nir/nir_builder.h"
#include "dev/intel_debug.h"
#include <memory>
using namespace brw;
static bool
brw_nir_lower_load_uniforms_filter(const nir_instr *instr,
UNUSED const void *data)
{
if (instr->type != nir_instr_type_intrinsic)
return false;
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
return intrin->intrinsic == nir_intrinsic_load_uniform;
}
static nir_def *
brw_nir_lower_load_uniforms_impl(nir_builder *b, nir_instr *instr,
UNUSED void *data)
{
assert(instr->type == nir_instr_type_intrinsic);
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
assert(intrin->intrinsic == nir_intrinsic_load_uniform);
/* Read the first few 32-bit scalars from InlineData. */
if (nir_src_is_const(intrin->src[0]) &&
intrin->def.bit_size == 32 &&
intrin->def.num_components == 1) {
unsigned off = nir_intrinsic_base(intrin) + nir_src_as_uint(intrin->src[0]);
unsigned off_dw = off / 4;
if (off % 4 == 0 && off_dw < BRW_TASK_MESH_PUSH_CONSTANTS_SIZE_DW) {
off_dw += BRW_TASK_MESH_PUSH_CONSTANTS_START_DW;
return nir_load_mesh_inline_data_intel(b, 32, off_dw);
}
}
return brw_nir_load_global_const(b, intrin,
nir_load_mesh_inline_data_intel(b, 64, 0), 0);
}
static bool
brw_nir_lower_load_uniforms(nir_shader *nir)
{
return nir_shader_lower_instructions(nir, brw_nir_lower_load_uniforms_filter,
brw_nir_lower_load_uniforms_impl, NULL);
}
static inline int
type_size_scalar_dwords(const struct glsl_type *type, bool bindless)
{
return glsl_count_dword_slots(type, bindless);
}
/* TODO(mesh): Make this a common function. */
static void
shared_type_info(const struct glsl_type *type, unsigned *size, unsigned *align)
{
assert(glsl_type_is_vector_or_scalar(type));
uint32_t comp_size = glsl_type_is_boolean(type)
? 4 : glsl_get_bit_size(type) / 8;
unsigned length = glsl_get_vector_elements(type);
*size = comp_size * length,
*align = comp_size * (length == 3 ? 4 : length);
}
static bool
brw_nir_lower_launch_mesh_workgroups_instr(nir_builder *b,
nir_intrinsic_instr *intrin,
void *data)
{
if (intrin->intrinsic != nir_intrinsic_launch_mesh_workgroups)
return false;
b->cursor = nir_before_instr(&intrin->instr);
nir_def *local_invocation_index = nir_load_local_invocation_index(b);
/* Make sure that the mesh workgroup size is taken from the first invocation
* (nir_intrinsic_launch_mesh_workgroups requirement)
*/
nir_def *cmp = nir_ieq_imm(b, local_invocation_index, 0);
nir_if *if_stmt = nir_push_if(b, cmp);
{
/* TUE header contains 4 words:
*
* - Word 0 for Task Count.
*
* - Words 1-3 used for "Dispatch Dimensions" feature, to allow mapping a
* 3D dispatch into the 1D dispatch supported by HW.
*/
nir_def *x = nir_channel(b, intrin->src[0].ssa, 0);
nir_def *y = nir_channel(b, intrin->src[0].ssa, 1);
nir_def *z = nir_channel(b, intrin->src[0].ssa, 2);
nir_def *task_count = nir_imul(b, x, nir_imul(b, y, z));
nir_def *tue_header = nir_vec4(b, task_count, x, y, z);
nir_store_task_payload(b, tue_header, nir_imm_int(b, 0));
}
nir_pop_if(b, if_stmt);
nir_instr_remove(&intrin->instr);
return true;
}
static bool
brw_nir_lower_launch_mesh_workgroups(nir_shader *nir)
{
return nir_shader_intrinsics_pass(nir,
brw_nir_lower_launch_mesh_workgroups_instr,
nir_metadata_none,
NULL);
}
static void
brw_nir_lower_tue_outputs(nir_shader *nir, brw_tue_map *map)
{
memset(map, 0, sizeof(*map));
NIR_PASS(_, nir, nir_lower_io, nir_var_shader_out,
type_size_scalar_dwords, nir_lower_io_lower_64bit_to_32);
/* From bspec: "It is suggested that SW reserve the 16 bytes following the
* TUE Header, and therefore start the SW-defined data structure at 32B
* alignment. This allows the TUE Header to always be written as 32 bytes
* with 32B alignment, the most optimal write performance case."
*/
map->per_task_data_start_dw = 8;
/* Lowering to explicit types will start offsets from task_payload_size, so
* set it to start after the header.
*/
nir->info.task_payload_size = map->per_task_data_start_dw * 4;
NIR_PASS(_, nir, nir_lower_vars_to_explicit_types,
nir_var_mem_task_payload, shared_type_info);
NIR_PASS(_, nir, nir_lower_explicit_io,
nir_var_mem_task_payload, nir_address_format_32bit_offset);
map->size_dw = ALIGN(DIV_ROUND_UP(nir->info.task_payload_size, 4), 8);
}
static void
brw_print_tue_map(FILE *fp, const struct brw_tue_map *map)
{
fprintf(fp, "TUE (%d dwords)\n\n", map->size_dw);
}
static bool
brw_nir_adjust_task_payload_offsets_instr(struct nir_builder *b,
nir_intrinsic_instr *intrin,
void *data)
{
switch (intrin->intrinsic) {
case nir_intrinsic_store_task_payload:
case nir_intrinsic_load_task_payload: {
nir_src *offset_src = nir_get_io_offset_src(intrin);
if (nir_src_is_const(*offset_src))
assert(nir_src_as_uint(*offset_src) % 4 == 0);
b->cursor = nir_before_instr(&intrin->instr);
/* Regular I/O uses dwords while explicit I/O used for task payload uses
* bytes. Normalize it to dwords.
*
* TODO(mesh): Figure out how to handle 8-bit, 16-bit.
*/
nir_def *offset = nir_ishr_imm(b, offset_src->ssa, 2);
nir_src_rewrite(offset_src, offset);
unsigned base = nir_intrinsic_base(intrin);
assert(base % 4 == 0);
nir_intrinsic_set_base(intrin, base / 4);
return true;
}
default:
return false;
}
}
static bool
brw_nir_adjust_task_payload_offsets(nir_shader *nir)
{
return nir_shader_intrinsics_pass(nir,
brw_nir_adjust_task_payload_offsets_instr,
nir_metadata_block_index |
nir_metadata_dominance,
NULL);
}
void
brw_nir_adjust_payload(nir_shader *shader, const struct brw_compiler *compiler)
{
/* Adjustment of task payload offsets must be performed *after* last pass
* which interprets them as bytes, because it changes their unit.
*/
bool adjusted = false;
NIR_PASS(adjusted, shader, brw_nir_adjust_task_payload_offsets);
if (adjusted) /* clean up the mess created by offset adjustments */
NIR_PASS(_, shader, nir_opt_constant_folding);
}
static bool
brw_nir_align_launch_mesh_workgroups_instr(nir_builder *b,
nir_intrinsic_instr *intrin,
void *data)
{
if (intrin->intrinsic != nir_intrinsic_launch_mesh_workgroups)
return false;
/* nir_lower_task_shader uses "range" as task payload size. */
unsigned range = nir_intrinsic_range(intrin);
/* This will avoid special case in nir_lower_task_shader dealing with
* not vec4-aligned payload when payload_in_shared workaround is enabled.
*/
nir_intrinsic_set_range(intrin, ALIGN(range, 16));
return true;
}
static bool
brw_nir_align_launch_mesh_workgroups(nir_shader *nir)
{
return nir_shader_intrinsics_pass(nir,
brw_nir_align_launch_mesh_workgroups_instr,
nir_metadata_block_index |
nir_metadata_dominance,
NULL);
}
const unsigned *
brw_compile_task(const struct brw_compiler *compiler,
struct brw_compile_task_params *params)
{
struct nir_shader *nir = params->base.nir;
const struct brw_task_prog_key *key = params->key;
struct brw_task_prog_data *prog_data = params->prog_data;
const bool debug_enabled = brw_should_print_shader(nir, DEBUG_TASK);
brw_nir_lower_tue_outputs(nir, &prog_data->map);
NIR_PASS(_, nir, brw_nir_align_launch_mesh_workgroups);
nir_lower_task_shader_options lower_ts_opt = {
.payload_to_shared_for_atomics = true,
.payload_to_shared_for_small_types = true,
/* The actual payload data starts after the TUE header and padding,
* so skip those when copying.
*/
.payload_offset_in_bytes = prog_data->map.per_task_data_start_dw * 4,
};
NIR_PASS(_, nir, nir_lower_task_shader, lower_ts_opt);
NIR_PASS(_, nir, brw_nir_lower_launch_mesh_workgroups);
prog_data->base.base.stage = MESA_SHADER_TASK;
prog_data->base.base.total_shared = nir->info.shared_size;
prog_data->base.base.total_scratch = 0;
prog_data->base.local_size[0] = nir->info.workgroup_size[0];
prog_data->base.local_size[1] = nir->info.workgroup_size[1];
prog_data->base.local_size[2] = nir->info.workgroup_size[2];
prog_data->uses_drawid =
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID);
brw_simd_selection_state simd_state{
.devinfo = compiler->devinfo,
.prog_data = &prog_data->base,
.required_width = brw_required_dispatch_width(&nir->info),
};
std::unique_ptr<fs_visitor> v[3];
for (unsigned simd = 0; simd < 3; simd++) {
if (!brw_simd_should_compile(simd_state, simd))
continue;
const unsigned dispatch_width = 8 << simd;
nir_shader *shader = nir_shader_clone(params->base.mem_ctx, nir);
brw_nir_apply_key(shader, compiler, &key->base, dispatch_width);
NIR_PASS(_, shader, brw_nir_lower_load_uniforms);
NIR_PASS(_, shader, brw_nir_lower_simd, dispatch_width);
brw_postprocess_nir(shader, compiler, debug_enabled,
key->base.robust_flags);
v[simd] = std::make_unique<fs_visitor>(compiler, &params->base,
&key->base,
&prog_data->base.base,
shader, dispatch_width,
params->base.stats != NULL,
debug_enabled);
if (prog_data->base.prog_mask) {
unsigned first = ffs(prog_data->base.prog_mask) - 1;
v[simd]->import_uniforms(v[first].get());
}
const bool allow_spilling = !brw_simd_any_compiled(simd_state);
if (v[simd]->run_task(allow_spilling))
brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers);
else
simd_state.error[simd] = ralloc_strdup(params->base.mem_ctx, v[simd]->fail_msg);
}
int selected_simd = brw_simd_select(simd_state);
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: "
"SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n",
simd_state.error[0], simd_state.error[1],
simd_state.error[2]);
return NULL;
}
fs_visitor *selected = v[selected_simd].get();
prog_data->base.prog_mask = 1 << selected_simd;
if (unlikely(debug_enabled)) {
fprintf(stderr, "Task Output ");
brw_print_tue_map(stderr, &prog_data->map);
}
fs_generator g(compiler, &params->base, &prog_data->base.base,
false, MESA_SHADER_TASK);
if (unlikely(debug_enabled)) {
g.enable_debug(ralloc_asprintf(params->base.mem_ctx,
"%s task shader %s",
nir->info.label ? nir->info.label
: "unnamed",
nir->info.name));
}
g.generate_code(selected->cfg, selected->dispatch_width, selected->shader_stats,
selected->performance_analysis.require(), params->base.stats);
g.add_const_data(nir->constant_data, nir->constant_data_size);
return g.get_assembly();
}
static void
brw_nir_lower_tue_inputs(nir_shader *nir, const brw_tue_map *map)
{
if (!map)
return;
nir->info.task_payload_size = map->per_task_data_start_dw * 4;
bool progress = false;
NIR_PASS(progress, nir, nir_lower_vars_to_explicit_types,
nir_var_mem_task_payload, shared_type_info);
if (progress) {
/* The types for Task Output and Mesh Input should match, so their sizes
* should also match.
*/
assert(map->size_dw == ALIGN(DIV_ROUND_UP(nir->info.task_payload_size, 4), 8));
} else {
/* Mesh doesn't read any input, to make it clearer set the
* task_payload_size to zero instead of keeping an incomplete size that
* just includes the header.
*/
nir->info.task_payload_size = 0;
}
NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_task_payload,
nir_address_format_32bit_offset);
}
/* Attribute types. Flat attributes have to be a separate class because
* flat and interpolated attributes can't share the same vec4 slot
* (see 3DSTATE_SBE.ConstantInterpolationEnable).
*/
enum {
PRIM, /* per primitive */
VERT, /* per vertex interpolated */
VERT_FLAT, /* per vertex flat */
};
struct attr_desc {
int location;
const struct glsl_type *type;
unsigned dwords;
unsigned slots;
};
struct attr_type_info {
/* order of attributes, negative values are holes */
std::list<struct attr_desc> *order;
/* attributes after which there's hole of size equal to array index */
std::list<int> holes[5];
};
static void
brw_mue_assign_position(const struct attr_desc *attr,
struct brw_mue_map *map,
unsigned start_dw)
{
bool is_array = glsl_type_is_array(attr->type);
int location = attr->location;
unsigned remaining = attr->dwords;
for (unsigned slot = 0; slot < attr->slots; ++slot) {
map->start_dw[location + slot] = start_dw;
unsigned sz;
if (is_array) {
assert(attr->dwords % attr->slots == 0);
sz = attr->dwords / attr->slots;
} else {
sz = MIN2(remaining, 4);
}
map->len_dw[location + slot] = sz;
start_dw += sz;
remaining -= sz;
}
}
static nir_variable *
brw_nir_find_complete_variable_with_location(nir_shader *shader,
nir_variable_mode mode,
int location)
{
nir_variable *best_var = NULL;
unsigned last_size = 0;
nir_foreach_variable_with_modes(var, shader, mode) {
if (var->data.location != location)
continue;
unsigned new_size = glsl_count_dword_slots(var->type, false);
if (new_size > last_size) {
best_var = var;
last_size = new_size;
}
}
return best_var;
}
static unsigned
brw_sum_size(const std::list<struct attr_desc> &orders)
{
unsigned sz = 0;
for (auto it = orders.cbegin(); it != orders.cend(); ++it)
sz += (*it).dwords;
return sz;
}
/* Finds order of outputs which require minimum size, without splitting
* of URB read/write messages (which operate on vec4-aligned memory).
*/
static void
brw_compute_mue_layout(const struct brw_compiler *compiler,
std::list<struct attr_desc> *orders,
uint64_t outputs_written,
struct nir_shader *nir,
bool *pack_prim_data_into_header,
bool *pack_vert_data_into_header)
{
const struct shader_info *info = &nir->info;
struct attr_type_info data[3];
if ((compiler->mesh.mue_header_packing & 1) == 0)
*pack_prim_data_into_header = false;
if ((compiler->mesh.mue_header_packing & 2) == 0)
*pack_vert_data_into_header = false;
for (unsigned i = PRIM; i <= VERT_FLAT; ++i)
data[i].order = &orders[i];
/* If packing into header is enabled, add a hole of size 4 and add
* a virtual location to keep the algorithm happy (it expects holes
* to be preceded by some location). We'll remove those virtual
* locations at the end.
*/
const gl_varying_slot virtual_header_location = VARYING_SLOT_POS;
assert((outputs_written & BITFIELD64_BIT(virtual_header_location)) == 0);
struct attr_desc d;
d.location = virtual_header_location;
d.type = NULL;
d.dwords = 0;
d.slots = 0;
struct attr_desc h;
h.location = -1;
h.type = NULL;
h.dwords = 4;
h.slots = 0;
if (*pack_prim_data_into_header) {
orders[PRIM].push_back(d);
orders[PRIM].push_back(h);
data[PRIM].holes[4].push_back(virtual_header_location);
}
if (*pack_vert_data_into_header) {
orders[VERT].push_back(d);
orders[VERT].push_back(h);
data[VERT].holes[4].push_back(virtual_header_location);
}
u_foreach_bit64(location, outputs_written) {
if ((BITFIELD64_BIT(location) & outputs_written) == 0)
continue;
/* At this point there are both complete and split variables as
* outputs. We need the complete variable to compute the required
* size.
*/
nir_variable *var =
brw_nir_find_complete_variable_with_location(nir,
nir_var_shader_out,
location);
d.location = location;
d.type = brw_nir_get_var_type(nir, var);
d.dwords = glsl_count_dword_slots(d.type, false);
d.slots = glsl_count_attribute_slots(d.type, false);
struct attr_type_info *type_data;
if (BITFIELD64_BIT(location) & info->per_primitive_outputs)
type_data = &data[PRIM];
else if (var->data.interpolation == INTERP_MODE_FLAT)
type_data = &data[VERT_FLAT];
else
type_data = &data[VERT];
std::list<struct attr_desc> *order = type_data->order;
std::list<int> *holes = type_data->holes;
outputs_written &= ~BITFIELD64_RANGE(location, d.slots);
/* special case to use hole of size 4 */
if (d.dwords == 4 && !holes[4].empty()) {
holes[4].pop_back();
assert(order->front().location == virtual_header_location);
order->pop_front();
assert(order->front().location == -1);
assert(order->front().dwords == 4);
order->front() = d;
continue;
}
int mod = d.dwords % 4;
if (mod == 0) {
order->push_back(d);
continue;
}
h.location = -1;
h.type = NULL;
h.dwords = 4 - mod;
h.slots = 0;
if (!compiler->mesh.mue_compaction) {
order->push_back(d);
order->push_back(h);
continue;
}
if (d.dwords > 4) {
order->push_back(d);
order->push_back(h);
holes[h.dwords].push_back(location);
continue;
}
assert(d.dwords < 4);
unsigned found = 0;
/* try to find the smallest hole big enough to hold this attribute */
for (unsigned sz = d.dwords; sz <= 4; sz++){
if (!holes[sz].empty()) {
found = sz;
break;
}
}
/* append at the end if not found */
if (found == 0) {
order->push_back(d);
order->push_back(h);
holes[h.dwords].push_back(location);
continue;
}
assert(found <= 4);
assert(!holes[found].empty());
int after_loc = holes[found].back();
holes[found].pop_back();
bool inserted_back = false;
for (auto it = order->begin(); it != order->end(); ++it) {
if ((*it).location != after_loc)
continue;
++it;
/* must be a hole */
assert((*it).location < 0);
/* and it must be big enough */
assert(d.dwords <= (*it).dwords);
if (d.dwords == (*it).dwords) {
/* exact size, just replace */
*it = d;
} else {
/* inexact size, shrink hole */
(*it).dwords -= d.dwords;
/* and insert new attribute before it */
order->insert(it, d);
/* Insert shrunk hole in a spot so that the order of attributes
* is preserved.
*/
std::list<int> &hole_list = holes[(*it).dwords];
std::list<int>::iterator insert_before = hole_list.end();
for (auto it2 = hole_list.begin(); it2 != hole_list.end(); ++it2) {
if ((*it2) >= (int)location) {
insert_before = it2;
break;
}
}
hole_list.insert(insert_before, location);
}
inserted_back = true;
break;
}
assert(inserted_back);
}
if (*pack_prim_data_into_header) {
if (orders[PRIM].front().location == virtual_header_location)
orders[PRIM].pop_front();
if (!data[PRIM].holes[4].empty()) {
*pack_prim_data_into_header = false;
assert(orders[PRIM].front().location == -1);
assert(orders[PRIM].front().dwords == 4);
orders[PRIM].pop_front();
}
if (*pack_prim_data_into_header) {
unsigned sz = brw_sum_size(orders[PRIM]);
if (sz % 8 == 0 || sz % 8 > 4)
*pack_prim_data_into_header = false;
}
}
if (*pack_vert_data_into_header) {
if (orders[VERT].front().location == virtual_header_location)
orders[VERT].pop_front();
if (!data[VERT].holes[4].empty()) {
*pack_vert_data_into_header = false;
assert(orders[VERT].front().location == -1);
assert(orders[VERT].front().dwords == 4);
orders[VERT].pop_front();
}
if (*pack_vert_data_into_header) {
unsigned sz = brw_sum_size(orders[VERT]) +
brw_sum_size(orders[VERT_FLAT]);
if (sz % 8 == 0 || sz % 8 > 4)
*pack_vert_data_into_header = false;
}
}
if (INTEL_DEBUG(DEBUG_MESH)) {
fprintf(stderr, "MUE attribute order:\n");
for (unsigned i = PRIM; i <= VERT_FLAT; ++i) {
if (!orders[i].empty())
fprintf(stderr, "%d: ", i);
for (auto it = orders[i].cbegin(); it != orders[i].cend(); ++it) {
fprintf(stderr, "%d(%d) ", (*it).location, (*it).dwords);
}
if (!orders[i].empty())
fprintf(stderr, "\n");
}
}
}
/* Mesh URB Entry consists of an initial section
*
* - Primitive Count
* - Primitive Indices (from 0 to Max-1)
* - Padding to 32B if needed
*
* optionally followed by a section for per-primitive data,
* in which each primitive (from 0 to Max-1) gets
*
* - Primitive Header (e.g. ViewportIndex)
* - Primitive Custom Attributes
*
* then followed by a section for per-vertex data
*
* - Vertex Header (e.g. Position)
* - Vertex Custom Attributes
*
* Each per-element section has a pitch and a starting offset. All the
* individual attributes offsets in start_dw are considering the first entry
* of the section (i.e. where the Position for first vertex, or ViewportIndex
* for first primitive). Attributes for other elements are calculated using
* the pitch.
*/
static void
brw_compute_mue_map(const struct brw_compiler *compiler,
struct nir_shader *nir, struct brw_mue_map *map,
enum brw_mesh_index_format index_format, bool compact_mue)
{
memset(map, 0, sizeof(*map));
memset(&map->start_dw[0], -1, sizeof(map->start_dw));
memset(&map->len_dw[0], 0, sizeof(map->len_dw));
unsigned vertices_per_primitive =
num_mesh_vertices_per_primitive(nir->info.mesh.primitive_type);
map->max_primitives = nir->info.mesh.max_primitives_out;
map->max_vertices = nir->info.mesh.max_vertices_out;
uint64_t outputs_written = nir->info.outputs_written;
/* One dword for primitives count then K extra dwords for each primitive. */
switch (index_format) {
case BRW_INDEX_FORMAT_U32:
map->per_primitive_indices_dw = vertices_per_primitive;
break;
case BRW_INDEX_FORMAT_U888X:
map->per_primitive_indices_dw = 1;
break;
default:
unreachable("invalid index format");
}
map->per_primitive_start_dw = ALIGN(map->per_primitive_indices_dw *
map->max_primitives + 1, 8);
/* Assign initial section. */
if (BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_COUNT) & outputs_written) {
map->start_dw[VARYING_SLOT_PRIMITIVE_COUNT] = 0;
map->len_dw[VARYING_SLOT_PRIMITIVE_COUNT] = 1;
outputs_written &= ~BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_COUNT);
}
if (BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_INDICES) & outputs_written) {
map->start_dw[VARYING_SLOT_PRIMITIVE_INDICES] = 1;
map->len_dw[VARYING_SLOT_PRIMITIVE_INDICES] =
map->per_primitive_indices_dw * map->max_primitives;
outputs_written &= ~BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_INDICES);
}
const uint64_t per_primitive_header_bits =
BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE) |
BITFIELD64_BIT(VARYING_SLOT_LAYER) |
BITFIELD64_BIT(VARYING_SLOT_VIEWPORT) |
BITFIELD64_BIT(VARYING_SLOT_CULL_PRIMITIVE);
const uint64_t per_vertex_header_bits =
BITFIELD64_BIT(VARYING_SLOT_PSIZ) |
BITFIELD64_BIT(VARYING_SLOT_POS) |
BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST0) |
BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST1);
std::list<struct attr_desc> orders[3];
uint64_t regular_outputs = outputs_written &
~(per_primitive_header_bits | per_vertex_header_bits);
/* packing into prim header is possible only if prim header is present */
map->user_data_in_primitive_header = compact_mue &&
(outputs_written & per_primitive_header_bits) != 0;
/* Packing into vert header is always possible, but we allow it only
* if full vec4 is available (so point size is not used) and there's
* nothing between it and normal vertex data (so no clip distances).
*/
map->user_data_in_vertex_header = compact_mue &&
(outputs_written & per_vertex_header_bits) ==
BITFIELD64_BIT(VARYING_SLOT_POS);
if (outputs_written & per_primitive_header_bits) {
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE)) {
map->start_dw[VARYING_SLOT_PRIMITIVE_SHADING_RATE] =
map->per_primitive_start_dw + 0;
map->len_dw[VARYING_SLOT_PRIMITIVE_SHADING_RATE] = 1;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_LAYER)) {
map->start_dw[VARYING_SLOT_LAYER] =
map->per_primitive_start_dw + 1; /* RTAIndex */
map->len_dw[VARYING_SLOT_LAYER] = 1;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_VIEWPORT)) {
map->start_dw[VARYING_SLOT_VIEWPORT] =
map->per_primitive_start_dw + 2;
map->len_dw[VARYING_SLOT_VIEWPORT] = 1;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_CULL_PRIMITIVE)) {
map->start_dw[VARYING_SLOT_CULL_PRIMITIVE] =
map->per_primitive_start_dw + 3;
map->len_dw[VARYING_SLOT_CULL_PRIMITIVE] = 1;
}
map->per_primitive_header_size_dw = 8;
outputs_written &= ~per_primitive_header_bits;
} else {
map->per_primitive_header_size_dw = 0;
}
map->per_primitive_data_size_dw = 0;
/* For fast linked libraries, we can't pack the MUE, as the fragment shader
* will be compiled without access to the MUE map and won't be able to find
* out where everything is.
* Instead, keep doing things as we did before the packing, just laying out
* everything in varying order, which is how the FS will expect them.
*/
if (compact_mue) {
brw_compute_mue_layout(compiler, orders, regular_outputs, nir,
&map->user_data_in_primitive_header,
&map->user_data_in_vertex_header);
unsigned start_dw = map->per_primitive_start_dw;
if (map->user_data_in_primitive_header)
start_dw += 4; /* first 4 dwords are used */
else
start_dw += map->per_primitive_header_size_dw;
unsigned header_used_dw = 0;
for (auto it = orders[PRIM].cbegin(); it != orders[PRIM].cend(); ++it) {
int location = (*it).location;
if (location < 0) {
start_dw += (*it).dwords;
if (map->user_data_in_primitive_header && header_used_dw < 4)
header_used_dw += (*it).dwords;
else
map->per_primitive_data_size_dw += (*it).dwords;
assert(header_used_dw <= 4);
continue;
}
assert(map->start_dw[location] == -1);
assert(location == VARYING_SLOT_PRIMITIVE_ID ||
location >= VARYING_SLOT_VAR0);
brw_mue_assign_position(&*it, map, start_dw);
start_dw += (*it).dwords;
if (map->user_data_in_primitive_header && header_used_dw < 4)
header_used_dw += (*it).dwords;
else
map->per_primitive_data_size_dw += (*it).dwords;
assert(header_used_dw <= 4);
outputs_written &= ~BITFIELD64_RANGE(location, (*it).slots);
}
} else {
unsigned start_dw = map->per_primitive_start_dw +
map->per_primitive_header_size_dw;
uint64_t per_prim_outputs = outputs_written & nir->info.per_primitive_outputs;
while (per_prim_outputs) {
uint64_t location = ffsll(per_prim_outputs) - 1;
assert(map->start_dw[location] == -1);
assert(location == VARYING_SLOT_PRIMITIVE_ID ||
location >= VARYING_SLOT_VAR0);
nir_variable *var =
brw_nir_find_complete_variable_with_location(nir,
nir_var_shader_out,
location);
struct attr_desc d;
d.location = location;
d.type = brw_nir_get_var_type(nir, var);
d.dwords = glsl_count_dword_slots(d.type, false);
d.slots = glsl_count_attribute_slots(d.type, false);
brw_mue_assign_position(&d, map, start_dw);
map->per_primitive_data_size_dw += ALIGN(d.dwords, 4);
start_dw += ALIGN(d.dwords, 4);
per_prim_outputs &= ~BITFIELD64_RANGE(location, d.slots);
}
}
map->per_primitive_pitch_dw = ALIGN(map->per_primitive_header_size_dw +
map->per_primitive_data_size_dw, 8);
map->per_vertex_start_dw = ALIGN(map->per_primitive_start_dw +
map->per_primitive_pitch_dw *
map->max_primitives, 8);
/* TODO(mesh): Multiview. */
unsigned fixed_header_size = 8;
map->per_vertex_header_size_dw = ALIGN(fixed_header_size +
nir->info.clip_distance_array_size +
nir->info.cull_distance_array_size, 8);
if (outputs_written & per_vertex_header_bits) {
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_PSIZ)) {
map->start_dw[VARYING_SLOT_PSIZ] = map->per_vertex_start_dw + 3;
map->len_dw[VARYING_SLOT_PSIZ] = 1;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_POS)) {
map->start_dw[VARYING_SLOT_POS] = map->per_vertex_start_dw + 4;
map->len_dw[VARYING_SLOT_POS] = 4;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST0)) {
map->start_dw[VARYING_SLOT_CLIP_DIST0] =
map->per_vertex_start_dw + fixed_header_size + 0;
map->len_dw[VARYING_SLOT_CLIP_DIST0] = 4;
}
if (outputs_written & BITFIELD64_BIT(VARYING_SLOT_CLIP_DIST1)) {
map->start_dw[VARYING_SLOT_CLIP_DIST1] =
map->per_vertex_start_dw + fixed_header_size + 4;
map->len_dw[VARYING_SLOT_CLIP_DIST1] = 4;
}
outputs_written &= ~per_vertex_header_bits;
}
/* cull distances should be lowered earlier */
assert(!(outputs_written & BITFIELD64_BIT(VARYING_SLOT_CULL_DIST0)));
assert(!(outputs_written & BITFIELD64_BIT(VARYING_SLOT_CULL_DIST1)));
map->per_vertex_data_size_dw = 0;
/* For fast linked libraries, we can't pack the MUE, as the fragment shader
* will be compiled without access to the MUE map and won't be able to find
* out where everything is.
* Instead, keep doing things as we did before the packing, just laying out
* everything in varying order, which is how the FS will expect them.
*/
if (compact_mue) {
unsigned start_dw = map->per_vertex_start_dw;
if (!map->user_data_in_vertex_header)
start_dw += map->per_vertex_header_size_dw;
unsigned header_used_dw = 0;
for (unsigned type = VERT; type <= VERT_FLAT; ++type) {
for (auto it = orders[type].cbegin(); it != orders[type].cend(); ++it) {
int location = (*it).location;
if (location < 0) {
start_dw += (*it).dwords;
if (map->user_data_in_vertex_header && header_used_dw < 4) {
header_used_dw += (*it).dwords;
assert(header_used_dw <= 4);
if (header_used_dw == 4)
start_dw += 4; /* jump over gl_position */
} else {
map->per_vertex_data_size_dw += (*it).dwords;
}
continue;
}
assert(map->start_dw[location] == -1);
assert(location >= VARYING_SLOT_VAR0);
brw_mue_assign_position(&*it, map, start_dw);
start_dw += (*it).dwords;
if (map->user_data_in_vertex_header && header_used_dw < 4) {
header_used_dw += (*it).dwords;
assert(header_used_dw <= 4);
if (header_used_dw == 4)
start_dw += 4; /* jump over gl_position */
} else {
map->per_vertex_data_size_dw += (*it).dwords;
}
outputs_written &= ~BITFIELD64_RANGE(location, (*it).slots);
}
}
} else {
unsigned start_dw = map->per_vertex_start_dw +
map->per_vertex_header_size_dw;
uint64_t per_vertex_outputs = outputs_written & ~nir->info.per_primitive_outputs;
while (per_vertex_outputs) {
uint64_t location = ffsll(per_vertex_outputs) - 1;
assert(map->start_dw[location] == -1);
assert(location >= VARYING_SLOT_VAR0);
nir_variable *var =
brw_nir_find_complete_variable_with_location(nir,
nir_var_shader_out,
location);
struct attr_desc d;
d.location = location;
d.type = brw_nir_get_var_type(nir, var);
d.dwords = glsl_count_dword_slots(d.type, false);
d.slots = glsl_count_attribute_slots(d.type, false);
brw_mue_assign_position(&d, map, start_dw);
map->per_vertex_data_size_dw += ALIGN(d.dwords, 4);
start_dw += ALIGN(d.dwords, 4);
per_vertex_outputs &= ~BITFIELD64_RANGE(location, d.slots);
}
}
map->per_vertex_pitch_dw = ALIGN(map->per_vertex_header_size_dw +
map->per_vertex_data_size_dw, 8);
map->size_dw =
map->per_vertex_start_dw + map->per_vertex_pitch_dw * map->max_vertices;
assert(map->size_dw % 8 == 0);
}
static void
brw_print_mue_map(FILE *fp, const struct brw_mue_map *map, struct nir_shader *nir)
{
fprintf(fp, "MUE map (%d dwords, %d primitives, %d vertices)\n",
map->size_dw, map->max_primitives, map->max_vertices);
fprintf(fp, " <%4d, %4d>: VARYING_SLOT_PRIMITIVE_COUNT\n",
map->start_dw[VARYING_SLOT_PRIMITIVE_COUNT],
map->start_dw[VARYING_SLOT_PRIMITIVE_COUNT] +
map->len_dw[VARYING_SLOT_PRIMITIVE_COUNT] - 1);
fprintf(fp, " <%4d, %4d>: VARYING_SLOT_PRIMITIVE_INDICES\n",
map->start_dw[VARYING_SLOT_PRIMITIVE_INDICES],
map->start_dw[VARYING_SLOT_PRIMITIVE_INDICES] +
map->len_dw[VARYING_SLOT_PRIMITIVE_INDICES] - 1);
fprintf(fp, " ----- per primitive (start %d, header_size %d, data_size %d, pitch %d)\n",
map->per_primitive_start_dw,
map->per_primitive_header_size_dw,
map->per_primitive_data_size_dw,
map->per_primitive_pitch_dw);
for (unsigned i = 0; i < VARYING_SLOT_MAX; i++) {
if (map->start_dw[i] < 0)
continue;
const unsigned offset = map->start_dw[i];
const unsigned len = map->len_dw[i];
if (offset < map->per_primitive_start_dw ||
offset >= map->per_primitive_start_dw + map->per_primitive_pitch_dw)
continue;
const char *name =
gl_varying_slot_name_for_stage((gl_varying_slot)i,
MESA_SHADER_MESH);
fprintf(fp, " <%4d, %4d>: %s (%d)\n", offset, offset + len - 1,
name, i);
}
fprintf(fp, " ----- per vertex (start %d, header_size %d, data_size %d, pitch %d)\n",
map->per_vertex_start_dw,
map->per_vertex_header_size_dw,
map->per_vertex_data_size_dw,
map->per_vertex_pitch_dw);
for (unsigned i = 0; i < VARYING_SLOT_MAX; i++) {
if (map->start_dw[i] < 0)
continue;
const unsigned offset = map->start_dw[i];
const unsigned len = map->len_dw[i];
if (offset < map->per_vertex_start_dw ||
offset >= map->per_vertex_start_dw + map->per_vertex_pitch_dw)
continue;
nir_variable *var =
nir_find_variable_with_location(nir, nir_var_shader_out, i);
bool flat = var->data.interpolation == INTERP_MODE_FLAT;
const char *name =
gl_varying_slot_name_for_stage((gl_varying_slot)i,
MESA_SHADER_MESH);
fprintf(fp, " <%4d, %4d>: %s (%d)%s\n", offset, offset + len - 1,
name, i, flat ? " (flat)" : "");
}
fprintf(fp, "\n");
}
static void
brw_nir_lower_mue_outputs(nir_shader *nir, const struct brw_mue_map *map)
{
nir_foreach_shader_out_variable(var, nir) {
int location = var->data.location;
assert(location >= 0);
assert(map->start_dw[location] != -1);
var->data.driver_location = map->start_dw[location];
}
NIR_PASS(_, nir, nir_lower_io, nir_var_shader_out,
type_size_scalar_dwords, nir_lower_io_lower_64bit_to_32);
}
static void
brw_nir_initialize_mue(nir_shader *nir,
const struct brw_mue_map *map,
unsigned dispatch_width)
{
assert(map->per_primitive_header_size_dw > 0);
nir_builder b;
nir_function_impl *entrypoint = nir_shader_get_entrypoint(nir);
b = nir_builder_at(nir_before_impl(entrypoint));
nir_def *dw_off = nir_imm_int(&b, 0);
nir_def *zerovec = nir_imm_vec4(&b, 0, 0, 0, 0);
/* TODO(mesh): can we write in bigger batches, generating fewer SENDs? */
assert(!nir->info.workgroup_size_variable);
const unsigned workgroup_size = nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
nir->info.workgroup_size[2];
/* Invocations from a single workgroup will cooperate in zeroing MUE. */
/* How many prims each invocation needs to cover without checking its index? */
unsigned prims_per_inv = map->max_primitives / workgroup_size;
/* Zero first 4 dwords of MUE Primitive Header:
* Reserved, RTAIndex, ViewportIndex, CullPrimitiveMask.
*/
nir_def *local_invocation_index = nir_load_local_invocation_index(&b);
/* Zero primitive headers distanced by workgroup_size, starting from
* invocation index.
*/
for (unsigned prim_in_inv = 0; prim_in_inv < prims_per_inv; ++prim_in_inv) {
nir_def *prim = nir_iadd_imm(&b, local_invocation_index,
prim_in_inv * workgroup_size);
nir_store_per_primitive_output(&b, zerovec, prim, dw_off,
.base = (int)map->per_primitive_start_dw,
.write_mask = WRITEMASK_XYZW,
.component = 0,
.src_type = nir_type_uint32);
}
/* How many prims are left? */
unsigned remaining = map->max_primitives % workgroup_size;
if (remaining) {
/* Zero "remaining" primitive headers starting from the last one covered
* by the loop above + workgroup_size.
*/
nir_def *cmp = nir_ilt_imm(&b, local_invocation_index, remaining);
nir_if *if_stmt = nir_push_if(&b, cmp);
{
nir_def *prim = nir_iadd_imm(&b, local_invocation_index,
prims_per_inv * workgroup_size);
nir_store_per_primitive_output(&b, zerovec, prim, dw_off,
.base = (int)map->per_primitive_start_dw,
.write_mask = WRITEMASK_XYZW,
.component = 0,
.src_type = nir_type_uint32);
}
nir_pop_if(&b, if_stmt);
}
/* If there's more than one subgroup, then we need to wait for all of them
* to finish initialization before we can proceed. Otherwise some subgroups
* may start filling MUE before other finished initializing.
*/
if (workgroup_size > dispatch_width) {
nir_barrier(&b, SCOPE_WORKGROUP, SCOPE_WORKGROUP,
NIR_MEMORY_ACQ_REL, nir_var_shader_out);
}
if (remaining) {
nir_metadata_preserve(entrypoint, nir_metadata_none);
} else {
nir_metadata_preserve(entrypoint, nir_metadata_block_index |
nir_metadata_dominance);
}
}
static void
brw_nir_adjust_offset(nir_builder *b, nir_intrinsic_instr *intrin, uint32_t pitch)
{
nir_src *index_src = nir_get_io_arrayed_index_src(intrin);
nir_src *offset_src = nir_get_io_offset_src(intrin);
b->cursor = nir_before_instr(&intrin->instr);
nir_def *offset =
nir_iadd(b,
offset_src->ssa,
nir_imul_imm(b, index_src->ssa, pitch));
nir_src_rewrite(offset_src, offset);
}
static bool
brw_nir_adjust_offset_for_arrayed_indices_instr(nir_builder *b,
nir_intrinsic_instr *intrin,
void *data)
{
const struct brw_mue_map *map = (const struct brw_mue_map *) data;
/* Remap per_vertex and per_primitive offsets using the extra source and
* the pitch.
*/
switch (intrin->intrinsic) {
case nir_intrinsic_load_per_vertex_output:
case nir_intrinsic_store_per_vertex_output:
brw_nir_adjust_offset(b, intrin, map->per_vertex_pitch_dw);
return true;
case nir_intrinsic_load_per_primitive_output:
case nir_intrinsic_store_per_primitive_output: {
struct nir_io_semantics sem = nir_intrinsic_io_semantics(intrin);
uint32_t pitch;
if (sem.location == VARYING_SLOT_PRIMITIVE_INDICES)
pitch = map->per_primitive_indices_dw;
else
pitch = map->per_primitive_pitch_dw;
brw_nir_adjust_offset(b, intrin, pitch);
return true;
}
default:
return false;
}
}
static bool
brw_nir_adjust_offset_for_arrayed_indices(nir_shader *nir, const struct brw_mue_map *map)
{
return nir_shader_intrinsics_pass(nir,
brw_nir_adjust_offset_for_arrayed_indices_instr,
nir_metadata_block_index |
nir_metadata_dominance,
(void *)map);
}
struct index_packing_state {
unsigned vertices_per_primitive;
nir_variable *original_prim_indices;
nir_variable *packed_prim_indices;
};
static bool
brw_can_pack_primitive_indices(nir_shader *nir, struct index_packing_state *state)
{
/* can single index fit into one byte of U888X format? */
if (nir->info.mesh.max_vertices_out > 255)
return false;
state->vertices_per_primitive =
num_mesh_vertices_per_primitive(nir->info.mesh.primitive_type);
/* packing point indices doesn't help */
if (state->vertices_per_primitive == 1)
return false;
state->original_prim_indices =
nir_find_variable_with_location(nir,
nir_var_shader_out,
VARYING_SLOT_PRIMITIVE_INDICES);
/* no indices = no changes to the shader, but it's still worth it,
* because less URB space will be used
*/
if (!state->original_prim_indices)
return true;
ASSERTED const struct glsl_type *type = state->original_prim_indices->type;
assert(type->is_array());
assert(type->without_array()->is_vector());
assert(type->without_array()->vector_elements == state->vertices_per_primitive);
nir_foreach_function_impl(impl, nir) {
nir_foreach_block(block, impl) {
nir_foreach_instr(instr, block) {
if (instr->type != nir_instr_type_intrinsic)
continue;
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
if (intrin->intrinsic != nir_intrinsic_store_deref) {
/* any unknown deref operation on primitive indices -> don't pack */
unsigned num_srcs = nir_intrinsic_infos[intrin->intrinsic].num_srcs;
for (unsigned i = 0; i < num_srcs; i++) {
nir_deref_instr *deref = nir_src_as_deref(intrin->src[i]);
if (!deref)
continue;
nir_variable *var = nir_deref_instr_get_variable(deref);
if (var == state->original_prim_indices)
return false;
}
continue;
}
nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]);
if (!deref)
continue;
nir_variable *var = nir_deref_instr_get_variable(deref);
if (var != state->original_prim_indices)
continue;
if (deref->deref_type != nir_deref_type_array)
return false; /* unknown chain of derefs */
nir_deref_instr *var_deref = nir_src_as_deref(deref->parent);
if (!var_deref || var_deref->deref_type != nir_deref_type_var)
return false; /* unknown chain of derefs */
assert (var_deref->var == state->original_prim_indices);
unsigned write_mask = nir_intrinsic_write_mask(intrin);
/* If only some components are written, then we can't easily pack.
* In theory we could, by loading current dword value, bitmasking
* one byte and storing back the whole dword, but it would be slow
* and could actually decrease performance. TODO: reevaluate this
* once there will be something hitting this.
*/
if (write_mask != BITFIELD_MASK(state->vertices_per_primitive))
return false;
}
}
}
return true;
}
static bool
brw_pack_primitive_indices_instr(nir_builder *b, nir_intrinsic_instr *intrin,
void *data)
{
if (intrin->intrinsic != nir_intrinsic_store_deref)
return false;
nir_deref_instr *array_deref = nir_src_as_deref(intrin->src[0]);
if (!array_deref || array_deref->deref_type != nir_deref_type_array)
return false;
nir_deref_instr *var_deref = nir_src_as_deref(array_deref->parent);
if (!var_deref || var_deref->deref_type != nir_deref_type_var)
return false;
struct index_packing_state *state =
(struct index_packing_state *)data;
nir_variable *var = var_deref->var;
if (var != state->original_prim_indices)
return false;
unsigned vertices_per_primitive = state->vertices_per_primitive;
b->cursor = nir_before_instr(&intrin->instr);
nir_deref_instr *new_var_deref =
nir_build_deref_var(b, state->packed_prim_indices);
nir_deref_instr *new_array_deref =
nir_build_deref_array(b, new_var_deref, array_deref->arr.index.ssa);
nir_src *data_src = &intrin->src[1];
nir_def *data_def =
data_src->ssa;
nir_def *new_data =
nir_ior(b, nir_ishl_imm(b, nir_channel(b, data_def, 0), 0),
nir_ishl_imm(b, nir_channel(b, data_def, 1), 8));
if (vertices_per_primitive >= 3) {
new_data =
nir_ior(b, new_data,
nir_ishl_imm(b, nir_channel(b, data_def, 2), 16));
}
nir_build_store_deref(b, &new_array_deref->def, new_data);
nir_instr_remove(&intrin->instr);
return true;
}
static bool
brw_pack_primitive_indices(nir_shader *nir, void *data)
{
struct index_packing_state *state = (struct index_packing_state *)data;
const struct glsl_type *new_type =
glsl_array_type(glsl_uint_type(),
nir->info.mesh.max_primitives_out,
0);
state->packed_prim_indices =
nir_variable_create(nir, nir_var_shader_out,
new_type, "gl_PrimitiveIndicesPacked");
state->packed_prim_indices->data.location = VARYING_SLOT_PRIMITIVE_INDICES;
state->packed_prim_indices->data.interpolation = INTERP_MODE_NONE;
state->packed_prim_indices->data.per_primitive = 1;
return nir_shader_intrinsics_pass(nir, brw_pack_primitive_indices_instr,
nir_metadata_block_index |
nir_metadata_dominance,
data);
}
const unsigned *
brw_compile_mesh(const struct brw_compiler *compiler,
struct brw_compile_mesh_params *params)
{
struct nir_shader *nir = params->base.nir;
const struct brw_mesh_prog_key *key = params->key;
struct brw_mesh_prog_data *prog_data = params->prog_data;
const bool debug_enabled = brw_should_print_shader(nir, DEBUG_MESH);
prog_data->base.base.stage = MESA_SHADER_MESH;
prog_data->base.base.total_shared = nir->info.shared_size;
prog_data->base.base.total_scratch = 0;
prog_data->base.local_size[0] = nir->info.workgroup_size[0];
prog_data->base.local_size[1] = nir->info.workgroup_size[1];
prog_data->base.local_size[2] = nir->info.workgroup_size[2];
prog_data->clip_distance_mask = (1 << nir->info.clip_distance_array_size) - 1;
prog_data->cull_distance_mask =
((1 << nir->info.cull_distance_array_size) - 1) <<
nir->info.clip_distance_array_size;
prog_data->primitive_type = nir->info.mesh.primitive_type;
struct index_packing_state index_packing_state = {};
if (brw_can_pack_primitive_indices(nir, &index_packing_state)) {
if (index_packing_state.original_prim_indices)
NIR_PASS(_, nir, brw_pack_primitive_indices, &index_packing_state);
prog_data->index_format = BRW_INDEX_FORMAT_U888X;
} else {
prog_data->index_format = BRW_INDEX_FORMAT_U32;
}
prog_data->uses_drawid =
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID);
brw_nir_lower_tue_inputs(nir, params->tue_map);
brw_compute_mue_map(compiler, nir, &prog_data->map,
prog_data->index_format, key->compact_mue);
brw_nir_lower_mue_outputs(nir, &prog_data->map);
brw_simd_selection_state simd_state{
.devinfo = compiler->devinfo,
.prog_data = &prog_data->base,
.required_width = brw_required_dispatch_width(&nir->info),
};
std::unique_ptr<fs_visitor> v[3];
for (int simd = 0; simd < 3; simd++) {
if (!brw_simd_should_compile(simd_state, simd))
continue;
const unsigned dispatch_width = 8 << simd;
nir_shader *shader = nir_shader_clone(params->base.mem_ctx, nir);
/*
* When Primitive Header is enabled, we may not generates writes to all
* fields, so let's initialize everything.
*/
if (prog_data->map.per_primitive_header_size_dw > 0)
NIR_PASS_V(shader, brw_nir_initialize_mue, &prog_data->map, dispatch_width);
brw_nir_apply_key(shader, compiler, &key->base, dispatch_width);
NIR_PASS(_, shader, brw_nir_adjust_offset_for_arrayed_indices, &prog_data->map);
/* Load uniforms can do a better job for constants, so fold before it. */
NIR_PASS(_, shader, nir_opt_constant_folding);
NIR_PASS(_, shader, brw_nir_lower_load_uniforms);
NIR_PASS(_, shader, brw_nir_lower_simd, dispatch_width);
brw_postprocess_nir(shader, compiler, debug_enabled,
key->base.robust_flags);
v[simd] = std::make_unique<fs_visitor>(compiler, &params->base,
&key->base,
&prog_data->base.base,
shader, dispatch_width,
params->base.stats != NULL,
debug_enabled);
if (prog_data->base.prog_mask) {
unsigned first = ffs(prog_data->base.prog_mask) - 1;
v[simd]->import_uniforms(v[first].get());
}
const bool allow_spilling = !brw_simd_any_compiled(simd_state);
if (v[simd]->run_mesh(allow_spilling))
brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers);
else
simd_state.error[simd] = ralloc_strdup(params->base.mem_ctx, v[simd]->fail_msg);
}
int selected_simd = brw_simd_select(simd_state);
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: "
"SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n",
simd_state.error[0], simd_state.error[1],
simd_state.error[2]);
return NULL;
}
fs_visitor *selected = v[selected_simd].get();
prog_data->base.prog_mask = 1 << selected_simd;
if (unlikely(debug_enabled)) {
if (params->tue_map) {
fprintf(stderr, "Mesh Input ");
brw_print_tue_map(stderr, params->tue_map);
}
fprintf(stderr, "Mesh Output ");
brw_print_mue_map(stderr, &prog_data->map, nir);
}
fs_generator g(compiler, &params->base, &prog_data->base.base,
false, MESA_SHADER_MESH);
if (unlikely(debug_enabled)) {
g.enable_debug(ralloc_asprintf(params->base.mem_ctx,
"%s mesh shader %s",
nir->info.label ? nir->info.label
: "unnamed",
nir->info.name));
}
g.generate_code(selected->cfg, selected->dispatch_width, selected->shader_stats,
selected->performance_analysis.require(), params->base.stats);
g.add_const_data(nir->constant_data, nir->constant_data_size);
return g.get_assembly();
}
static unsigned
component_from_intrinsic(nir_intrinsic_instr *instr)
{
if (nir_intrinsic_has_component(instr))
return nir_intrinsic_component(instr);
else
return 0;
}
static void
adjust_handle_and_offset(const fs_builder &bld,
fs_reg &urb_handle,
unsigned &urb_global_offset)
{
/* Make sure that URB global offset is below 2048 (2^11), because
* that's the maximum possible value encoded in Message Descriptor.
*/
unsigned adjustment = (urb_global_offset >> 11) << 11;
if (adjustment) {
fs_builder ubld8 = bld.group(8, 0).exec_all();
/* Allocate new register to not overwrite the shared URB handle. */
fs_reg new_handle = ubld8.vgrf(BRW_REGISTER_TYPE_UD);
ubld8.ADD(new_handle, urb_handle, brw_imm_ud(adjustment));
urb_handle = new_handle;
urb_global_offset -= adjustment;
}
}
static void
emit_urb_direct_vec4_write(const fs_builder &bld,
unsigned urb_global_offset,
const fs_reg &src,
fs_reg urb_handle,
unsigned dst_comp_offset,
unsigned comps,
unsigned mask)
{
for (unsigned q = 0; q < bld.dispatch_width() / 8; q++) {
fs_builder bld8 = bld.group(8, q);
fs_reg payload_srcs[8];
unsigned length = 0;
for (unsigned i = 0; i < dst_comp_offset; i++)
payload_srcs[length++] = reg_undef;
for (unsigned c = 0; c < comps; c++)
payload_srcs[length++] = quarter(offset(src, bld, c), q);
fs_reg srcs[URB_LOGICAL_NUM_SRCS];
srcs[URB_LOGICAL_SRC_HANDLE] = urb_handle;
srcs[URB_LOGICAL_SRC_CHANNEL_MASK] = brw_imm_ud(mask << 16);
srcs[URB_LOGICAL_SRC_DATA] = fs_reg(VGRF, bld.shader->alloc.allocate(length),
BRW_REGISTER_TYPE_F);
bld8.LOAD_PAYLOAD(srcs[URB_LOGICAL_SRC_DATA], payload_srcs, length, 0);
fs_inst *inst = bld8.emit(SHADER_OPCODE_URB_WRITE_LOGICAL,
reg_undef, srcs, ARRAY_SIZE(srcs));
inst->mlen = 2 + length;
inst->offset = urb_global_offset;
assert(inst->offset < 2048);
}
}
static void
emit_urb_direct_writes(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &src, fs_reg urb_handle)
{
assert(nir_src_bit_size(instr->src[0]) == 32);
nir_src *offset_nir_src = nir_get_io_offset_src(instr);
assert(nir_src_is_const(*offset_nir_src));
const unsigned comps = nir_src_num_components(instr->src[0]);
assert(comps <= 4);
const unsigned offset_in_dwords = nir_intrinsic_base(instr) +
nir_src_as_uint(*offset_nir_src) +
component_from_intrinsic(instr);
/* URB writes are vec4 aligned but the intrinsic offsets are in dwords.
* We can write up to 8 dwords, so single vec4 write is enough.
*/
const unsigned comp_shift = offset_in_dwords % 4;
const unsigned mask = nir_intrinsic_write_mask(instr) << comp_shift;
unsigned urb_global_offset = offset_in_dwords / 4;
adjust_handle_and_offset(bld, urb_handle, urb_global_offset);
emit_urb_direct_vec4_write(bld, urb_global_offset, src, urb_handle,
comp_shift, comps, mask);
}
static void
emit_urb_indirect_vec4_write(const fs_builder &bld,
const fs_reg &offset_src,
unsigned base,
const fs_reg &src,
fs_reg urb_handle,
unsigned dst_comp_offset,
unsigned comps,
unsigned mask)
{
for (unsigned q = 0; q < bld.dispatch_width() / 8; q++) {
fs_builder bld8 = bld.group(8, q);
/* offset is always positive, so signedness doesn't matter */
assert(offset_src.type == BRW_REGISTER_TYPE_D ||
offset_src.type == BRW_REGISTER_TYPE_UD);
fs_reg off = bld8.vgrf(offset_src.type, 1);
bld8.MOV(off, quarter(offset_src, q));
bld8.ADD(off, off, brw_imm_ud(base));
bld8.SHR(off, off, brw_imm_ud(2));
fs_reg payload_srcs[8];
unsigned length = 0;
for (unsigned i = 0; i < dst_comp_offset; i++)
payload_srcs[length++] = reg_undef;
for (unsigned c = 0; c < comps; c++)
payload_srcs[length++] = quarter(offset(src, bld, c), q);
fs_reg srcs[URB_LOGICAL_NUM_SRCS];
srcs[URB_LOGICAL_SRC_HANDLE] = urb_handle;
srcs[URB_LOGICAL_SRC_PER_SLOT_OFFSETS] = off;
srcs[URB_LOGICAL_SRC_CHANNEL_MASK] = brw_imm_ud(mask << 16);
srcs[URB_LOGICAL_SRC_DATA] = fs_reg(VGRF, bld.shader->alloc.allocate(length),
BRW_REGISTER_TYPE_F);
bld8.LOAD_PAYLOAD(srcs[URB_LOGICAL_SRC_DATA], payload_srcs, length, 0);
fs_inst *inst = bld8.emit(SHADER_OPCODE_URB_WRITE_LOGICAL,
reg_undef, srcs, ARRAY_SIZE(srcs));
inst->mlen = 3 + length;
inst->offset = 0;
}
}
static void
emit_urb_indirect_writes_mod(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &src, const fs_reg &offset_src,
fs_reg urb_handle, unsigned mod)
{
assert(nir_src_bit_size(instr->src[0]) == 32);
const unsigned comps = nir_src_num_components(instr->src[0]);
assert(comps <= 4);
const unsigned base_in_dwords = nir_intrinsic_base(instr) +
component_from_intrinsic(instr);
const unsigned comp_shift = mod;
const unsigned mask = nir_intrinsic_write_mask(instr) << comp_shift;
emit_urb_indirect_vec4_write(bld, offset_src, base_in_dwords, src,
urb_handle, comp_shift, comps, mask);
}
static void
emit_urb_indirect_writes(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &src, const fs_reg &offset_src,
fs_reg urb_handle)
{
assert(nir_src_bit_size(instr->src[0]) == 32);
const unsigned comps = nir_src_num_components(instr->src[0]);
assert(comps <= 4);
const unsigned base_in_dwords = nir_intrinsic_base(instr) +
component_from_intrinsic(instr);
/* Use URB write message that allow different offsets per-slot. The offset
* is in units of vec4s (128 bits), so we use a write for each component,
* replicating it in the sources and applying the appropriate mask based on
* the dword offset.
*/
for (unsigned c = 0; c < comps; c++) {
if (((1 << c) & nir_intrinsic_write_mask(instr)) == 0)
continue;
fs_reg src_comp = offset(src, bld, c);
for (unsigned q = 0; q < bld.dispatch_width() / 8; q++) {
fs_builder bld8 = bld.group(8, q);
/* offset is always positive, so signedness doesn't matter */
assert(offset_src.type == BRW_REGISTER_TYPE_D ||
offset_src.type == BRW_REGISTER_TYPE_UD);
fs_reg off = bld8.vgrf(offset_src.type, 1);
bld8.MOV(off, quarter(offset_src, q));
bld8.ADD(off, off, brw_imm_ud(c + base_in_dwords));
fs_reg mask = bld8.vgrf(BRW_REGISTER_TYPE_UD, 1);
bld8.AND(mask, off, brw_imm_ud(0x3));
fs_reg one = bld8.vgrf(BRW_REGISTER_TYPE_UD, 1);
bld8.MOV(one, brw_imm_ud(1));
bld8.SHL(mask, one, mask);
bld8.SHL(mask, mask, brw_imm_ud(16));
bld8.SHR(off, off, brw_imm_ud(2));
fs_reg payload_srcs[4];
unsigned length = 0;
for (unsigned j = 0; j < 4; j++)
payload_srcs[length++] = quarter(src_comp, q);
fs_reg srcs[URB_LOGICAL_NUM_SRCS];
srcs[URB_LOGICAL_SRC_HANDLE] = urb_handle;
srcs[URB_LOGICAL_SRC_PER_SLOT_OFFSETS] = off;
srcs[URB_LOGICAL_SRC_CHANNEL_MASK] = mask;
srcs[URB_LOGICAL_SRC_DATA] = fs_reg(VGRF, bld.shader->alloc.allocate(length),
BRW_REGISTER_TYPE_F);
bld8.LOAD_PAYLOAD(srcs[URB_LOGICAL_SRC_DATA], payload_srcs, length, 0);
fs_inst *inst = bld8.emit(SHADER_OPCODE_URB_WRITE_LOGICAL,
reg_undef, srcs, ARRAY_SIZE(srcs));
inst->mlen = 3 + length;
inst->offset = 0;
}
}
}
static void
emit_urb_direct_reads(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &dest, fs_reg urb_handle)
{
assert(instr->def.bit_size == 32);
unsigned comps = instr->def.num_components;
if (comps == 0)
return;
nir_src *offset_nir_src = nir_get_io_offset_src(instr);
assert(nir_src_is_const(*offset_nir_src));
const unsigned offset_in_dwords = nir_intrinsic_base(instr) +
nir_src_as_uint(*offset_nir_src) +
component_from_intrinsic(instr);
unsigned urb_global_offset = offset_in_dwords / 4;
adjust_handle_and_offset(bld, urb_handle, urb_global_offset);
const unsigned comp_offset = offset_in_dwords % 4;
const unsigned num_regs = comp_offset + comps;
fs_builder ubld8 = bld.group(8, 0).exec_all();
fs_reg data = ubld8.vgrf(BRW_REGISTER_TYPE_UD, num_regs);
fs_reg srcs[URB_LOGICAL_NUM_SRCS];
srcs[URB_LOGICAL_SRC_HANDLE] = urb_handle;
fs_inst *inst = ubld8.emit(SHADER_OPCODE_URB_READ_LOGICAL, data,
srcs, ARRAY_SIZE(srcs));
inst->mlen = 1;
inst->offset = urb_global_offset;
assert(inst->offset < 2048);
inst->size_written = num_regs * REG_SIZE;
for (unsigned c = 0; c < comps; c++) {
fs_reg dest_comp = offset(dest, bld, c);
fs_reg data_comp = horiz_stride(offset(data, ubld8, comp_offset + c), 0);
bld.MOV(retype(dest_comp, BRW_REGISTER_TYPE_UD), data_comp);
}
}
static void
emit_urb_indirect_reads(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &dest, const fs_reg &offset_src, fs_reg urb_handle)
{
assert(instr->def.bit_size == 32);
unsigned comps = instr->def.num_components;
if (comps == 0)
return;
fs_reg seq_ud;
{
fs_builder ubld8 = bld.group(8, 0).exec_all();
seq_ud = ubld8.vgrf(BRW_REGISTER_TYPE_UD, 1);
fs_reg seq_uw = ubld8.vgrf(BRW_REGISTER_TYPE_UW, 1);
ubld8.MOV(seq_uw, fs_reg(brw_imm_v(0x76543210)));
ubld8.MOV(seq_ud, seq_uw);
ubld8.SHL(seq_ud, seq_ud, brw_imm_ud(2));
}
const unsigned base_in_dwords = nir_intrinsic_base(instr) +
component_from_intrinsic(instr);
for (unsigned c = 0; c < comps; c++) {
for (unsigned q = 0; q < bld.dispatch_width() / 8; q++) {
fs_builder bld8 = bld.group(8, q);
/* offset is always positive, so signedness doesn't matter */
assert(offset_src.type == BRW_REGISTER_TYPE_D ||
offset_src.type == BRW_REGISTER_TYPE_UD);
fs_reg off = bld8.vgrf(offset_src.type, 1);
bld8.MOV(off, quarter(offset_src, q));
bld8.ADD(off, off, brw_imm_ud(base_in_dwords + c));
STATIC_ASSERT(IS_POT(REG_SIZE) && REG_SIZE > 1);
fs_reg comp = bld8.vgrf(BRW_REGISTER_TYPE_UD, 1);
bld8.AND(comp, off, brw_imm_ud(0x3));
bld8.SHL(comp, comp, brw_imm_ud(ffs(REG_SIZE) - 1));
bld8.ADD(comp, comp, seq_ud);
bld8.SHR(off, off, brw_imm_ud(2));
fs_reg srcs[URB_LOGICAL_NUM_SRCS];
srcs[URB_LOGICAL_SRC_HANDLE] = urb_handle;
srcs[URB_LOGICAL_SRC_PER_SLOT_OFFSETS] = off;
fs_reg data = bld8.vgrf(BRW_REGISTER_TYPE_UD, 4);
fs_inst *inst = bld8.emit(SHADER_OPCODE_URB_READ_LOGICAL,
data, srcs, ARRAY_SIZE(srcs));
inst->mlen = 2;
inst->offset = 0;
inst->size_written = 4 * REG_SIZE;
fs_reg dest_comp = offset(dest, bld, c);
bld8.emit(SHADER_OPCODE_MOV_INDIRECT,
retype(quarter(dest_comp, q), BRW_REGISTER_TYPE_UD),
data,
comp,
brw_imm_ud(4 * REG_SIZE));
}
}
}
void
fs_visitor::emit_task_mesh_store(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &urb_handle)
{
fs_reg src = get_nir_src(instr->src[0]);
nir_src *offset_nir_src = nir_get_io_offset_src(instr);
if (nir_src_is_const(*offset_nir_src)) {
emit_urb_direct_writes(bld, instr, src, urb_handle);
} else {
bool use_mod = false;
unsigned mod;
/* Try to calculate the value of (offset + base) % 4. If we can do
* this, then we can do indirect writes using only 1 URB write.
*/
use_mod = nir_mod_analysis(nir_get_scalar(offset_nir_src->ssa, 0), nir_type_uint, 4, &mod);
if (use_mod) {
mod += nir_intrinsic_base(instr) + component_from_intrinsic(instr);
mod %= 4;
}
if (use_mod) {
emit_urb_indirect_writes_mod(bld, instr, src, get_nir_src(*offset_nir_src), urb_handle, mod);
} else {
emit_urb_indirect_writes(bld, instr, src, get_nir_src(*offset_nir_src), urb_handle);
}
}
}
void
fs_visitor::emit_task_mesh_load(const fs_builder &bld, nir_intrinsic_instr *instr,
const fs_reg &urb_handle)
{
fs_reg dest = get_nir_def(instr->def);
nir_src *offset_nir_src = nir_get_io_offset_src(instr);
/* TODO(mesh): for per_vertex and per_primitive, if we could keep around
* the non-array-index offset, we could use to decide if we can perform
* a single large aligned read instead one per component.
*/
if (nir_src_is_const(*offset_nir_src))
emit_urb_direct_reads(bld, instr, dest, urb_handle);
else
emit_urb_indirect_reads(bld, instr, dest, get_nir_src(*offset_nir_src), urb_handle);
}
void
fs_visitor::nir_emit_task_intrinsic(const fs_builder &bld,
nir_intrinsic_instr *instr)
{
assert(stage == MESA_SHADER_TASK);
const task_mesh_thread_payload &payload = task_mesh_payload();
switch (instr->intrinsic) {
case nir_intrinsic_store_output:
case nir_intrinsic_store_task_payload:
emit_task_mesh_store(bld, instr, payload.urb_output);
break;
case nir_intrinsic_load_output:
case nir_intrinsic_load_task_payload:
emit_task_mesh_load(bld, instr, payload.urb_output);
break;
default:
nir_emit_task_mesh_intrinsic(bld, instr);
break;
}
}
void
fs_visitor::nir_emit_mesh_intrinsic(const fs_builder &bld,
nir_intrinsic_instr *instr)
{
assert(stage == MESA_SHADER_MESH);
const task_mesh_thread_payload &payload = task_mesh_payload();
switch (instr->intrinsic) {
case nir_intrinsic_store_per_primitive_output:
case nir_intrinsic_store_per_vertex_output:
case nir_intrinsic_store_output:
emit_task_mesh_store(bld, instr, payload.urb_output);
break;
case nir_intrinsic_load_per_vertex_output:
case nir_intrinsic_load_per_primitive_output:
case nir_intrinsic_load_output:
emit_task_mesh_load(bld, instr, payload.urb_output);
break;
case nir_intrinsic_load_task_payload:
emit_task_mesh_load(bld, instr, payload.task_urb_input);
break;
default:
nir_emit_task_mesh_intrinsic(bld, instr);
break;
}
}
void
fs_visitor::nir_emit_task_mesh_intrinsic(const fs_builder &bld,
nir_intrinsic_instr *instr)
{
assert(stage == MESA_SHADER_MESH || stage == MESA_SHADER_TASK);
const task_mesh_thread_payload &payload = task_mesh_payload();
fs_reg dest;
if (nir_intrinsic_infos[instr->intrinsic].has_dest)
dest = get_nir_def(instr->def);
switch (instr->intrinsic) {
case nir_intrinsic_load_mesh_inline_data_intel: {
fs_reg data = offset(payload.inline_parameter, 1, nir_intrinsic_align_offset(instr));
bld.MOV(dest, retype(data, dest.type));
break;
}
case nir_intrinsic_load_draw_id:
dest = retype(dest, BRW_REGISTER_TYPE_UD);
bld.MOV(dest, payload.extended_parameter_0);
break;
case nir_intrinsic_load_local_invocation_id:
unreachable("local invocation id should have been lowered earlier");
break;
case nir_intrinsic_load_local_invocation_index:
dest = retype(dest, BRW_REGISTER_TYPE_UD);
bld.MOV(dest, payload.local_index);
break;
case nir_intrinsic_load_num_workgroups:
dest = retype(dest, BRW_REGISTER_TYPE_UD);
bld.MOV(offset(dest, bld, 0), brw_uw1_grf(0, 13)); /* g0.6 >> 16 */
bld.MOV(offset(dest, bld, 1), brw_uw1_grf(0, 8)); /* g0.4 & 0xffff */
bld.MOV(offset(dest, bld, 2), brw_uw1_grf(0, 9)); /* g0.4 >> 16 */
break;
case nir_intrinsic_load_workgroup_index:
dest = retype(dest, BRW_REGISTER_TYPE_UD);
bld.MOV(dest, retype(brw_vec1_grf(0, 1), BRW_REGISTER_TYPE_UD));
break;
default:
nir_emit_cs_intrinsic(bld, instr);
break;
}
}
Reference in New Issue View Git Blame Copy Permalink