Add a lowering pass extracting all control barriers embedded in scoped
barriers into proper control barriers so we can get rid of the logic
inserting control barriers when an SpvOpControlBarrier with WorkGroup
scope is parsed in spirv_to_nir().
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4900>
SPIRV OpControlBarrier can have both a memory and a control barrier
which some hardware can handle with a single instruction. Let's
turn the scoped_memory_barrier into a scoped barrier which can embed
both barrier types. Note that control-only or memory-only barriers can
be supported through this new intrinsic by passing NIR_SCOPE_NONE to the
unused barrier type.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Suggested-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4900>
This allows us to do API specific checks before removing variable
without filling nir_remove_dead_variables() with API specific code.
In the following patches we will use this to support the removal
of dead uniforms in GLSL.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4797>
This uses a meson builtin to handle -fvisibility=hidden. This is nice
because we don't need to track which languages are used, if C++ is
suddenly added meson just does the right thing.
Acked-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Eric Engestrom <eric@engestrom.ch>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4740>
We should be passing _gt instead of 1 to GEN12_FEATURES or else all TGL
SKUs will be considered as gt1 only.
Fixes: 54996ad492 ("intel/dev: Split .num_subslices out of GEN12_FEATURES macro")
Signed-off-by: Satyeshwar Singh <satyeshwar.singh@intel.com>
Reviewed-by: Eric Engestrom <eric@engestrom.ch>
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5261>
Using HALT to immediately jump to the end of the shader is required to
implement GL_EXT_gpu_shader4 and OpenGL 3.0. However, vanilla OpenGL
1.2 doesn't forbid it and it likely makes something somewhere faster.
We should be consistent and implement the same discard behavior on all
hardware if we can.
The rules for HALT on Gen4-5 are a bit different from Gen6+. On the
older hardware, there is no stack for HALT; instead it's up to software
to save and restore mask registers. However, there's no real saving
needed since we only use HALT to jump to the end of the program where
we're about about to do our FB writes. All we need to do is reset AMask
to DMask, the value it was initialized to at the start of the thread.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5244>
This extension only allows HLSL shader compilers to optionally embed
unambiguous type information which can be safely ignored by the driver.
This fixes a crash with the recent Vulkan backend of Path Of Exile
(it uses the extension without checking if it's supported).
Cc: <mesa-stable@lists.freedesktop.org>
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Tested-by: Edmondo Tommasina <edmondo.tommasina@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5237>
This will make the GL drivers pick the right SIMD variant for a given
group size set during dispatch. The heuristic implemented in
brw_cs_simd_size_for_group_size() is the same as in brw_compile_cs().
The cs_prog_data::simd_size field was removed. The generated SIMD
sizes are marked in a bitmask, which is already used via
brw_cs_simd_size_for_group_size() by the drivers.
When in variable group size, it is OK if larger SIMD shader spill,
since we'd need it for the cases where the smaller one can't hold all
the invocations.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5142>
Move the decision one level up, let brw_compile_*() functions use the
spilling information to decide whether or not a certain width
compilation can spill (passed via run_*() functions).
The min_dispatch_width was used to compare with the dispatch_width and
decide whether "a previous shader is already available, so don't
accept spill".
This is replaced by:
- Not calling run_*() functions if it is know beforehand a smaller width
already spilled -- since the larger width will spill and fail;
- Explicitly passing whether or not a shader is allowed to spill. For
the cases where the smaller width is available and haven't spilled,
the larger width will be compiled but is only useful if it won't
spill.
Moving the decision to this level will be useful later for variable
group size, which is a case where we want all the widths to be allowed
to spill.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5142>
It was found that dual-source blending hangs with SIMD16 dispatch in some
specific but unknown situation. Which in the wild happen when rgba
anti-aliasing is enabled for fonts.
Cc: <mesa-stable@lists.freedesktop.org>
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/2183
Signed-off-by: Danylo Piliaiev <danylo.piliaiev@globallogic.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5037>
After 2663759af0 ("intel/fs: Add and use a new load_simd_width_intel
intrinsic") the local_workgroup_size is not used anymore except for
assertions at the pass' start, so drop it from state struct.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5213>
The nir_intrinsic_load_simd_width_intel is always lowered by the
brw_nir_lower_simd() pass before the emission happens. This is likely
a "leftover" from patch rewriting/squashing that happened when this
intrinsic was added.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5213>
When running in a virtual context, the timestamp register is unreadable
on Gen12+.
While we could work around this, that would result in very inaccurate
results for an extension where the whole point is accuracy, so let's
just disable the extension.
Signed-off-by: Eric Engestrom <eric.engestrom@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2797>
This allows its use without the need for an anv_device.
Signed-off-by: Eric Engestrom <eric.engestrom@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2797>
Issue description from Matt's commit e7c376ad:
"var_range_end(v, n) loops over the n components of variable number v and
finds the maximum value, giving the last use of any component of v.
Therefore it expects v to correspond to the variable associated with the
.x channel of the VGRF.
var_from_reg() however returns the variable for the first channel of the
VGRF, post-swizzle.
So, if the last register had a swizzle with y, z, or w in the swizzle
component, we would read out of bounds. For any other register, we would
read liveness information from the next register.
The fix is to convert the src_reg to a dst_reg in order to call the
dst_reg version of var_from_reg() that doesn't consider the swizzle."
Closes: #3003
Fixes: 48dfb30f ('intel/compiler: Move all live interval analysis results into vec4_live_variables')
Reviewed-by: Matt Turner <mattst88@gmail.com>
Signed-off-by: Andrii Simiklit <asimiklit.work@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4941>
An unknown issue is causing vs push constants to become corrupted
during object-level preemption. For now, restrict to command
buffer level preemption to avoid rendering corruption.
Cc: <mesa-stable@lists.freedesktop.org>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5110>
A buffer added to all execbufs so that we can attribute a batch that
caused a hang to a particular driver.
v2: Reuse workaround BO
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3203>
We initially used this debug option to mean "don't bother registering
the OA configuration into the kernel".
This change makes this option suppress any interaction with the
i915/perf interface. This is useful when debugging self modifying
batches with performance queries while running on the intel_mi_runner.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
This has the same kernel requirements are VK_INTEL_performance_query
v2: Fix empty queue submit (Lionel)
v3: Fix autotool build issue (Piotr Byszewski)
v4: Fix Reset & Begin/End in same command buffer, using soft-pin &
relocation on the same buffer won't work currently. This version
uses a somewhat dirty trick in anv_execbuf_add_bo (Piotr Byszewski)
v5: Fix enumeration with null pointers for either pCounters or
pCounterDescriptions (Piotr)
Fix return condition on enumeration (Lionel)
Set counter uuid using sha1 hashes (Lionel)
v6: Fix counters scope, should be COMMAND_KHR not COMMAND_BUFFER_KHR (Lionel)
v7: Rebase (Lionel)
v8: Rework checking for loaded queries (Lionel)
v9: Use new i915-perf interface
v10: Use anv_multialloc (Jason)
v11: Implement perf query passes using self modifying batches (Lionel)
Limit support to softpin/gen8
v12: Remove spurious changes (Jason)
v13: Drop relocs (Jason)
v14: Avoid overwritting .sType in
VkPerformanceCounterKHR/VkPerformanceCounterDescriptionKHR (Lionel)
v15: Don't copy the entire
VkPerformanceCounterKHR/VkPerformanceCounterDescriptionKHR (Jason)
Reuse anv_batch rather than custom packing (Jason)
v16: Fix missing MI_BB_END in reconfiguration batch
Only report the extension with kernel support (perf_version >= 3)
v17: Some cleanup of unused stuff
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
The current code relies on the order of the function
gen_perf_query_result_accumulate() to match the descriptions written
by gen_perf.py. Let's just reuse the offset specified in the python
script.
v2: Use accumlator offsets more (Jason)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
We're about to use the offset fields from the query object. We can't
just use a made up object.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
The produced array tells use what metric to enable for a given pass.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
We want to compute the number of passes required to gather performance
data about a set of counters.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
For a future extension we want to be able to list the counters. Our
existing sets counters might contain the same counters multiple times.
This is a side effect of the fixed OA counters in the HW. We track
thoses with a mask so that we know when a counter is available from
multiple metrics.
v2: Use BITFIELD64_BIT() (Jason)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
Previously counter descriptions as well register values were written
in global static variables. This isn't really thread safe so instead
ralloc all the data back under the gen_perf_config object.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
On Vulkan most of those are already covered by standard queries so
add the ability to skip them.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
This new bit invalidates the cache/prefetch of commands in the command
streamer. This will be useful for self modifying batches.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
This change adds a call/return execution mode for secondary command
buffer rather than the existing copy into the primary batch mode.
v2: Rework convention to avoid burning an ALU register (Jason)
v3: Use anv_address_add() (Jason)
v4: Move command emissions to anv_batch_chain.c (Jason)
v5: Also move last MI_BBS emission in secondary command buffer to
anv_batch_chain.c (Jason)
v6: Fix end secondary command buffer end (Jason)
v7: Refactor anv_batch_address() to remove additional emit functions
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2775>
