mesa-clc: vendor as full-fork recipe (path=source, patches baked)

local/recipes/dev/mesa-clc/source/docs/nir/alu.rst

@@ -0,0 +1,72 @@
+NIR ALU Instructions
+====================
+
+ALU instructions represent simple operations, such as addition, multiplication,
+comparison, etc., that take a certain number of arguments and return a result
+that only depends on the arguments.  ALU instructions in NIR must be pure in
+the sense that they have no side effect and that identical inputs yields an
+identical output.  A good rule of thumb is that only things which can be
+constant folded should be ALU operations.  If it can't be constant folded, then
+it should probably be an intrinsic instead.
+
+Each ALU instruction has an opcode, which is a member of the :c:enum:`nir_op`
+enum, that describes what it does as well as how many arguments it takes.
+Associated with each opcode is an metadata structure,
+:c:struct:`nir_op_info`, which shows how many arguments the opcode takes,
+information about data types, and algebraic properties such as associativity
+and commutativity. The info structure for each opcode may be accessed through
+a global :c:var:`nir_op_infos` array that’s indexed by the opcode.
+
+ALU operations are typeless, meaning that they're only defined to convert
+a certain bit-pattern input to another bit-pattern output.  The only concrete
+notion of types for a NIR SSA value or register is that each value has a number
+of vector components and a bit-size.  How that data is interpreted is entirely
+controlled by the opcode.  NIR doesn't have opcodes for ``intBitsToFloat()``
+and friends because they are implicit.
+
+Even though ALU operations are typeless, each opcode also has an "ALU type"
+metadata for each of the sources and the destination which can be
+floating-point, boolean, integer, or unsigned integer.  The ALU type mainly
+helps back-ends which want to handle all conversion instructions, for instance,
+in a single switch case.  They're also important when a back-end requests the
+absolute value, negate, and saturate modifiers (not used by core NIR).  In that
+case, modifiers are interpreted with respect to the ALU type on the source or
+destination of the instruction.  In addition, if an operation takes a boolean
+argument, then the argument may be assumed to be either ``0`` for false or
+``~0`` (a.k.a ``-1``) for true even if it is not a 1-bit value.  If an
+operation’s result has a boolean type, then it may only produce only ``0`` or ``~0``.
+
+Most of the common ALU ops in NIR operate per-component, meaning that the
+operation is defined by what it does on a single scalar value and, when
+performed on vectors, it performs the same operation on each component.  Things
+like add, multiply, etc. fall into this category.  Per-component operations
+naturally scale to as many components as necessary.  Non-per-component ALU ops
+are things like :nir:alu-op:`vec4` or :nir:alu-op:`pack_64_2x32` where any
+given component in the result value may be a combination of any component in
+any source.  These ops have a number of destination components and a number of
+components required by each source which is fixed by the opcode.
+
+While most instruction types in NIR require vector sizes to perfectly match on
+inputs and outputs, ALU instruction sources have an additional
+:c:member:`nir_alu_src.swizzle` field which allows them to act on vectors
+which are not the native vector size of the instruction.  This is ideal for
+hardware with a native data type of :nir:alu-op:`vec4` but also means that ALU
+instructions are often used (and required) for packing/unpacking vectors for
+use in other instruction types like intrinsics or texture ops.
+
+.. c:autostruct:: nir_op_info
+   :file: src/compiler/nir/nir.h
+   :members:
+
+.. c:autovar:: nir_op_infos
+
+.. c:autostruct:: nir_alu_instr
+   :members:
+
+.. c:autostruct:: nir_alu_src
+   :members:
+
+NIR ALU Opcode Reference:
+-------------------------
+
+.. nir:alu-opcodes::

local/recipes/dev/mesa-clc/source/docs/nir/index.rst

@@ -0,0 +1,15 @@
+NIR Intermediate Representation (NIR)
+=====================================
+
+The NIR Intermediate Representation (NIR) is the optimizing compiler stack that
+sits at the core of most Mesa drivers' shader compilers.  It consists of a set
+of enums and data structures that make up the IR as well as a suite of helper
+functions, optimization passes, and lowering passes for building a compiler
+stack.
+
+.. toctree::
+   :maxdepth: 2
+
+   alu
+   tex
+   unit-testing

local/recipes/dev/mesa-clc/source/docs/nir/tex.rst

@@ -0,0 +1,83 @@
+NIR Texture Instructions
+========================
+
+Even though texture instructions *could* be supported as intrinsics, the vast
+number of combinations mean that doing so is practically impossible. Instead,
+NIR has a dedicated texture instruction.  There are several texture operations:
+
+.. c:autoenum:: nir_texop
+   :file: src/compiler/nir/nir.h
+   :members:
+
+As with other instruction types, there is still an array of sources, except
+that each source also has a *type* associated with it.  There are various
+source types, each corresponding to a piece of information that the different
+texture operations require.
+
+.. c:autoenum:: nir_tex_src_type
+   :members:
+
+Of particular interest are the texture/sampler deref/index/handle source types.
+First, note that textures and samplers are specified separately in NIR.  While
+not required for OpenGL, this is required for Vulkan and OpenCL.  Some
+OpenGL [ES] drivers have to deal with hardware that does not have separate
+samplers and textures.  While not recommended, an OpenGL-only driver may assume
+that the texture and sampler derefs will always point to the same resource, if
+needed.  Note that this pretty well paints your compiler into a corner and
+makes any future port to Vulkan or OpenCL harder, so such assumptions should
+really only be made if targeting OpenGL ES 2.0 era hardware.
+
+Also, like a lot of other resources, there are multiple ways to represent a
+texture in NIR. It can be referenced by a variable dereference, an index, or a
+bindless handle. When using an index or a bindless handle, the texture type
+information is generally not available.  To handle this, various information
+from the type is redundantly stored in the :c:struct:`nir_tex_instr` itself.
+
+.. c:autostruct:: nir_tex_instr
+   :members:
+
+.. c:autostruct:: nir_tex_src
+   :members:
+
+Texture instruction helpers
+---------------------------
+
+There are a number of helper functions for working with NIR texture
+instructions.  They are documented here in no particular order.
+
+.. c:autofunction:: nir_tex_instr_create
+
+.. c:autofunction:: nir_tex_instr_need_sampler
+
+.. c:autofunction:: nir_tex_instr_result_size
+
+.. c:autofunction:: nir_tex_instr_dest_size
+
+.. c:autofunction:: nir_tex_instr_is_query
+
+.. c:autofunction:: nir_tex_instr_has_implicit_derivative
+
+.. c:autofunction:: nir_tex_instr_src_type
+
+.. c:autofunction:: nir_tex_instr_src_size
+
+.. c:autofunction:: nir_tex_instr_src_index
+
+.. c:autofunction:: nir_tex_instr_add_src
+
+.. c:autofunction:: nir_tex_instr_remove_src
+
+Texture instruction lowering
+----------------------------
+
+Because most hardware only supports some subset of all possible GLSL/SPIR-V
+texture operations, NIR provides a quite powerful lowering pass which is able
+to implement more complex texture operations in terms of simpler ones.
+
+.. c:autofunction:: nir_lower_tex
+
+.. c:autostruct:: nir_lower_tex_options
+   :members:
+
+.. c:autoenum:: nir_lower_tex_packing
+   :members:

local/recipes/dev/mesa-clc/source/docs/nir/unit-testing.rst

@@ -0,0 +1,80 @@
+Writing and running NIR unit tests
+==================================
+
+NIR uses `gtest <https://github.com/google/googletest>`__
+for unit testing lowering and optimization passes. Tests
+should declare a class to use for all test cases:
+
+.. code:: c++
+
+   class nir_my_pass_test : public nir_test {
+   protected:
+      nir_my_pass_test();
+
+      void run_pass(nir_reference_shader expected);
+
+      /* Resources used by the test */
+   }
+
+   nir_my_pass_test::nir_my_pass_test()
+      : nir_test::nir_test("nir_my_pass_test")
+   {
+      /* Create resources used by the test */
+   }
+
+   void nir_my_pass_test::run_pass(nir_reference_shader expected)
+   {
+      nir_validate_shader(b->shader, "before nir_my_pass");
+      NIR_PASS(_, b->shader, nir_my_pass);
+      check_nir_string(expected);
+   }
+
+With this setup, the individual test cases can use ``nir_builder``
+to initialize a shader, run the pass that should be tested on it
+and compare it against a string, containing the expected pass
+output:
+
+.. code:: c++
+
+   TEST_F(nir_my_pass_test, basic)
+   {
+      run_pass(NIR_REFERENCE_SHADER(R"(
+         shader: MESA_SHADER_COMPUTE
+         name: nir_my_pass_test
+         workgroup_size: 1, 1, 1
+         subgroup_size: 0
+         decl_function main () (entrypoint)
+
+         impl main {
+             block b0:  // preds:
+         }
+      )"));
+   }
+
+The expected string can be managed using the
+``bin/nir-test-runner.py`` script which builds and runs NIR
+unit tests. The output of the tests are compared against the
+expectations. The runner can then optionally update the
+expectations.
+
+.. code::
+
+   user@distro:~/mesa$ python bin/nir-test-runner.py -Bbuild
+   INFO: autodetecting backend as ninja
+   INFO: calculating backend command to run: /usr/bin/ninja -C /home/konstantin/dev/mesa/build
+   ninja: Entering directory `/home/konstantin/dev/mesa/build'
+   [1/1] Generating src/git_sha1.h with a custom command
+   diff --git a/src/compiler/nir/tests/opt_loop_tests.cpp b/src/compiler/nir/tests/opt_loop_tests.cpp
+   index 05d3c6357c6..0d4810c5f85 100644
+   --- a/src/compiler/nir/tests/opt_loop_tests.cpp
+   +++ b/src/compiler/nir/tests/opt_loop_tests.cpp
+   @@ -136,6 +136,7 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_break_in_then)
+       check_nir_string(NIR_REFERENCE_SHADER(R"(
+          shader: MESA_SHADER_FRAGMENT
+          name: nir_opt_loop_test
+   +      subgroup_size: 0
+          decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0)
+          decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0)
+          decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0)
+
+   Apply the changes listed above? [Y/n]y