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

intel/brw: Simplify enabling brw_fs_test_dispatch_packing

Browse Source 
Signed-off-by: Jordan Justen <jordan.l.justen@intel.com>
Reviewed-by: Francisco Jerez <currojerez@riseup.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29273>
This commit is contained in:
Jordan Justen
2024-05-17 01:20:33 -07:00
committed by Marge Bot
parent aa152ef431
commit 739613ec70
3 changed files with 40 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/intel/compiler/brw_compiler.h
+3 -1
View File
@@ -1641,8 +1641,10 @@ brw_stage_has_packed_dispatch(ASSERTED const struct intel_device_info *devinfo,
/* The code below makes assumptions about the hardware's thread dispatch
* behavior that could be proven wrong in future generations -- Make sure
* to do a full test run with brw_fs_test_dispatch_packing() hooked up to
* the NIR front-end before changing this assertion.
* the NIR front-end before changing this assertion. It can be temporarily
* enabled by setting the macro below to true.
*/
#define ENABLE_FS_TEST_DISPATCH_PACKING false
assert(devinfo->ver <= 12);
switch (stage) {
src/intel/compiler/brw_fs.cpp
-34
View File
@@ -4499,40 +4499,6 @@ brw_compile_bs(const struct brw_compiler *compiler,
return g.get_assembly();
}
/**
* Test the dispatch mask packing assumptions of
* brw_stage_has_packed_dispatch(). Call this from e.g. the top of
* fs_visitor::emit_nir_code() to cause a GPU hang if any shader invocation is
* executed with an unexpected dispatch mask.
*/
static UNUSED void
brw_fs_test_dispatch_packing(const fs_builder &bld)
{
const fs_visitor *shader = bld.shader;
const gl_shader_stage stage = shader->stage;
const bool uses_vmask =
stage == MESA_SHADER_FRAGMENT &&
brw_wm_prog_data(shader->prog_data)->uses_vmask;
if (brw_stage_has_packed_dispatch(shader->devinfo, stage,
shader->max_polygons,
shader->prog_data)) {
const fs_builder ubld = bld.exec_all().group(1, 0);
const fs_reg tmp = component(bld.vgrf(BRW_TYPE_UD), 0);
const fs_reg mask = uses_vmask ? brw_vmask_reg() : brw_dmask_reg();
ubld.ADD(tmp, mask, brw_imm_ud(1));
ubld.AND(tmp, mask, tmp);
/* This will loop forever if the dispatch mask doesn't have the expected
* form '2^n-1', in which case tmp will be non-zero.
*/
bld.emit(BRW_OPCODE_DO);
bld.CMP(bld.null_reg_ud(), tmp, brw_imm_ud(0), BRW_CONDITIONAL_NZ);
set_predicate(BRW_PREDICATE_NORMAL, bld.emit(BRW_OPCODE_WHILE));
}
}
unsigned
fs_visitor::workgroup_size() const
{
src/intel/compiler/brw_fs_nir.cpp
+37
View File
@@ -8446,6 +8446,40 @@ emit_shader_float_controls_execution_mode(nir_to_brw_state &ntb)
brw_imm_d(mode), brw_imm_d(mask));
}
/**
* Test the dispatch mask packing assumptions of
* brw_stage_has_packed_dispatch(). Call this from e.g. the top of
* nir_to_brw() to cause a GPU hang if any shader invocation is
* executed with an unexpected dispatch mask.
*/
static UNUSED void
brw_fs_test_dispatch_packing(const fs_builder &bld)
{
const fs_visitor *shader = bld.shader;
const gl_shader_stage stage = shader->stage;
const bool uses_vmask =
stage == MESA_SHADER_FRAGMENT &&
brw_wm_prog_data(shader->prog_data)->uses_vmask;
if (brw_stage_has_packed_dispatch(shader->devinfo, stage,
shader->max_polygons,
shader->prog_data)) {
const fs_builder ubld = bld.exec_all().group(1, 0);
const fs_reg tmp = component(bld.vgrf(BRW_TYPE_UD), 0);
const fs_reg mask = uses_vmask ? brw_vmask_reg() : brw_dmask_reg();
ubld.ADD(tmp, mask, brw_imm_ud(1));
ubld.AND(tmp, mask, tmp);
/* This will loop forever if the dispatch mask doesn't have the expected
* form '2^n-1', in which case tmp will be non-zero.
*/
bld.emit(BRW_OPCODE_DO);
bld.CMP(bld.null_reg_ud(), tmp, brw_imm_ud(0), BRW_CONDITIONAL_NZ);
set_predicate(BRW_PREDICATE_NORMAL, bld.emit(BRW_OPCODE_WHILE));
}
}
void
nir_to_brw(fs_visitor *s)
{
@@ -8457,6 +8491,9 @@ nir_to_brw(fs_visitor *s)
.bld = fs_builder(s).at_end(),
};
if (ENABLE_FS_TEST_DISPATCH_PACKING)
brw_fs_test_dispatch_packing(ntb.bld);
for (unsigned i = 0; i < s->nir->printf_info_count; i++) {
brw_stage_prog_data_add_printf(s->prog_data,
s->mem_ctx,