nvc0: fix some compute texture validation bits on kepler

(a) Make sure to update the TIC in case of an updated buffer address
(b) Mark newly-inactive textures dirty so that we update the handle in
set_tex_handles.

Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>

mesa/xfb: report calculated size for XFB buffer objects.

This fixes:
GL45-CTS.direct_state_access.xfb_buffers

This test looks correct to me, we should work out the
size value and report it rather than using only the size
from the Range interface.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Signed-off-by: Dave Airlie <airlied@redhat.com>

swr: automake: silence the python invocation

Cc: Tim Rowley <timothy.o.rowley@intel.com>
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

swr: automake: attempt to fix the out-of-tree build

Make sure that the output folder is created otherwise the python scripts
yells at us.

Cc: 0xe2.0x9a.0x9b@gmail.com
Cc: Tim Rowley <timothy.o.rowley@intel.com>
Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=96238
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

swr: remove LLVM dependency from source generation rules.

The dependencies should not mention any files external to the project.
If we want to do sanity checks for the LLVM installed on the system we
should do that in configure, yet again where is the merit which header
gets checked and which doesn't ?

Cc: Tim Rowley <timothy.o.rowley@intel.com>
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

swr: add all the generators to the release tarball.

Namely the python scripts and the knobs.template.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

anv: automake: don't forget to cleanup dev_icd.json

Otherwise `make distcheck' will barf at us as the file is dangling.

Ideally this should be part of the clean-local hook, although we include
install-lib-links.mk which already has one.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

anv: automake: bring back VULKAN_ENTRYPOINT_CPPFLAGS

We should not have removed them in the first place. There's a subtle
difference between generating the complete sources and using them which
was not obvious as we nuked them.

Without this, the release tarball ends up without various hunks of the
generated sources, thus things fail at a later stage as we attempt to
build them.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

anv: automake: ship the json files in the release tarball

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

softpipe: add sp_buffer.h to the sources list (release tarball)

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

freedreno: make sure we pick up ir3_nir_trig.py in the release tarball

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

isl: add isl_priv.h to the sources list

Otherwise it will be missing from the release tarball.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

isl: move the sources lists to Makefile.sources

[Emil Velikov: use the file in the autoconf build]
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

isl: automake: list builddir before srcdir in the includes list

As seen elsewhere - we want to include the freshly built sources as
opposed the the (likely) stale ones in the srcdir.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

isl: automake: flatten the tests rules

Fold the unneeded extra variable tests_ldadd, the explicit sources
section (single file with the default extension) and flip the
check_PROGRAMS <> TESTS order (TESTS includes scripts, while
check_PROGRAMS is binaries only).

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

isl: automake: remove unneeded install-lib-links.mk include

One uses the makefile to create compatibility symlinks (to
$top_builddir/libs) for shared libraries/modules. As we don't create any
here, there's no need to include the file.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

isl: automake: remove unneeded SUBDIRS

As we do not include any other subdirs but self, we don't need to set
it.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

genxml: move the sources (headers) list to Makefile.sources

[Emil Velikov: use the file in the autoconf build]
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>

anv: bail out if anv_wsi_init() fails

Otherwise we'll end up setting up a device with no winsys integration.

Signed-off-by: Emil Velikov <emil.velikov@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
---
Hard-coding the rendernode name in anv_physical_device_init() is a bad
idea really. We could/should be using drmGetDevices() to get info on all
the devices (master/render/etc. node names, pci location etc.) and apply
our heuristics on top of that.

That can come up as a follow up change.

anv: resolve wayland-only build

Ensure that the final X11/XCB hunk is guarded by the correct macro.
Otherwise we'll require the symbol even when building without said
platform.

Cc: Cedric Sodhi <manday@openmail.cc>
Reported-by: Cedric Sodhi <manday@openmail.cc>
Signed-off-by: Emil Velikov <emil.l.velikov@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

anv: Fix use of uninitialized variable.

The return variable was not set for failure paths.
It has now been changed to VK_ERROR_INITIALIZATION_FAILED
for failure paths.

Coverity: 1358944
Reviewed-by: Eric Engestrom <eric.engestrom@imgtec.com>
Signed-off-by: Robert Foss <robert.foss@collabora.com>
Reviewed-by: Emil Velikov <emil.velikov@collabora.com>
[Emil Velikov: rebase against master, s/vulkan/anv/]
Signed-off-by: Emil Velikov <emil.velikov@collabora.com>

gallium: push offset down to driver

Push offset down to drivers when importing dmabuf. This is needed
to more fully support EGL_EXT_image_dma_buf_import when a non-zero
offset is specified.

Tesing has been done for freedreno, and compile tested following
gallium drivers:
nouveau,svga,virgl,r600,r300,radeonsi,swrast,i915,ilo

Signed-off-by: Stanimir Varbanov <stanimir.varbanov@linaro.org>
Reviewed-by: Emil Velikov <emil.l.velikov@gmail.com>

st/dri: cleanup image_from_fd/dma_buf paths

Signed-off-by: Stanimir Varbanov <stanimir.varbanov@linaro.org>
Reviewed-by: Emil Velikov <emil.l.velikov@gmail.com>

st/dri: add handling of R8 and GR88 DRI fourcc formats

This helps to import dmabuf buffers from DRM_FORMAT_R8 and
DRM_FORMAT_GR88 used for example by GStreamer for YUV to RGB
conversion using shaders.

Signed-off-by: Stanimir Varbanov <stanimir.varbanov@linaro.org>
Reviewed-by: Emil Velikov <emil.l.velikov@gmail.com>

radeonsi: Don't offset OFFCHIP_BUFFERING on pre-VI cards.

Signed-off-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=96239
Reviewed-by: Marek Olšák <marek.olsak@amd.com>

i965: Expose GL 4.3 on Gen8+.

ARB_compute_shader was the last feature missing.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Skip gen4 pre/post-send dependency workaronds for the first/last block.

We know that there cannot be any destination dependency race if we
reach the beginning or end of the program without having found any
other instruction the send could possibly race with.  This avoids
emitting a pile of useless moves at the beginning or end of the
program in the most common case in which the program has a single
basic block only.

On the original i965 I get the following shader-db results:

 total instructions in shared programs: 3354165 -> 3215637 (-4.13%)
 instructions in affected programs: 3183065 -> 3044537 (-4.35%)
 helped: 13498
 HURT: 0

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Skip SIMD lowering source unzipping for regular scalar regions.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Factor out region zipping and unzipping from the SIMD lowering pass.

Just to make sure we keep the SIMD lowering pass tidy when we
introduce additional logic to try to optimize out the copy
instructions used to zip and unzip the destination and source regions
into multiple packed regions of the lowered instruction width.
Shouldn't cause any functional changes.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Generalize regions_overlap() from copy propagation to handle non-VGRF files.

This will be useful in several places.  The only externally visible
difference (other than non-VGRF files being supported now) is that the
region sizes are now passed in byte units instead of in GRF units
because the loss of precision would have become a problem in the SIMD
lowering pass.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Refactor offset() into a separate function taking the width as argument.

This will be useful in the SIMD lowering pass to avoid having to
construct a builder object of the known region width just to pass it
as argument to offset(), which doesn't do anything with it other than
taking the builder dispatch_width as region width.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Implement opt_sampler_eot() in terms of logical sends.

This makes the whole LOAD_PAYLOAD munging unnecessary which simplifies
the code and will allow the optimization to succeed in more cases
independent of whether the LOAD_PAYLOAD instruction can be found or
not.

The following patch is squashed in:

SQUASH: i965/fs: Add basic dataflow check to opt_sampler_eot().

The sampler EOT optimization pass naively assumes that the texturing
instruction provides all the data used by the FB write just because
they're standing next to each other.  The least we should be checking
is whether the source and destination regions of the FB write and
texturing instructions match.  Without this the previous seemingly
harmless patch would have caused opt_sampler_eot() to misoptimize a
shader from dota-2 causing DCE to eliminate all of its 78 instructions
except for the final sampler EOT message (!).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Fix UB list sentinel dereference in opt_sampler_eot().

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Take opt_redundant_discard_jumps out of the optimization loop.

No shader-db regressions.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Run SIMD and logical send lowering after the optimization loop.

There are two reasons why this is useful:

 - It avoids the introduction of an amount of partial writes emitted
   by the SIMD lowering pass to zip and unzip register regions early
   during optimization, which can make subsequent optimization less
   effective.

 - It substantially reduces the burden on the compiler when a large
   fraction of the instructions in the program need to be split (e.g.
   during SIMD32 builds).  Individual halves of split instructions
   will be optimized identically (if they can still be optimized at
   all), so doing it up front can duplicate the amount of instructions
   the optimizer has to deal with which causes the compilation time to
   explode in some cases due to the worse-than-linear runtime
   behaviour of the back-end.

It seems helpful to re-run a few optimization passes in cases where
any of the lowering passes was able to make progress.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Add FS_OPCODE_FB_WRITE_LOGICAL to has_side_effects().

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Allow constant propagation into logical send sources.

Logical sends are eventually lowered into a series of copies so they
can take almost anything as source.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Let CSE handle logical sampler sends as expressions.

This will prevent some shader-db regressions when we start plumbing
logical sends through the optimizer.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Pass a BAD_FILE register to the logical FB write when oMask is unused.

This will let the optimizer know that the sample mask value is unused
so its definition can be DCE'ed.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

glsl: fix xfb_offset unsized array validation

This partially fixes CTS test:
GL44-CTS.enhanced_layouts.xfb_get_program_resource_api

The test now fails at a tes evaluation shader with unsized output arrays.

The ARB_enhanced_layouts spec says:

   "It is a compile-time error to apply xfb_offset to the declaration of an
   unsized array."

So this seems like a bug in the CTS.

Reviewed-by: Dave Airlie <airlied@redhat.com>

glsl: dont crash when attempting to assign a value to a builtin define

For example GL_ARB_enhanced_layouts = 3;

Fixes:
GL44-CTS.enhanced_layouts.glsl_contant_immutablity

Reviewed-by: Dave Airlie <airlied@redhat.com>

egl/dri3: don't crash on no context.

Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=94925

Pointed out by Karol Herbst on irc.

Signed-off-by: Dave Airlie <airlied@redhat.com>
Cc: "11.1 11.2" <mesa-stable@lists.freedesktop.org>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>

mesa/program_interface_query: fix transform feedback varyings.

The spec says gl_NextBuffer and gl_SkipComponents need to be
returned to userspace in the program interface queries.

We currently throw those away, this requires a complete piglit
run to make sure no drivers fallover due to the extra varyings.

This fixes:
GL45-CTS.program_interface_query.transform-feedback-built-in

Reviewed-by: Timothy Arceri <timothy.arceri@collabora.com>
Signed-off-by: Dave Airlie <airlied@redhat.com>

glsl/ast: subroutineTypes can't be returned from functions.

These types can't be returned.

This fixes:
GL43-CTS.shader_subroutine.subroutines_not_allowed_as_variables_constructors_and_argument_or_return_types
for the return type case.

Reviewed-by: Chris Forbes <chrisforbes@google.com>
Signed-off-by: Dave Airlie <airlied@redhat.com>

glsl: use has_double() helper

Reviewed-by: Eduardo Lima Mitev <elima@igalia.com>

glsl: fix explicit uniform block alignment

This stops the offset being bumped again when and an explicit
alignment has already been applied.

Fixes alignment issues in:
GL44-CTS.enhanced_layouts.uniform_block_alignment

Note the test still fails due to unrelated issues with doubles.

Reviewed-by: Eduardo Lima Mitev <elima@igalia.com>

i965: Shrink stage_prog_data param array length

It appears we were over-allocating these arrays.

Previously we would use nir->num_uniforms directly for scalar
programs, and multiply it by 4 for vec4 programs.

Instead we should have been dividing by 4 in both cases to convert
from bytes to a gl_constant_value count. The size of gl_constant_value
is 4 bytes.

Signed-off-by: Jordan Justen <jordan.l.justen@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

nv50,nvc0: fix the max_vertices=0 case

This is apparently legal. Drop any emit/restarts, and pass a 1 to the
hardware.

Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>

st/mesa: fix setting of point_size_per_vertex in ES contexts

GL ES 2.0+ does not have a GL_PROGRAM_POINT_SIZE enable, unlike desktop
GL. So we have to go and check the last pre-rasterizer stage to see
whether it outputs a point size or not.

This fixes a number of dEQP tests that use a geometry or tessellation
shader to emit points primitives.

Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Cc: "11.1 11.2" <mesa-stable@lists.freedesktop.org>

mesa: skip level checking for FramebufferTexture*D if texture is zero

From the OpenGL 4.5 core spec:
  "An INVALID_VALUE error is generated if texture is not zero and level is
  not a supported texture level for textarget, as described above."

Other FramebufferTexture functions already do the right thing.

This fixes the main menu in F1 2015.

Cc: 11.1 11.2 <mesa-stable@lists.freedesktop.org>
Reviewed-by: Dave Airlie <airlied@redhat.com>

st/mesa: expose OES_shader_io_blocks when we have enough for ES 3.1

Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>
Reviewed-by: Matt Turner <mattst88@gmail.com>

swr: [rasterizer] Do not define _mm256_storeu2_m128i with icc.

Fix build error with icc.

  CXX      libswrAVX_la-swr_clear.lo
icpc: command line warning #10006: ignoring unknown option '-Wdelete-non-virtual-dtor'
In file included from ./rasterizer/jitter/jit_api.h(31),
                 from swr_context.h(30),
                 from swr_clear.cpp(24):
./rasterizer/common/os.h(135): error: expected an identifier
  void _mm256_storeu2_m128i(__m128i *hi, __m128i *lo, __m256i a)
       ^

Signed-off-by: Vinson Lee <vlee@freedesktop.org>
Reviewed-by: Tim Rowley <timothy.o.rowley@intel.com>

i965: add missing return in if statement

Re-add the "return false" that was removed in 0c02d7002d6c005b4c1fe997b5ef5916978dd183

It seems that something went wrong when merging the patch. The patch
sent to the mailing list does not directly match what was committed.
https://lists.freedesktop.org/archives/mesa-dev/2016-May/118198.html

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

gk110/ir: fix unspilling of predicates from registers

Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=96258
Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>
Cc: "11.2 11.1" <mesa-stable@lists.freedesktop.org>

nvc0: remove outdated surfaces validation code for GK104

This code was used for validating surfaces with compute but now we use
pipe_image_view instead. Anyway, surfaces support should be
re-introduced properly once OpenCL happens.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Ilia Mirkin <imirkin@alum.mit.edu>

nvc0: do not always invalidate 3D CBs when using compute

Constant buffers are aliased between 3D and CP on Fermi, but we should
only invalidate them when a compute shader actually uses CBs and not
all the time after a lauching grid.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Ilia Mirkin <imirkin@alum.mit.edu>

i965: Update compute workgroup size limit calculation for SIMD32.

This should have the side effect of enabling the ARB_compute_shader
extension on Gen8+ hardware and all Gen7 platforms that didn't
previously expose it (VLV and IVB GT1) due to the number of hardware
threads per subslice being insufficient in SIMD16 mode.

v2: Bump workgroup size limit for GLES too (Jordan).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>

i965: Add do32 debug option.

The do32 INTEL_DEBUG option causes the back-end to try to generate a
SIMD32 program when compiling a compute shader regardless of the
specified compute shader workgroup size, which will be useful for
testing SIMD32 code generation in the most common case in which the
workgroup size doesn't exceed the SIMD16 limit so SIMD32 codegen
wouldn't be automatically enabled.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Build 32-wide compute shader when needed.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Extend back-end interface for limiting the shader dispatch width.

This replaces the current fs_visitor::no16() interface with
fs_visitor::limit_dispatch_width(), which takes an additional
parameter allowing the caller to specify the maximum dispatch width a
shader can be compiled with.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Implement SIMD32 register allocation support.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Remove pre-Gen7 register allocation class micro-optimization.

This was trying to save some one-time init on pre-Gen7 hardware under
the assumption that one would only ever need 1, 2, 4 and 8-wide
registers on those platforms.  However nothing guarantees that those
will be the only VGRF sizes used after lowering and optimization.  In
some cases we may end up with a temporary of different size being
allocated (e.g. by SIMD lowering to zip or unzip a multi-component
register region of a logical send instruction), and there is no
guarantee that they will be optimized away before register allocation
(especially since the compute_to_mrf coalescing pass is
rather... lacking...).  Instead just allocate classes for all possible
VGRF sizes up to MAX_VGRF_SIZE to avoid a crash in pq_test() when we
encounter a variable of any other size.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Don't mutate multi-component arguments in sampler payload set-up.

The Gen5+ sampler message payload construction code steps through the
coordinate and derivative components by induction like 'coordinate =
offset(coordinate, bld, 1)', the problem is that while doing that it
may step one past the end of the coordinate vector causing an
assertion failure in offset() if it happens to be a (single component)
immediate.  Right now coordinates and derivatives are typically passed
as actual registers but that will no longer be the case when we start
propagating constants into logical messages.

Instead express coordinate components in closed form like
'offset(coordinate, bld, i)' -- The end result seems slightly more
readable that way and it allows passing the coordinate and derivative
registers by const reference instead of by value, so it seems like a
clean-up in its own right.

v2: Fold a few post-increment operators into the last MOV
    statement. (Jason)

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Fix multiple ACP interference during copy propagation.

This is more fallout from cf375a3333e54a01462f192202d609436e5fbec8.
It's possible for multiple ACP entries to interfere with a given VGRF
write, so we need to continue iterating even if an overlapping entry
has already been found.

Cc: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Fix cmod propagation not to propagate non-identity cmod into CMP(N).

The conditional mod of these instructions determines the semantics of
the comparison itself (rather than being evaluated based on the result
of the instruction as is usually the case for most other instructions
that allow conditional mods), so it's in general not legal to
propagate a conditional mod into a CMP instruction.  This prevents
cmod propagation from (mis)optimizing:

 cmp.z.f0 tmp, ...
 mov.z.f0 null, tmp

into:

 cmp.z.f0 tmp, ...

which gives the negation of the flag result of the original sequence.
I could reproduce this easily with SIMD32 but I don't see any reason
why the problem would be SIMD32-specific, it was most likely working
by luck.

Cc: mesa-stable@lists.freedesktop.org
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Estimate number of registers written correctly in opt_register_renaming.

The current estimate is incorrect for non-32b types.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Add (sub)reg_offset asserts to brw_reg_from_fs_reg.

These are completely ignored by the conversion to brw_reg, so they
better be zero.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Reset reg_offset of the original destination to zero in compute_to_mrf().

Prevents an assertion failure in the following commit.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Skip remove_duplicate_mrf_writes() during SIMD32 runs.

The pass is disabled in SIMD16 dispatch mode for the same reason, it
cannot handle instructions that write multiple MRF registers at once.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Use SIMD8 SSBO GET_BUFFER_SIZE message regardless of the dispatch width.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Don't emit duplicated SSBO GET_BUFFER_SIZE instruction unnecessarily.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Emit fixed width memory fence opcode regardless of the dispatch width.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Return 32 bit mask from fs_builder::sample_mask().

This doesn't actually handle the FS case, just add an assertion for
the moment so I don't forget to update it later on for SIMD32 fragment
shader dispatch.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Emit fixed-width null register regardless of the dispatch width.

brw_null_vec() cannot handle widths over 16 but it doesn't really
matter what width we specify for null registers because destination
regions have no width field at the hardware level.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Fix half() to handle more exotic register files.

horiz_offset() is able to deal with a superset of the register files
currently special-cased in half().  Just call horiz_offset() in all
cases.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Fix horiz_offset() to handle ARF and HW GRF register files.

We'll hit these in some cases during SIMD lowering in 32-wide
programs.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Clean up remaining uses of fs_inst::reads_flag and ::writes_flag.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Keep track of flag dependencies with byte granularity during scheduling.

This prevents false dependencies from being created between
instructions that write disjoint 8-bit portions of the flag register
and OTOH should make sure that the scheduler considers dependencies
between instructions that write or read multiple flag subregisters
at once (e.g. 32-wide predication or conditional mods).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Track flag register liveness with byte granularity.

This is required for correctness in presence of multiple 8-wide flag
writes (e.g. 8-wide instructions with a conditional mod set) which
update a different portion of the same 16-bit flag subregister.  Right
now we keep track of flag dataflow with 16-bit granularity and
consider flag writes to have killed any previous definition of the
same subregister even if the write was less than 16 channels wide,
which can cause live flag register updates to be dead code-eliminated
incorrectly.

Additionally this makes sure that we handle 32-wide flag writes and
reads which may span multiple flag subregisters so the current
approach of just setting/testing a single bit from the live set
wouldn't have worked.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Define methods to calculate the flag subset read or written by an fs_inst.

v2: Codestyle fixes (Jason).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Expose arbitrary channel execution groups to the IR.

This generalizes the current fs_inst::force_sechalf flag to allow
specifying channel enable groups other than 0 or 8.  At some point it
will likely make sense to fix the vec4 generator to support arbitrary
execution groups and then move the definition of fs_inst::group into
backend_instruction (e.g. so we can do FP64 in the VEC4 back-end).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/ir: Make BROADCAST emit an unmasked single-channel move.

Alternatively we could have extended the current semantics to 32-wide
mode by changing brw_broadcast() to emit multiple indexed MOV
instructions in the generator copying the selected value to all
destination registers, but it seemed rather silly to waste EU cycles
unnecessarily copying the exact same value 32 times in the GRF.

The vstride change in the Align16 path is required to avoid assertions
in validate_reg() since the change causes the execution size of the
MOV and SEL instructions to be equal to the source region width.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Allow specifying arbitrary quarter control to FIND_LIVE_CHANNEL.

This makes FIND_LIVE_CHANNEL behave like a normal instruction for
non-zero quarter control.  On Gen8+ we just leave the quarter control
field of the emitted FBL instruction set to the default value so the
hardware applies the expected shift to the execution mask signals.  On
Gen7 we apply the offset manually by specifying a non-zero subregister
offset in the source region of the FBL instruction.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Allow specifying arbitrary execution sizes up to 32 to FIND_LIVE_CHANNEL.

Due to a Gen7-specific hardware bug native 32-wide instructions get
the lower 16 bits of the execution mask applied incorrectly to both
halves of the instruction, so the MOV trick we currently use wouldn't
work.  Instead emit multiple 16-wide MOV instructions in 32-wide mode
in order to cover the whole execution mask.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Lower 32-wide scratch writes in the generator.

The hardware has messages that can write 32 32bit components at once
but the channel enable mask gets messed up.  We need to split them
into several 16-wide scratch writes for the channel enables to be
applied correctly.  The SIMD lowering pass cannot be used for this
because scratch writes are emitted rather late during register
allocation long after SIMD lowering has been done.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Implement scratch reads and writes of 4 GRFs at a time.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Fix Gen7+ DP scratch message size calculation on Gen7.

Gen7 hardware expects the block size field in the message descriptor
to be the number of registers minus one instead of the log2 of the
number of registers.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Set execution size explicitly for memory fence send message.

We don't want to emit a 32-wide send message in 32-wide programs.  The
memory fence message should have the same effect regardless of the
execution size (as long as it's valid) so just set it to one.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Consider QtrCtrl 3Q-4Q in typed surface message descriptor setup.

In SIMD32 programs the compiler is responsible for providing the
appropriate half of the sample mask in the message header, so the
first and third quarters both map to the first slot group of the
provided 16-bit half, while the second and fourth quarters map to the
second slot group -- IOW they should be equivalent to 1Q and 2Q modulo
two.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Clean up remaining uses of dispatch_width in the generator.

Most of these are bugs because the intended execution size of an
instruction and the dispatch width of the shader aren't necessarily
the same (especially in SIMD32 programs).

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Remove brw_codegen::compressed and ::compressed_stack.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Use current exec size instead of p->compressed in surface message generation.

This was kind of an abuse of p->compressed, dataport send message
instructions are always uncompressed.  Use the current execution size
instead since p->compressed is on its way out.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: No need to reset predicate control after emitting some instructions.

Trivial clean-up.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Pass current execution size to brw_IF() and brw_DO().

This gets IF and DO instructions working in SIMD32 programs.  brw_IF()
and brw_DO() should probably behave in the same way as other generator
functions that emit control flow instructions and just figure out the
right execution size by themselves from the current execution controls
specified through the brw_codegen argument.  Changing that will
require updating lots of Gen4-5 clipper code though, so for the moment
just pass the current value redundantly from the FS generator.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Stop using p->compressed to specify the exec size of control flow instructions.

p->compressed won't work for SIMD32, we should just be using the
execution size value specified via p->current instead.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Extend region width calculation to allow arbitrary execution sizes.

Instead of just halving the execution size when the instruction is
compressed hoping that it will give a legal source region width, we
can calculate the maximum legal width value in closed form from the
component size and stride.  This makes sure that brw_reg_from_fs_reg()
always returns a valid hardware region even for virtual 32-wide
instructions (e.g. send-like instructions) that would seem to exceed
the hardware region width limit after halving.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Pass the compression mode to brw_reg_from_fs_reg().

Curro is planning to eliminate p->compressed, so let's avoid using it
here and just pass in the value directly.

Signed-off-by: Kenneth Graunke <kenneth@whitecape.org>
[ Francisco Jerez: Pass boolean flag instead of brw_compression enum. ]
Reviewed-by: Francisco Jerez <currojerez@riseup.net>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: Simplify per-instruction compression control setup in generator.

By using the new compression/group control interface.  This will allow
easier extension to support arbitrary channel enable groups at the IR
level.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/fs: No need to set compression control at the top of generate_code().

The right value is dependent on the specific IR instruction being
generated so it has to be reset in every iteration of the loop anyway.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>

i965/eu: Fix a bunch of compression control bugs in the generator.

Most of these were resetting quarter control to zero incorrectly even
though everything they needed to do was disable instruction
compression -- The brw_SAMPLE() case was doing the right thing but it
can be simplified slightly by using the new compression control
interface.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>