Merge pull request #28 from blueskythlikesclouds/input-attachment-to-srv

Convert input attachments to textures when lowering them.

.github/workflows/ci.yaml

@@ -0,0 +1,179 @@
+name: Godot Mesa/NIR static libs
+
+on:
+  push:
+  pull_request:
+  workflow_dispatch:
+
+jobs:
+  build:
+    name: ${{ matrix.name }}
+    strategy:
+      matrix:
+        include:
+        # MinGW/LLVM libs using UCRT
+        - name: 🏁 Windows - MinGW/LLVM (UCRT) x86_64
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_64-llvm-release
+          artifact-path: bin/libNIR.windows.x86_64.a
+          flags: use_mingw=yes arch=x86_64 use_llvm=yes mingw_prefix=$HOME/llvm-mingw
+          llvm: yes
+
+        - name: 🏁 Windows - MinGW/LLVM (UCRT) x86_32
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_32-llvm-release
+          artifact-path: bin/libNIR.windows.x86_32.a
+          flags: use_mingw=yes arch=x86_32 use_llvm=yes mingw_prefix=$HOME/llvm-mingw
+          llvm: yes
+
+        - name: 🏁 Windows - MinGW/LLVM (UCRT) arm64
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-arm64-llvm-release
+          artifact-path: bin/libNIR.windows.arm64.a
+          flags: use_mingw=yes arch=arm64 use_llvm=yes mingw_prefix=$HOME/llvm-mingw
+          llvm: yes
+
+        # MSVC libs
+        - name: 🏁 Windows - MSVC x86_64
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_64-msvc-release
+          artifact-path: bin/libNIR.windows.x86_64.lib
+          flags: use_mingw=no arch=x86_64
+
+        - name: 🏁 Windows - MSVC x86_32
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_32-msvc-release
+          artifact-path: bin/libNIR.windows.x86_32.lib
+          flags: use_mingw=no arch=x86_32
+
+        - name: 🏁 Windows - MSVC arm64
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-arm64-msvc-release
+          artifact-path: bin/libNIR.windows.arm64.lib
+          flags: use_mingw=no arch=arm64
+
+        # MinGW/GCC libs using MSVCRT
+        - name: 🏁 Windows - MinGW/GCC (MSVCRT) x86_64
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_64-gcc-release
+          artifact-path: bin/libNIR.windows.x86_64.a
+          flags: use_mingw=yes arch=x86_64
+          mingw: yes
+          msys: mingw64
+          msysenv: x86_64
+
+        - name: 🏁 Windows - MinGW/GCC (MSVCRT) x86_32
+          platform: windows
+          os: windows-2022
+          artifact-name: godot-nir-static-x86_32-gcc-release
+          artifact-path: bin/libNIR.windows.x86_32.a
+          flags: use_mingw=yes arch=x86_32
+          mingw: yes
+          msys: mingw32
+          msysenv: i686
+
+    runs-on: ${{ matrix.os }}
+
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          submodules: 'true'
+
+      - name: Set up Python
+        uses: actions/setup-python@v5
+        with:
+          python-version: '3.x'
+
+      - name: Install SCons (Native env)
+        if: ${{ matrix.mingw != 'yes' }}
+        run: |
+          python -m pip install scons==4.9.0
+
+      - name: Install mako
+        run: |
+          python -m pip install mako
+
+      - name: Install pyyaml
+        run: |
+          python -m pip install pyyaml
+
+      - name: Setup MinGW/LLVM
+        if: ${{ matrix.platform == 'windows' && matrix.llvm == 'yes' }}
+        run: |
+          curl -L -O https://github.com/mstorsjo/llvm-mingw/releases/download/20240619/llvm-mingw-20240619-ucrt-x86_64.zip
+          unzip -q llvm-mingw-*.zip
+          rm llvm-mingw-*.zip
+          mv llvm-mingw-* "$HOME/llvm-mingw"
+          echo "$HOME/llvm-mingw/bin" >> $GITHUB_PATH
+
+      - name: Setup MinGW/MSYS2
+        if: ${{ matrix.mingw == 'yes' }}
+        uses: msys2/setup-msys2@v2
+        with:
+          update: true
+          msystem: ${{matrix.msys}}
+          install: mingw-w64-${{matrix.msysenv}}-gcc mingw-w64-${{matrix.msysenv}}-scons
+
+      - name: Prepare Mesa source
+        shell: bash
+        run: |
+          ./update_mesa.sh
+
+      - name: Build Mesa (Native env)
+        if: ${{ matrix.mingw != 'yes' }}
+        run: |
+          scons platform=${{ matrix.platform }} ${{ matrix.flags }} optimize=speed
+
+      - name: Build Mesa (MSYS2 env)
+        if: ${{ matrix.mingw == 'yes' }}
+        shell: msys2 {0}
+        run: |
+          scons platform=${{ matrix.platform }} ${{ matrix.flags }} optimize=speed
+
+      - name: Upload artifact
+        uses: actions/upload-artifact@v4
+        with:
+          name: ${{ matrix.artifact-name }}
+          path: |
+            ${{ matrix.artifact-path }}
+            godot-mesa/VERSION.info
+            godot-mesa/**/*.h
+            godot-mesa/**/*.hpp
+          if-no-files-found: error
+
+  release:
+    name: 'Release'
+
+    needs: [build]
+
+    runs-on: ubuntu-latest
+
+    if: ${{ startsWith(github.ref, 'refs/tags/') }}
+
+    permissions:
+      contents: write
+
+    steps:
+      - name: Download Artifacts
+        uses: actions/download-artifact@v4
+
+      - name: ZIP Artifacts
+        run: |
+          for i in */; do (cd "${i%/}"; zip -r "../${i%/}.zip" .; cd ..); done
+
+      - name: Create Release
+        uses: ncipollo/release-action@v1
+        with:
+          allowUpdates: true
+          removeArtifacts: true
+          omitNameDuringUpdate: true
+          omitBodyDuringUpdate: true
+          artifacts: "*.zip"
+          artifactErrorsFailBuild: true

.gitignore

@@ -1,3 +1,4 @@
 .sconsign.dblite
 bin/
 godot-mesa/
+*.pyc

.gitmodules

@@ -1,3 +1,3 @@
 [submodule "mesa"]
 	path = mesa
-	url = https://github.com/Mesa3D/mesa
+	url = https://gitlab.freedesktop.org/mesa/mesa

LICENSE.md

@@ -0,0 +1,21 @@
+# MIT License
+
+Copyright (c) 2017-present Godot Engine contributors.
+
+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 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.

SConstruct

@@ -3,9 +3,7 @@
 import os
 import platform
 import sys
-import subprocess
 from pathlib import Path
-from SCons.Errors import UserError
 
 EnsureSConsVersion(4, 0)
 
@@ -113,6 +111,14 @@ opts.Add(
     )
 )
 
+opts.Add("cppdefines", "Custom defines for the pre-processor")
+opts.Add("ccflags", "Custom flags for both the C and C++ compilers")
+opts.Add("cxxflags", "Custom flags for the C++ compiler")
+opts.Add("cflags", "Custom flags for the C compiler")
+opts.Add("linkflags", "Custom flags for the linker")
+opts.Add("extra_suffix", "Custom extra suffix added to the base filename of all generated binary files.", "")
+opts.Add(BoolVariable("use_asan", "Use address sanitizer (ASAN) in MSVC", False))
+
 # Targets flags tool (optimizations, debug symbols)
 target_tool = Tool("targets", toolpath=["godot-tools"])
 target_tool.options(opts)
@@ -157,6 +163,12 @@ if unknown:
 
 print("Building for architecture " + env["arch"] + " on platform " + env["platform"])
 
+env.Append(CPPDEFINES=env.get("cppdefines", "").split())
+env.Append(CCFLAGS=env.get("ccflags", "").split())
+env.Append(CXXFLAGS=env.get("cxxflags", "").split())
+env.Append(CFLAGS=env.get("cflags", "").split())
+env.Append(LINKFLAGS=env.get("linkflags", "").split())
+
 # Require C++17
 if env.get("is_msvc", False):
     env.Append(CXXFLAGS=["/std:c++17"])
@@ -178,6 +190,12 @@ if env["platform"] == "macos":
 elif env["platform"] == "windows":
     env.AppendUnique(CPPDEFINES=["WINVER=0x0603", "_WIN32_WINNT=0x0603"])
 
+# Sanitizers.
+if env.get("use_asan", False) and env.get("is_msvc", False):
+    env["extra_suffix"] += ".san"
+    env.Append(LINKFLAGS=["/INFERASANLIBS"])
+    env.Append(CCFLAGS=["/fsanitize=address"])
+
 
 scons_cache_path = os.environ.get("SCONS_CACHE")
 if scons_cache_path is not None:
@@ -197,12 +215,12 @@ custom_build_steps = [
     ],
     ["src/compiler/nir", "nir_builder_opcodes_h.py > %s/nir_builder_opcodes.h", "nir_builder_opcodes.h"],
     ["src/compiler/nir", "nir_constant_expressions.py > %s/nir_constant_expressions.c", "nir_constant_expressions.c"],
-    ["src/compiler/nir", "nir_intrinsics_h.py --outdir %s", "nir_intrinsics.h"],
-    ["src/compiler/nir", "nir_intrinsics_c.py --outdir %s", "nir_intrinsics.c"],
-    ["src/compiler/nir", "nir_intrinsics_indices_h.py --outdir %s", "nir_intrinsics_indices.h"],
+    ["src/compiler/nir", "nir_intrinsics_h.py --out %s/nir_intrinsics.h", "nir_intrinsics.h"],
+    ["src/compiler/nir", "nir_intrinsics_c.py --out %s/nir_intrinsics.c", "nir_intrinsics.c"],
+    ["src/compiler/nir", "nir_intrinsics_indices_h.py --out %s/nir_intrinsics_indices.h", "nir_intrinsics_indices.h"],
     ["src/compiler/nir", "nir_opcodes_h.py > %s/nir_opcodes.h", "nir_opcodes.h"],
     ["src/compiler/nir", "nir_opcodes_c.py > %s/nir_opcodes.c", "nir_opcodes.c"],
-    ["src/compiler/nir", "nir_opt_algebraic.py > %s/nir_opt_algebraic.c", "nir_opt_algebraic.c"],
+    ["src/compiler/nir", "nir_opt_algebraic.py --out %s/nir_opt_algebraic.c", "nir_opt_algebraic.c"],
     ["src/compiler/spirv", "vtn_generator_ids_h.py spir-v.xml %s/vtn_generator_ids.h", "vtn_generator_ids.h"],
     [
         "src/microsoft/compiler",
@@ -210,8 +228,11 @@ custom_build_steps = [
         "dxil_nir_algebraic.c",
     ],
     ["src/util", "format_srgb.py > %s/format_srgb.c", "format_srgb.c"],
-    ["src/util/format", "u_format_table.py u_format.csv --header > %s/u_format_pack.h", "u_format_pack.h"],
-    ["src/util/format", "u_format_table.py u_format.csv > %s/u_format_table.c", "u_format_table.c"],
+    ["src/util/format", "u_format_table.py u_format.yaml --enums > %s/u_format_gen.h", "u_format_gen.h"],
+    ["src/util/format", "u_format_table.py u_format.yaml --header > %s/u_format_pack.h", "u_format_pack.h"],
+    ["src/util/format", "u_format_table.py u_format.yaml > %s/u_format_table.c", "u_format_table.c"],
+    ["src/compiler", "builtin_types_h.py %s/builtin_types.h", "builtin_types.h"],
+    ["src/compiler", "builtin_types_c.py %s/builtin_types.c", "builtin_types.c"],
 ]
 
 mesa_sources = []
@@ -274,15 +295,22 @@ extra_defines = [
     "WINDOWS_NO_FUTEX",
 ]
 
+if env["optimize"] != "debug":
+    extra_defines += ["NDEBUG"]
+
 # These defines are inspired by the Meson build scripts in the original repo.
 extra_defines += [
     "__STDC_CONSTANT_MACROS",
     "__STDC_FORMAT_MACROS",
     "__STDC_LIMIT_MACROS",
-    ("PACKAGE_VERSION", '\\"' + Path(mesa_absdir + "/VERSION").read_text().strip() + '\\"'),
+    ("PACKAGE_VERSION", '\\"' + Path(mesa_absdir + "/VERSION.info").read_text().strip() + '\\"'),
     ("PACKAGE_BUGREPORT", '\\"https://gitlab.freedesktop.org/mesa/mesa/-/issues\\"'),
     "PIPE_SUBSYSTEM_WINDOWS_USER",
     "_USE_MATH_DEFINES",
+    "BLAKE3_NO_SSE2",
+    "BLAKE3_NO_SSE41",
+    "BLAKE3_NO_AVX2",
+    "BLAKE3_NO_AVX512",
 ]
 
 if env.get("is_msvc", False):
@@ -297,23 +325,33 @@ if env.get("is_msvc", False):
         ("_HAS_EXCEPTIONS", 0),
         "NOMINMAX",
         "HAVE_STRUCT_TIMESPEC",
+        "HAVE_TIMESPEC_GET",
         ("_Static_assert", "static_assert"),
     ]
     env.Append(CFLAGS=["/std:c11"])
 else:
     env.Append(
         CPPDEFINES=[
-            ("__MSVCRT_VERSION__", 0x0700),
             "HAVE_STRUCT_TIMESPEC",
         ]
     )
     env.Append(CFLAGS=["-std=c11"])
+    env.Append(CXXFLAGS=["-fno-exceptions"])
+
+if env.get("use_llvm", False):
+    extra_defines += [
+        ("__MSVCRT_VERSION__", 0x0700),
+        "HAVE_TIMESPEC_GET",
+        "_UCRT",
+    ]
 
 env.Append(CPPDEFINES=extra_defines)
 env.Append(CPPPATH=".")
 env.Append(CPPPATH="#vulkan/include")
 
 suffix = ".{}.{}".format(env["platform"], env["arch"])
+if env.get("extra_suffix", "") != "":
+    suffix += "." + env["extra_suffix"]
 
 # Expose it when included from another project
 env["suffix"] = suffix

drivers/d3d12/d3d12_godot_nir_bridge.h

@@ -53,6 +53,10 @@ typedef struct GodotNirCallbacks {
 	void (*report_bitcode_bit_offset_fn)(uint64_t p_bit_offset, void *p_data);
 } GodotNirCallbacks;
 
+extern void *godot_nir_malloc(size_t p_size);
+extern void *godot_nir_realloc(void *p_block, size_t p_size);
+extern void godot_nir_free(void *p_block);
+
 #ifdef __cplusplus
 }
 #endif

godot-patches/01_godot_nir_goodies.patch

@@ -1,47 +1,171 @@
+diff --git a/godot-mesa/src/compiler/nir/nir_dominance_lca.c b/godot-mesa/src/compiler/nir/nir_dominance_lca.c
+index 84a7f1d..40c0811 100644
+--- a/godot-mesa/src/compiler/nir/nir_dominance_lca.c
++++ b/godot-mesa/src/compiler/nir/nir_dominance_lca.c
+@@ -167,7 +167,7 @@ nir_dominance_lca(nir_block *b1, nir_block *b2)
+    uint32_t i1 = dom_lca_representative(b1);
+    uint32_t i2 = dom_lca_representative(b2);
+    if (i1 > i2)
+-      SWAP(i1, i2);
++      MESA_SWAP(i1, i2);
+    uint32_t index = range_minimum_query(&impl->dom_lca_info.table, i1, i2 + 1);
+    nir_block *result = impl->dom_lca_info.block_from_idx[index];
+ 
 diff --git a/godot-mesa/src/compiler/nir/nir_intrinsics.py b/godot-mesa/src/compiler/nir/nir_intrinsics.py
-index fcb8a9885c..54e199cc43 100644
+index 41371db..dd3454f 100644
 --- a/godot-mesa/src/compiler/nir/nir_intrinsics.py
 +++ b/godot-mesa/src/compiler/nir/nir_intrinsics.py
-@@ -1046,6 +1046,9 @@ load("push_constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINAT
+@@ -1262,6 +1262,9 @@ load("push_constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINAT
  # src[] = { offset }.
- load("constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET],
+ load("constant", [1], [BASE, RANGE, ACCESS, ALIGN_MUL, ALIGN_OFFSET],
       [CAN_ELIMINATE, CAN_REORDER])
 +# src[] = { offset }.
-+load("constant_non_opt", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET],
++load("constant_non_opt", [1], [BASE, RANGE, ACCESS, ALIGN_MUL, ALIGN_OFFSET],
 +     [CAN_ELIMINATE, CAN_REORDER])
  # src[] = { address }.
  load("global", [1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
- # src[] = { address }.
+ # src[] = { base_address, offset, bound }.
+diff --git a/godot-mesa/src/compiler/nir/nir_loop_analyze.c b/godot-mesa/src/compiler/nir/nir_loop_analyze.c
+index aa7369c..34b8a99 100644
+--- a/godot-mesa/src/compiler/nir/nir_loop_analyze.c
++++ b/godot-mesa/src/compiler/nir/nir_loop_analyze.c
+@@ -1020,7 +1020,7 @@ try_find_trip_count_vars_in_logical_op(nir_scalar *cond,
+ 
+       if (!nir_scalar_is_alu(logical_op) || !nir_scalar_is_const(zero)) {
+          /* Maybe we had it the wrong way, flip things around */
+-         SWAP(zero, logical_op);
++         MESA_SWAP(zero, logical_op);
+ 
+          /* If we still didn't find what we need then return */
+          if (!nir_scalar_is_const(zero))
+diff --git a/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c b/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
+index 830746b..74c8df6 100644
+--- a/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
++++ b/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
+@@ -530,7 +530,7 @@ split_cmat_load_store(nir_builder *b,
+          col_offset = (i / split->num_col_splits) * desc.rows;
+ 
+          if (layout == GLSL_MATRIX_LAYOUT_ROW_MAJOR)
+-            SWAP(row_offset, col_offset);
++            MESA_SWAP(row_offset, col_offset);
+ 
+          ptr_deref = nir_build_deref_cast(b, &addr_deref->def, addr_deref->modes, scalar_type, elem_size);
+          stride = nir_udiv_imm(b, nir_imul_imm(b, stride, deref_bytes_size), elem_size);
+diff --git a/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c b/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
+index 323b188..09460ff 100644
+--- a/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
++++ b/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
+@@ -106,12 +106,53 @@ load_coord(nir_builder *b, nir_deref_instr *deref,
+    }
+ }
+ 
++static const struct glsl_type *
++get_texture_type_for_image(const struct glsl_type *type)
++{
++   if (glsl_type_is_array(type)) {
++      const struct glsl_type *elem_type =
++         get_texture_type_for_image(glsl_get_array_element(type));
++      return glsl_array_type(elem_type, glsl_get_length(type), 0 /*explicit size*/);
++   }
++
++   assert((glsl_type_is_image(type)));
++   return glsl_texture_type(glsl_get_sampler_dim(type),
++                            glsl_sampler_type_is_array(type),
++                            glsl_get_sampler_result_type(type));
++}
++
++static bool
++replace_image_type_with_texture(nir_deref_instr *deref)
++{
++   const struct glsl_type *type = deref->type;
++
++   /* If we've already chased up the deref chain this far from a different intrinsic, we're done */
++   if (!glsl_type_is_image(glsl_without_array(type)))
++      return false;
++
++   deref->type = get_texture_type_for_image(type);
++   deref->modes = nir_var_uniform;
++   if (deref->deref_type == nir_deref_type_var) {
++      type = deref->var->type;
++      if (glsl_type_is_image(glsl_without_array(type))) {
++         deref->var->type = get_texture_type_for_image(type);
++         deref->var->data.mode = nir_var_uniform;
++         memset(&deref->var->data.sampler, 0, sizeof(deref->var->data.sampler));
++      }
++   } else {
++      nir_deref_instr *parent = nir_deref_instr_parent(deref);
++      if (parent)
++         replace_image_type_with_texture(parent);
++   }
++
++   return true;
++}
++
+ static bool
+ try_lower_input_load(nir_builder *b, nir_intrinsic_instr *load,
+                      const nir_input_attachment_options *options)
+ {
+    nir_deref_instr *deref = nir_src_as_deref(load->src[0]);
+-   assert(glsl_type_is_image(deref->type));
+ 
+    enum glsl_sampler_dim image_dim = glsl_get_sampler_dim(deref->type);
+    if (image_dim != GLSL_SAMPLER_DIM_SUBPASS &&
+@@ -172,6 +213,8 @@ try_lower_input_load(nir_builder *b, nir_intrinsic_instr *load,
+                            &tex->def);
+    }
+ 
++   replace_image_type_with_texture(deref);
++
+    return true;
+ }
+ 
+@@ -204,6 +247,8 @@ try_lower_input_texop(nir_builder *b, nir_tex_instr *tex,
+ 
+    nir_src_rewrite(&tex->src[coord_src_idx].src, coord);
+ 
++   replace_image_type_with_texture(deref);
++
+    return true;
+ }
+ 
+diff --git a/godot-mesa/src/compiler/nir/nir_opt_reassociate.c b/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
+index 09566a5..15c5ad1 100644
+--- a/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
++++ b/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
+@@ -118,7 +118,7 @@ get_pair_key(nir_op op, nir_scalar a, nir_scalar b)
+    if ((a.def->index > b.def->index) ||
+        ((a.def->index == b.def->index) && (a.comp > b.comp))) {
+ 
+-      SWAP(a, b);
++      MESA_SWAP(a, b);
+    }
+ 
+    return (struct pair_key){
+@@ -460,8 +460,8 @@ reassociate_chain(struct chain *c, void *pair_freq)
+ 
+       if (best_pair.i != best_pair.j) {
+          /* Pin the best pair at the front. The rest is sorted by rank. */
+-         SWAP(c->srcs[0], c->srcs[best_pair.i]);
+-         SWAP(c->srcs[1], c->srcs[best_pair.j]);
++         MESA_SWAP(c->srcs[0], c->srcs[best_pair.i]);
++         MESA_SWAP(c->srcs[1], c->srcs[best_pair.j]);
+          pinned = 2;
+       }
+    }
 diff --git a/godot-mesa/src/compiler/spirv/spirv_to_nir.c b/godot-mesa/src/compiler/spirv/spirv_to_nir.c
-index 5cb7691506..7620fef9f5 100644
+index 7150419..61b0b7c 100644
 --- a/godot-mesa/src/compiler/spirv/spirv_to_nir.c
 +++ b/godot-mesa/src/compiler/spirv/spirv_to_nir.c
-@@ -39,6 +39,8 @@
+@@ -45,6 +45,8 @@
  
  #include <stdio.h>
  
 +#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
 +
- #ifndef NDEBUG
- uint32_t mesa_spirv_debug = 0;
- 
-@@ -1121,6 +1123,7 @@ struct_member_decoration_cb(struct vtn_builder *b,
-    case SpvDecorationPerPrimitiveNV:
-    case SpvDecorationPerTaskNV:
-    case SpvDecorationPerViewNV:
-+   case SpvDecorationInvariant: /* Silence this one to avoid warning spam. */
-       break;
- 
-    case SpvDecorationSpecId:
-@@ -1129,7 +1132,6 @@ struct_member_decoration_cb(struct vtn_builder *b,
-    case SpvDecorationArrayStride:
-    case SpvDecorationGLSLShared:
-    case SpvDecorationGLSLPacked:
--   case SpvDecorationInvariant:
-    case SpvDecorationAliased:
-    case SpvDecorationConstant:
-    case SpvDecorationIndex:
-@@ -1938,7 +1940,7 @@ vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
+ #ifndef PATH_MAX
+ #define PATH_MAX 4096
+ #endif
+@@ -2481,7 +2483,7 @@ vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
  }
  
  static void
@@ -50,7 +174,7 @@ index 5cb7691506..7620fef9f5 100644
                              ASSERTED int member,
                              const struct vtn_decoration *dec, void *data)
  {
-@@ -1946,13 +1948,8 @@ spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
+@@ -2489,13 +2491,8 @@ spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
     if (dec->decoration != SpvDecorationSpecId)
        return;
  
@@ -66,7 +190,7 @@ index 5cb7691506..7620fef9f5 100644
  }
  
  static void
-@@ -1976,6 +1973,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
+@@ -2519,6 +2516,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
                      const uint32_t *w, unsigned count)
  {
     struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
@@ -79,7 +203,7 @@ index 5cb7691506..7620fef9f5 100644
     val->constant = rzalloc(b, nir_constant);
     switch (opcode) {
     case SpvOpConstantTrue:
-@@ -1993,7 +1996,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
+@@ -2536,7 +2539,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
  
        if (opcode == SpvOpSpecConstantTrue ||
            opcode == SpvOpSpecConstantFalse)
@@ -88,7 +212,7 @@ index 5cb7691506..7620fef9f5 100644
  
        val->constant->values[0].b = u32val.u32 != 0;
        break;
-@@ -2024,11 +2027,10 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
+@@ -2567,14 +2570,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
  
        if (opcode == SpvOpSpecConstant)
           vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
@@ -98,10 +222,14 @@ index 5cb7691506..7620fef9f5 100644
     }
  
 -   case SpvOpSpecConstantComposite:
-    case SpvOpConstantComposite: {
-       unsigned elem_count = count - 3;
-       vtn_fail_if(elem_count != val->type->length,
-@@ -2076,218 +2078,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
+    case SpvOpConstantComposite:
+-   case SpvOpConstantCompositeReplicateEXT:
+-   case SpvOpSpecConstantCompositeReplicateEXT: {
++   case SpvOpConstantCompositeReplicateEXT: {
+       const unsigned elem_count =
+          val->type->base_type == vtn_base_type_cooperative_matrix ?
+          1 : val->type->length;
+@@ -2648,315 +2649,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
        break;
     }
  
@@ -110,6 +238,54 @@ index 5cb7691506..7620fef9f5 100644
 -      vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
 -      SpvOp opcode = u32op.u32;
 -      switch (opcode) {
+-      case SpvOpBitcast: {
+-         struct vtn_value *src = &b->values[w[4]];
+-
+-         vtn_assert(src->value_type == vtn_value_type_constant ||
+-                    src->value_type == vtn_value_type_undef);
+-
+-         unsigned src_len = glsl_get_vector_elements(src->type->type);
+-         unsigned dst_len = glsl_get_vector_elements(val->type->type);
+-
+-         unsigned src_bit_size = glsl_get_bit_size(src->type->type);
+-         unsigned dst_bit_size = glsl_get_bit_size(val->type->type);
+-
+-         vtn_assert(src_len * src_bit_size == dst_len * dst_bit_size);
+-
+-         /* This will end up being zero */
+-         if (src->value_type == vtn_value_type_undef)
+-            break;
+-
+-         if (src_bit_size == dst_bit_size) {
+-            /* This is just a copy */
+-            for (unsigned i = 0; i < src_len; i++)
+-               val->constant->values[i] = src->constant->values[i];
+-         } else {
+-            /* You can't non-trivially bitcast booleans */
+-            vtn_assert(src_bit_size >= 8 && dst_bit_size >= 8);
+-            const unsigned src_byte_size = src_bit_size / 8;
+-            const unsigned dst_byte_size = dst_bit_size / 8;
+-
+-            vtn_assert(src_len <= NIR_MAX_VEC_COMPONENTS &&
+-                       dst_len <= NIR_MAX_VEC_COMPONENTS);
+-
+-            uint8_t bits[NIR_MAX_VEC_COMPONENTS * sizeof(nir_const_value)];
+-
+-            for (unsigned i = 0; i < src_len; i++) {
+-               uint64_t v = nir_const_value_as_int(src->constant->values[i],
+-                                                   src_bit_size);
+-               memcpy(bits + i * src_byte_size, &v, src_byte_size);
+-            }
+-
+-            for (unsigned i = 0; i < dst_len; i++) {
+-               uint64_t v = 0;
+-               memcpy(&v, bits + i * dst_byte_size, dst_byte_size);
+-               val->constant->values[i] =
+-                  nir_const_value_for_uint(v, dst_bit_size);
+-            }
+-         }
+-         break;
+-      }
 -      case SpvOpVectorShuffle: {
 -         struct vtn_value *v0 = &b->values[w[4]];
 -         struct vtn_value *v1 = &b->values[w[5]];
@@ -172,39 +348,45 @@ index 5cb7691506..7620fef9f5 100644
 -         } else {
 -            comp = vtn_value(b, w[5], vtn_value_type_constant);
 -            deref_start = 6;
--            val->constant = nir_constant_clone(comp->constant,
--                                               (nir_variable *)b);
+-            val->constant = nir_constant_clone(comp->constant, b->shader);
 -            c = &val->constant;
 -         }
 -
 -         int elem = -1;
 -         const struct vtn_type *type = comp->type;
 -         for (unsigned i = deref_start; i < count; i++) {
--            vtn_fail_if(w[i] > type->length,
--                        "%uth index of %s is %u but the type has only "
--                        "%u elements", i - deref_start,
--                        spirv_op_to_string(opcode), w[i], type->length);
+-            if (type->base_type == vtn_base_type_cooperative_matrix) {
+-               /* Cooperative matrices are always scalar constants.  We don't
+-                * care about the index w[i] because it's always replicated.
+-                */
+-               type = type->component_type;
+-            } else {
+-               vtn_fail_if(w[i] > type->length,
+-                           "%uth index of %s is %u but the type has only "
+-                           "%u elements", i - deref_start,
+-                           spirv_op_to_string(opcode), w[i], type->length);
 -
--            switch (type->base_type) {
--            case vtn_base_type_vector:
--               elem = w[i];
--               type = type->array_element;
--               break;
+-               switch (type->base_type) {
+-               case vtn_base_type_vector:
+-                  elem = w[i];
+-                  type = type->array_element;
+-                  break;
 -
--            case vtn_base_type_matrix:
--            case vtn_base_type_array:
--               c = &(*c)->elements[w[i]];
--               type = type->array_element;
--               break;
+-               case vtn_base_type_matrix:
+-               case vtn_base_type_array:
+-                  c = &(*c)->elements[w[i]];
+-                  type = type->array_element;
+-                  break;
 -
--            case vtn_base_type_struct:
--               c = &(*c)->elements[w[i]];
--               type = type->members[w[i]];
--               break;
+-               case vtn_base_type_struct:
+-                  c = &(*c)->elements[w[i]];
+-                  type = type->members[w[i]];
+-                  break;
 -
--            default:
--               vtn_fail("%s must only index into composite types",
--                        spirv_op_to_string(opcode));
+-               default:
+-                  vtn_fail("%s must only index into composite types",
+-                           spirv_op_to_string(opcode));
+-               }
 -            }
 -         }
 -
@@ -233,10 +415,12 @@ index 5cb7691506..7620fef9f5 100644
 -
 -      default: {
 -         bool swap;
--         nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
--         nir_alu_type src_alu_type = dst_alu_type;
+-
+-         const glsl_type *org_dst_type = val->type->type;
+-         const glsl_type *org_src_type = org_dst_type;
+-
+-         const bool saturate = vtn_has_decoration(b, val, SpvDecorationSaturatedToLargestFloat8NormalConversionEXT);
 -         unsigned num_components = glsl_get_vector_elements(val->type->type);
--         unsigned bit_size;
 -
 -         vtn_assert(count <= 7);
 -
@@ -244,26 +428,31 @@ index 5cb7691506..7620fef9f5 100644
 -         case SpvOpSConvert:
 -         case SpvOpFConvert:
 -         case SpvOpUConvert:
--            /* We have a source in a conversion */
--            src_alu_type =
--               nir_get_nir_type_for_glsl_type(vtn_get_value_type(b, w[4])->type);
--            /* We use the bitsize of the conversion source to evaluate the opcode later */
--            bit_size = glsl_get_bit_size(vtn_get_value_type(b, w[4])->type);
+-            /* We have a different source type in a conversion. */
+-            org_src_type = vtn_get_value_type(b, w[4])->type;
 -            break;
 -         default:
--            bit_size = glsl_get_bit_size(val->type->type);
+-            break;
 -         };
 -
+-         const glsl_type *dst_type = org_dst_type;
+-         if (glsl_type_is_bfloat_16(dst_type) || glsl_type_is_e4m3fn(dst_type) || glsl_type_is_e5m2(dst_type))
+-            dst_type = glsl_float_type();
+-
+-         const glsl_type *src_type = org_src_type;
+-         if (glsl_type_is_bfloat_16(src_type) || glsl_type_is_e4m3fn(src_type) || glsl_type_is_e5m2(src_type))
+-            src_type = glsl_float_type();
+-
 -         bool exact;
 -         nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap, &exact,
--                                                     nir_alu_type_get_type_size(src_alu_type),
--                                                     nir_alu_type_get_type_size(dst_alu_type));
+-                                                     src_type, dst_type);
 -
 -         /* No SPIR-V opcodes handled through this path should set exact.
 -          * Since it is ignored, assert on it.
 -          */
 -         assert(!exact);
 -
+-         unsigned bit_size = glsl_get_bit_size(dst_type);
 -         nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
 -
 -         for (unsigned i = 0; i < count - 4; i++) {
@@ -273,16 +462,30 @@ index 5cb7691506..7620fef9f5 100644
 -            /* If this is an unsized source, pull the bit size from the
 -             * source; otherwise, we'll use the bit size from the destination.
 -             */
--            if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
--               bit_size = glsl_get_bit_size(src_val->type->type);
+-            if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i])) {
+-               if (org_src_type != src_type) {
+-                  /* Small float conversion. */
+-                  assert(i == 0);
+-                  bit_size = glsl_get_bit_size(src_type);
+-               } else {
+-                  bit_size = glsl_get_bit_size(src_val->type->type);
+-               }
+-            }
 -
 -            unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
 -                                 nir_op_infos[op].input_sizes[i] :
 -                                 num_components;
 -
 -            unsigned j = swap ? 1 - i : i;
--            for (unsigned c = 0; c < src_comps; c++)
+-            for (unsigned c = 0; c < src_comps; c++) {
 -               src[j][c] = src_val->constant->values[c];
+-               if (glsl_type_is_bfloat_16(org_src_type))
+-                  src[j][c].f32 = _mesa_bfloat16_bits_to_float(src[j][c].u16);
+-               else if (glsl_type_is_e4m3fn(org_src_type))
+-                  src[j][c].f32 = _mesa_e4m3fn_to_float(src[j][c].u8);
+-               else if (glsl_type_is_e5m2(org_src_type))
+-                  src[j][c].f32 = _mesa_e5m2_to_float(src[j][c].u8);
+-            }
 -         }
 -
 -         /* fix up fixed size sources */
@@ -311,6 +514,28 @@ index 5cb7691506..7620fef9f5 100644
 -         nir_eval_const_opcode(op, val->constant->values,
 -                               num_components, bit_size, srcs,
 -                               b->shader->info.float_controls_execution_mode);
+-
+-         for (int i = 0; i < num_components; i++) {
+-            uint16_t conv;
+-            if (glsl_type_is_bfloat_16(org_dst_type)) {
+-               conv = _mesa_float_to_bfloat16_bits_rte(val->constant->values[i].f32);
+-            } else if (glsl_type_is_e4m3fn(org_dst_type)) {
+-               if (saturate)
+-                  conv = _mesa_float_to_e4m3fn_sat(val->constant->values[i].f32);
+-               else
+-                  conv = _mesa_float_to_e4m3fn(val->constant->values[i].f32);
+-            } else if (glsl_type_is_e5m2(org_dst_type)) {
+-               if (saturate)
+-                  conv = _mesa_float_to_e5m2_sat(val->constant->values[i].f32);
+-               else
+-                  conv = _mesa_float_to_e5m2(val->constant->values[i].f32);
+-            } else {
+-               continue;
+-            }
+-
+-            val->constant->values[i] = nir_const_value_for_raw_uint(conv, glsl_get_bit_size(org_dst_type));
+-         }
+-
 -         break;
 -      } /* default */
 -      }
@@ -320,7 +545,7 @@ index 5cb7691506..7620fef9f5 100644
     case SpvOpConstantNull:
        val->constant = vtn_null_constant(b, val->type);
        val->is_null_constant = true;
-@@ -6370,6 +6160,96 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
+@@ -7018,6 +6710,93 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
     return true;
  }
  
@@ -351,46 +576,43 @@ index 5cb7691506..7620fef9f5 100644
 +      val->value_type = vtn_value_type_ssa;
 +      val->ssa = vtn_create_ssa_value(b, val->type->type);
 +
-+      nir_ssa_def *sc_imm = nir_imm_int(&b->nb, GODOT_NIR_SC_SENTINEL_MAGIC | val->sc_id);
-+      nir_ssa_def *non_opt_const = nir_build_load_constant_non_opt(&b->nb, 1, 32, sc_imm);
++      nir_def *sc_imm = nir_imm_int(&b->nb, GODOT_NIR_SC_SENTINEL_MAGIC | val->sc_id);
++      nir_def *non_opt_const = nir_load_constant_non_opt(&b->nb, 1, 32, sc_imm);
 +
 +      vtn_assert(b->nb.cursor.option == nir_cursor_after_instr);
 +      vtn_assert(b->nb.cursor.instr->type == nir_instr_type_intrinsic);
 +
-+      nir_ssa_def *temp = nir_build_alu(
-+         &b->nb,
-+         nir_op_iand,
-+         non_opt_const,
-+         nir_imm_int(&b->nb, 0xbfffffff),
-+         NULL,
-+         NULL);
 +      val->ssa = vtn_create_ssa_value(b, val->type->type);
 +      if (val->type->type == glsl_uint_type()) {
-+         val->ssa->def = temp;
++         val->ssa->def = non_opt_const;
 +      } else if (val->type->type == glsl_bool_type()) {
 +         val->ssa->def = nir_build_alu(
 +            &b->nb,
 +            nir_op_ine,
-+            temp,
++            non_opt_const,
 +            nir_imm_int(&b->nb, 0),
 +            NULL,
 +            NULL);
 +      } else if (val->type->type == glsl_float_type()) {
-+         val->ssa->def = nir_build_alu(
-+            &b->nb,
-+            nir_op_ishl,
-+            temp,
-+            nir_imm_int(&b->nb, 1),
-+            NULL,
-+            NULL);
++         val->ssa->def = non_opt_const;
 +      } else {
 +         vtn_assert(false);
 +      }
 +   } break;
 +
-+   case SpvOpSpecConstantComposite:
-+      abort(); // Unimplemented
++   case SpvOpSpecConstantComposite: {
++      unsigned elem_count = count - 3;
++      vtn_fail_if(elem_count != val->type->length,
++                  "%s has %u constituents, expected %u",
++                  spirv_op_to_string(opcode), elem_count, val->type->length);
++
++      vtn_assert(b->values[w[2]].value_type == vtn_value_type_ssa);
++      if (!b->values[w[2]].ssa) {
++         b->values[w[2]].value_type = vtn_value_type_invalid; /* Pretend not yet set */
++         vtn_handle_composite(b, SpvOpCompositeConstruct, w, count);
++      }
 +      break;
++   }
 +
 +   case SpvOpSpecConstantOp: {
 +      vtn_assert(val->value_type == vtn_value_type_ssa);
@@ -417,7 +639,7 @@ index 5cb7691506..7620fef9f5 100644
  static bool
  is_glslang(const struct vtn_builder *b)
  {
-@@ -6609,6 +6489,8 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
+@@ -7376,6 +7155,8 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
     /* Skip the SPIR-V header, handled at vtn_create_builder */
     words+= 5;
  
@@ -426,20 +648,20 @@ index 5cb7691506..7620fef9f5 100644
     /* Handle all the preamble instructions */
     words = vtn_foreach_instruction(b, words, word_end,
                                     vtn_handle_preamble_instruction);
-@@ -6690,7 +6572,7 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
+@@ -7461,7 +7242,7 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
+       vtn_foreach_function(func, &b->functions) {
           if ((options->create_library || func->referenced) && !func->emitted) {
-             b->const_table = _mesa_pointer_hash_table_create(b);
- 
+             _mesa_hash_table_clear(b->strings, NULL);
 -            vtn_function_emit(b, func, vtn_handle_body_instruction);
 +            vtn_function_emit(b, func, vtn_handle_spec_constant_instructions, preamble_words, vtn_handle_body_instruction);
              progress = true;
           }
        }
 diff --git a/godot-mesa/src/compiler/spirv/vtn_cfg.c b/godot-mesa/src/compiler/spirv/vtn_cfg.c
-index b02c11b8d7..7d0c647f79 100644
+index 83d3228..78aee17 100644
 --- a/godot-mesa/src/compiler/spirv/vtn_cfg.c
 +++ b/godot-mesa/src/compiler/spirv/vtn_cfg.c
-@@ -1425,6 +1425,8 @@ vtn_emit_cf_func_unstructured(struct vtn_builder *b, struct vtn_function *func,
+@@ -763,6 +763,8 @@ vtn_emit_cf_func_unstructured(struct vtn_builder *b, struct vtn_function *func,
  
  void
  vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
@@ -448,7 +670,7 @@ index b02c11b8d7..7d0c647f79 100644
                    vtn_instruction_handler instruction_handler)
  {
     static int force_unstructured = -1;
-@@ -1440,6 +1442,9 @@ vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
+@@ -777,6 +779,9 @@ vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
     b->nb.exact = b->exact;
     b->phi_table = _mesa_pointer_hash_table_create(b);
  
@@ -459,10 +681,10 @@ index b02c11b8d7..7d0c647f79 100644
        impl->structured = false;
        vtn_emit_cf_func_unstructured(b, func, instruction_handler);
 diff --git a/godot-mesa/src/compiler/spirv/vtn_private.h b/godot-mesa/src/compiler/spirv/vtn_private.h
-index bd65a60d9b..389f477681 100644
+index 5d601f9..4ab1a22 100644
 --- a/godot-mesa/src/compiler/spirv/vtn_private.h
 +++ b/godot-mesa/src/compiler/spirv/vtn_private.h
-@@ -307,6 +307,8 @@ typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
+@@ -243,6 +243,8 @@ typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
  void vtn_build_cfg(struct vtn_builder *b, const uint32_t *words,
                     const uint32_t *end);
  void vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
@@ -471,7 +693,7 @@ index bd65a60d9b..389f477681 100644
                         vtn_instruction_handler instruction_handler);
  void vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
                                const uint32_t *w, unsigned count);
-@@ -563,6 +565,8 @@ struct vtn_variable {
+@@ -521,6 +523,8 @@ struct vtn_variable {
     unsigned descriptor_set;
     unsigned binding;
     bool explicit_binding;
@@ -480,7 +702,7 @@ index bd65a60d9b..389f477681 100644
     unsigned offset;
     unsigned input_attachment_index;
  
-@@ -633,6 +637,9 @@ struct vtn_value {
+@@ -594,6 +598,9 @@ struct vtn_value {
        struct vtn_ssa_value *ssa;
        vtn_instruction_handler ext_handler;
     };
@@ -491,7 +713,7 @@ index bd65a60d9b..389f477681 100644
  
  #define VTN_DEC_DECORATION -1
 diff --git a/godot-mesa/src/compiler/spirv/vtn_variables.c b/godot-mesa/src/compiler/spirv/vtn_variables.c
-index 8db61c8e61..8545903165 100644
+index 55557e7..16866df 100644
 --- a/godot-mesa/src/compiler/spirv/vtn_variables.c
 +++ b/godot-mesa/src/compiler/spirv/vtn_variables.c
 @@ -26,6 +26,8 @@
@@ -503,7 +725,7 @@ index 8db61c8e61..8545903165 100644
  static struct vtn_pointer*
  vtn_align_pointer(struct vtn_builder *b, struct vtn_pointer *ptr,
                    unsigned alignment)
-@@ -1407,13 +1409,17 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
+@@ -1620,13 +1622,17 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
  
     /* Handle decorations that apply to a vtn_variable as a whole */
     switch (dec->decoration) {
@@ -526,10 +748,10 @@ index 8db61c8e61..8545903165 100644
        vtn_var->input_attachment_index = dec->operands[0];
        vtn_var->access |= ACCESS_NON_WRITEABLE;
 diff --git a/godot-mesa/src/microsoft/compiler/dxil_container.c b/godot-mesa/src/microsoft/compiler/dxil_container.c
-index 209973459d..487cda7a40 100644
+index 77a14cd..24d8362 100644
 --- a/godot-mesa/src/microsoft/compiler/dxil_container.c
 +++ b/godot-mesa/src/microsoft/compiler/dxil_container.c
-@@ -338,7 +338,8 @@ dxil_container_add_state_validation(struct dxil_container *c,
+@@ -331,7 +331,8 @@ dxil_container_add_state_validation(struct dxil_container *c,
  
  bool
  dxil_container_add_module(struct dxil_container *c,
@@ -539,7 +761,7 @@ index 209973459d..487cda7a40 100644
  {
     assert(m->buf.buf_bits == 0); // make sure the module is fully flushed
     uint32_t version = (m->shader_kind << 16) |
-@@ -352,18 +353,22 @@ dxil_container_add_module(struct dxil_container *c,
+@@ -345,18 +346,22 @@ dxil_container_add_module(struct dxil_container *c,
     uint32_t bitcode_offset = 16;
     uint32_t bitcode_size = m->buf.blob.size;
  
@@ -566,7 +788,7 @@ index 209973459d..487cda7a40 100644
  {
     assert(blob->size == 0);
     if (!blob_write_bytes(blob, &DXIL_DXBC, sizeof(DXIL_DXBC)))
-@@ -394,8 +399,12 @@ dxil_container_write(struct dxil_container *c, struct blob *blob)
+@@ -387,8 +392,12 @@ dxil_container_write(struct dxil_container *c, struct blob *blob)
     }
  
     if (!blob_write_bytes(blob, &c->num_parts, sizeof(c->num_parts)) ||
@@ -582,7 +804,7 @@ index 209973459d..487cda7a40 100644
  
     return true;
 diff --git a/godot-mesa/src/microsoft/compiler/dxil_container.h b/godot-mesa/src/microsoft/compiler/dxil_container.h
-index b3279ee108..08ab970cb8 100644
+index 2c3f17c..3eb684f 100644
 --- a/godot-mesa/src/microsoft/compiler/dxil_container.h
 +++ b/godot-mesa/src/microsoft/compiler/dxil_container.h
 @@ -123,10 +123,11 @@ dxil_container_add_state_validation(struct dxil_container *c,
@@ -600,7 +822,7 @@ index b3279ee108..08ab970cb8 100644
  #ifdef __cplusplus
  }
 diff --git a/godot-mesa/src/microsoft/compiler/dxil_module.c b/godot-mesa/src/microsoft/compiler/dxil_module.c
-index 773e5640a6..57b6a62ffd 100644
+index c70c67c..0634b52 100644
 --- a/godot-mesa/src/microsoft/compiler/dxil_module.c
 +++ b/godot-mesa/src/microsoft/compiler/dxil_module.c
 @@ -32,6 +32,8 @@
@@ -612,7 +834,7 @@ index 773e5640a6..57b6a62ffd 100644
  void
  dxil_module_init(struct dxil_module *m, void *ralloc_ctx)
  {
-@@ -2630,6 +2632,12 @@ emit_consts(struct dxil_module *m)
+@@ -2669,6 +2671,12 @@ emit_consts(struct dxil_module *m)
           continue;
        }
  
@@ -626,7 +848,7 @@ index 773e5640a6..57b6a62ffd 100644
        case TYPE_INTEGER:
           if (!emit_int_value(m, c->int_value))
 diff --git a/godot-mesa/src/microsoft/compiler/dxil_module.h b/godot-mesa/src/microsoft/compiler/dxil_module.h
-index ca170c5d2f..7fab628a22 100644
+index 08ba263..d4c412d 100644
 --- a/godot-mesa/src/microsoft/compiler/dxil_module.h
 +++ b/godot-mesa/src/microsoft/compiler/dxil_module.h
 @@ -29,6 +29,8 @@
@@ -638,7 +860,7 @@ index ca170c5d2f..7fab628a22 100644
  #ifdef __cplusplus
  extern "C" {
  #endif
-@@ -247,6 +249,8 @@ struct dxil_module {
+@@ -268,6 +270,8 @@ struct dxil_module {
     struct rb_tree *functions;
  
     struct dxil_func_def *cur_emitting_func;
@@ -647,60 +869,11 @@ index ca170c5d2f..7fab628a22 100644
  };
  
  struct dxil_instr;
-diff --git a/godot-mesa/src/microsoft/compiler/dxil_validator.cpp b/godot-mesa/src/microsoft/compiler/dxil_validator.cpp
-index 4b68957a7d..a926ec991b 100644
---- a/godot-mesa/src/microsoft/compiler/dxil_validator.cpp
-+++ b/godot-mesa/src/microsoft/compiler/dxil_validator.cpp
-@@ -25,6 +25,7 @@ struct dxil_validator {
- 
- extern "C" {
- extern IMAGE_DOS_HEADER __ImageBase;
-+extern char godot_nir_arch_name[32];
- }
- 
- static HMODULE
-@@ -36,7 +37,10 @@ load_dxil_mod()
- #elif defined (_GAMING_XBOX)
-    HMODULE mod = LoadLibraryA("dxcompiler_x.dll");
- #else
--   HMODULE mod = LoadLibraryA("DXIL.dll");
-+   HMODULE mod = NULL;
-+   if (!godot_nir_arch_name[0]) {
-+      mod = LoadLibraryA("DXIL.dll");
-+   }
- #endif
-    if (mod)
-       return mod;
-@@ -60,12 +64,23 @@ load_dxil_mod()
-    }
- 
-    *(last_slash + 1) = '\0';
-+
-+   if (godot_nir_arch_name[0]) {
-+      strcat_s(self_path, godot_nir_arch_name);
-+      strcat_s(self_path, "\\");
-+   }
-+
-    if (strcat_s(self_path, "DXIL.dll") != 0) {
-       debug_printf("DXIL: Unable to get path to DXIL.dll next to self");
-       return NULL;
-    }
- 
--   return LoadLibraryA(self_path);
-+   mod = LoadLibraryA(self_path);
-+   if (mod)
-+      return mod;
-+
-+   godot_nir_arch_name[0] = '\0';
-+   return load_dxil_mod();
- }
- 
- static IDxcValidator *
 diff --git a/godot-mesa/src/microsoft/compiler/nir_to_dxil.c b/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
-index 245c5140f9..d6c9d06c30 100644
+index f99cf69..fa5bec7 100644
 --- a/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
 +++ b/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
-@@ -43,6 +43,8 @@
+@@ -44,6 +44,8 @@
  
  #include <stdint.h>
  
@@ -709,7 +882,7 @@ index 245c5140f9..d6c9d06c30 100644
  int debug_dxil = 0;
  
  static const struct debug_named_value
-@@ -1154,6 +1156,8 @@ add_resource(struct ntd_context *ctx, enum dxil_resource_type type,
+@@ -1226,6 +1228,8 @@ add_resource(struct ntd_context *ctx, enum dxil_resource_type type,
        /* No flags supported yet */
        resource_v1->resource_flags = 0;
     }
@@ -718,19 +891,20 @@ index 245c5140f9..d6c9d06c30 100644
  }
  
  static const struct dxil_value *
-@@ -5079,6 +5083,11 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
-    case nir_intrinsic_exclusive_scan:
-       return emit_reduce(ctx, intr);
+@@ -4963,6 +4967,12 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
+       return emit_load_unary_external_function(ctx, intr, "dx.op.startInstanceLocation",
+                                                DXIL_INTR_START_INSTANCE_LOCATION, nir_type_int);
  
-+   case nir_intrinsic_load_constant_non_opt:
++   case nir_intrinsic_load_constant_non_opt: {
 +      const struct dxil_value* value = get_src(ctx, &intr->src[0], 0, nir_type_uint);
-+      store_dest_value(ctx, &intr->dest, 0, value);
++      store_def(ctx, &intr->def, 0, value);
 +      return true;
++   }
 +
     case nir_intrinsic_load_num_workgroups:
     case nir_intrinsic_load_workgroup_size:
     default:
-@@ -6644,6 +6653,7 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
+@@ -6618,6 +6628,7 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
        MIN2(opts->shader_model_max & 0xffff, validator_version & 0xffff);
     ctx->mod.major_validator = validator_version >> 16;
     ctx->mod.minor_validator = validator_version & 0xffff;
@@ -738,7 +912,7 @@ index 245c5140f9..d6c9d06c30 100644
  
     if (s->info.stage <= MESA_SHADER_FRAGMENT) {
        uint64_t in_mask =
-@@ -6762,19 +6772,23 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
+@@ -6773,19 +6784,23 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
        goto out;
     }
  
@@ -765,7 +939,7 @@ index 245c5140f9..d6c9d06c30 100644
        static int shader_id = 0;
        char buffer[64];
 diff --git a/godot-mesa/src/microsoft/compiler/nir_to_dxil.h b/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
-index bdfbe23953..b95ca1d79f 100644
+index 9890115..abe9e9d 100644
 --- a/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
 +++ b/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
 @@ -29,6 +29,8 @@
@@ -785,8 +959,44 @@ index bdfbe23953..b95ca1d79f 100644
  };
  
  typedef void (*dxil_msg_callback)(void *priv, const char *msg);
+diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c b/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
+index 66a996e..4d8eff1 100644
+--- a/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
++++ b/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
+@@ -1058,30 +1058,7 @@ dxil_spirv_nir_passes(nir_shader *nir,
+    NIR_PASS(_, nir, nir_lower_alu_to_scalar, NULL, NULL);
+    NIR_PASS(_, nir, nir_opt_dce);
+    NIR_PASS(_, nir, dxil_nir_lower_double_math);
+-
+-   {
+-      bool progress;
+-      do
+-      {
+-         progress = false;
+-         NIR_PASS(progress, nir, nir_copy_prop);
+-         NIR_PASS(progress, nir, nir_opt_copy_prop_vars);
+-         NIR_PASS(progress, nir, nir_opt_deref);
+-         NIR_PASS(progress, nir, nir_opt_dce);
+-         NIR_PASS(progress, nir, nir_opt_undef);
+-         NIR_PASS(progress, nir, nir_opt_constant_folding);
+-         NIR_PASS(progress, nir, nir_opt_cse);
+-         if (nir_opt_loop(nir)) {
+-            progress = true;
+-            NIR_PASS(progress, nir, nir_copy_prop);
+-            NIR_PASS(progress, nir, nir_opt_dce);
+-         }
+-         NIR_PASS(progress, nir, nir_lower_vars_to_ssa);
+-         NIR_PASS(progress, nir, nir_opt_algebraic);
+-         NIR_PASS(progress, nir, nir_opt_dead_cf);
+-         NIR_PASS(progress, nir, nir_opt_remove_phis);
+-      } while (progress);
+-   }
++   NIR_PASS(_, nir, nir_lower_vars_to_ssa);
+ 
+    NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
+    NIR_PASS(_, nir, nir_split_struct_vars, nir_var_function_temp);
 diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
-index 76cf6b007e..819339adf5 100644
+index 1a8e6e2..1b3b9d3 100644
 --- a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
 +++ b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
 @@ -32,6 +32,8 @@
@@ -795,9 +1005,9 @@ index 76cf6b007e..819339adf5 100644
  
 +#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
 +
- static_assert(DXIL_SPIRV_SHADER_NONE == (int)MESA_SHADER_NONE, "must match");
- static_assert(DXIL_SPIRV_SHADER_VERTEX == (int)MESA_SHADER_VERTEX, "must match");
- static_assert(DXIL_SPIRV_SHADER_TESS_CTRL == (int)MESA_SHADER_TESS_CTRL, "must match");
+ static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_NONE == MESA_SHADER_NONE, "must match");
+ static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_VERTEX == MESA_SHADER_VERTEX, "must match");
+ static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_TESS_CTRL == MESA_SHADER_TESS_CTRL, "must match");
 @@ -50,6 +52,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
                const struct dxil_spirv_debug_options *dgb_opts,
                const struct dxil_spirv_runtime_conf *conf,
@@ -815,7 +1025,7 @@ index 76cf6b007e..819339adf5 100644
  
     const struct spirv_to_nir_options *spirv_opts = dxil_spirv_nir_get_spirv_options();
 diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
-index 40adf761ef..30efbd16dd 100644
+index 00b1884..0d10644 100644
 --- a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
 +++ b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
 @@ -30,6 +30,8 @@
@@ -827,7 +1037,7 @@ index 40adf761ef..30efbd16dd 100644
  #ifdef __cplusplus
  extern "C" {
  #endif
-@@ -215,6 +217,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
+@@ -234,6 +236,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
                const struct dxil_spirv_debug_options *debug_options,
                const struct dxil_spirv_runtime_conf *conf,
                const struct dxil_spirv_logger *logger,
@@ -835,3 +1045,81 @@ index 40adf761ef..30efbd16dd 100644
                struct dxil_spirv_object *out_dxil);
  
  /**
+diff --git a/godot-mesa/src/util/macros.h b/godot-mesa/src/util/macros.h
+index 0070067..72410ef 100644
+--- a/godot-mesa/src/util/macros.h
++++ b/godot-mesa/src/util/macros.h
+@@ -543,14 +543,14 @@ typedef int lock_cap_t;
+  * SWAP - swap value of @a and @b
+  */
+ #if !defined(_MSC_VER) || _MSC_VER >= 1939 /* MSVC 17.9 or later for __typeof__ */
+-#define SWAP(a, b)                                                             \
++#define MESA_SWAP(a, b)                                                        \
+    do {                                                                        \
+       __typeof__(a) __tmp = (a);                                               \
+       (a) = (b);                                                               \
+       (b) = __tmp;                                                             \
+    } while (0)
+ #else
+-#define SWAP(a, b)                                                             \
++#define MESA_SWAP(a, b)                                                        \
+    do {                                                                        \
+       /* NOLINTBEGIN(bugprone-sizeof-expression) */                            \
+       char __tmp[sizeof(a) == sizeof(b) ? (ptrdiff_t)sizeof(a) : -1];          \
+diff --git a/godot-mesa/src/util/ralloc.c b/godot-mesa/src/util/ralloc.c
+index ba560c8..07379ae 100644
+--- a/godot-mesa/src/util/ralloc.c
++++ b/godot-mesa/src/util/ralloc.c
+@@ -35,6 +35,8 @@
+ 
+ #include "ralloc.h"
+ 
++#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
++
+ #define CANARY 0x5A1106
+ 
+ #if defined(__LP64__) || defined(_WIN64)
+@@ -115,8 +117,8 @@ ralloc_size(const void *ctx, size_t size)
+     *  - Allocations of a size that rounds up to a multiple of 8 bytes and
+     *    not 16 bytes, are only required to have at least 8 byte alignment.
+     */
+-   void *block = malloc(align64(size + sizeof(ralloc_header),
+-                                alignof(ralloc_header)));
++   void *block = godot_nir_malloc(align64(size + sizeof(ralloc_header),
++                                          alignof(ralloc_header)));
+    ralloc_header *info;
+    ralloc_header *parent;
+ 
+@@ -164,8 +166,8 @@ resize(void *ptr, size_t size)
+    ralloc_header *child, *old, *info;
+ 
+    old = get_header(ptr);
+-   info = realloc(old, align64(size + sizeof(ralloc_header),
+-                               alignof(ralloc_header)));
++   info = godot_nir_realloc(old, align64(size + sizeof(ralloc_header),
++                                         alignof(ralloc_header)));
+ 
+    if (info == NULL)
+       return NULL;
+@@ -323,7 +325,7 @@ unsafe_free(ralloc_header *info)
+    if (info->destructor != NULL)
+       info->destructor(PTR_FROM_HEADER(info));
+ 
+-   free(info);
++   godot_nir_free(info);
+ }
+ 
+ void
+diff --git a/godot-mesa/src/util/set.c b/godot-mesa/src/util/set.c
+index d6c0623..1eac1e8 100644
+--- a/godot-mesa/src/util/set.c
++++ b/godot-mesa/src/util/set.c
+@@ -691,7 +691,7 @@ _mesa_set_intersects(struct set *a, struct set *b)
+ 
+    /* iterate over the set with less entries */
+    if (b->entries < a->entries) {
+-      SWAP(a, b);
++      MESA_SWAP(a, b);
+    }
+ 
+    set_foreach(a, entry) {

godot-tools/macos.py

@@ -1,4 +1,3 @@
-import os
 import sys
 import macos_osxcross
 

godot-tools/targets.py

@@ -1,7 +1,5 @@
-import os
-import sys
 from SCons.Script import ARGUMENTS
-from SCons.Variables import *
+from SCons.Variables import BoolVariable, EnumVariable
 from SCons.Variables.BoolVariable import _text2bool
 
 

godot-tools/windows.py

@@ -1,17 +1,21 @@
+import os
 import sys
 
 import my_spawn
 
 from SCons.Tool import msvc, mingw
-from SCons.Variables import *
+from SCons.Variables import BoolVariable
 
 
 def options(opts):
+    mingw = os.getenv("MINGW_PREFIX", "")
+
     opts.Add(BoolVariable("use_mingw", "Use the MinGW compiler instead of MSVC - only effective on Windows", False))
     opts.Add(BoolVariable("use_clang_cl", "Use the clang driver instead of MSVC - only effective on Windows", False))
     opts.Add(BoolVariable("use_static_cpp", "Link MinGW/MSVC C++ runtime libraries statically", True))
     opts.Add(BoolVariable("debug_crt", "Compile with MSVC's debug CRT (/MDd)", False))
     opts.Add(BoolVariable("use_llvm", "Use the LLVM compiler", False))
+    opts.Add("mingw_prefix", "MinGW prefix", mingw)
 
 
 def exists(env):
@@ -23,12 +27,22 @@ def generate(env):
     if not env["use_mingw"] and msvc.exists(env):
         if env["arch"] == "x86_64":
             env["TARGET_ARCH"] = "amd64"
+        elif env["arch"] == "arm64":
+            env["TARGET_ARCH"] = "arm64"
+        elif env["arch"] == "arm32":
+            env["TARGET_ARCH"] = "arm"
         elif env["arch"] == "x86_32":
             env["TARGET_ARCH"] = "x86"
+
+        env["MSVC_SETUP_RUN"] = False  # Need to set this to re-run the tool
+        env["MSVS_VERSION"] = None
+        env["MSVC_VERSION"] = None
+
         env["is_msvc"] = True
 
         # MSVC, linker, and archiver.
         msvc.generate(env)
+        env.Tool("msvc")
         env.Tool("mslib")
         env.Tool("mslink")
 
@@ -48,7 +62,7 @@ def generate(env):
             env["CC"] = "clang-cl"
             env["CXX"] = "clang-cl"
 
-    elif sys.platform == "win32" or sys.platform == "msys":
+    elif (sys.platform == "win32" or sys.platform == "msys") and not env["mingw_prefix"]:
         env["use_mingw"] = True
         mingw.generate(env)
         env.Append(CPPDEFINES=["MINGW_ENABLED"])
@@ -63,17 +77,29 @@ def generate(env):
     else:
         env["use_mingw"] = True
         # Cross-compilation using MinGW
-        prefix = "i686" if env["arch"] == "x86_32" else env["arch"]
+        prefix = ""
+        if env["mingw_prefix"]:
+            prefix = env["mingw_prefix"] + "/bin/"
+
+        if env["arch"] == "x86_64":
+            prefix += "x86_64"
+        elif env["arch"] == "arm64":
+            prefix += "aarch64"
+        elif env["arch"] == "arm32":
+            prefix += "armv7"
+        elif env["arch"] == "x86_32":
+            prefix += "i686"
+
         if env["use_llvm"]:
-            env["CXX"] = prefix + "-w64-mingw32-clang"
-            env["CC"] = prefix + "-w64-mingw32-clang++"
+            env["CXX"] = prefix + "-w64-mingw32-clang++"
+            env["CC"] = prefix + "-w64-mingw32-clang"
             env["AR"] = prefix + "-w64-mingw32-ar"
             env["RANLIB"] = prefix + "-w64-mingw32-ranlib"
             env["LINK"] = prefix + "-w64-mingw32-clang"
         else:
             env["CXX"] = prefix + "-w64-mingw32-g++"
             env["CC"] = prefix + "-w64-mingw32-gcc"
-            env["AR"] = prefix + "-w64-mingw32-ar"
+            env["AR"] = prefix + "-w64-mingw32-gcc-ar"
             env["RANLIB"] = prefix + "-w64-mingw32-ranlib"
             env["LINK"] = prefix + "-w64-mingw32-g++"
 
@@ -92,3 +118,6 @@ def generate(env):
                 "-Wl,--no-undefined",
             ]
         )
+
+        if (sys.platform == "win32" or sys.platform == "msys"):
+            my_spawn.configure(env)

update_mesa.sh

@@ -36,7 +36,7 @@ run_custom_steps_at_source() {
     }
 
     run_step bin 'git_sha1_gen.py --output $OUTDIR/git_sha1.h'
-    run_step src/compiler/spirv 'spirv_info_c.py spirv.core.grammar.json $OUTDIR/spirv_info.c'
+    run_step src/compiler/spirv 'spirv_info_gen.py --json spirv.core.grammar.json --out-h $OUTDIR/spirv_info.h --out-c $OUTDIR/spirv_info.c'
     run_step src/compiler/spirv 'vtn_gather_types_c.py spirv.core.grammar.json $OUTDIR/vtn_gather_types.c'
 }
 
@@ -49,6 +49,9 @@ copy_file() {
 }
 
 copy_custom_steps_sources() {
+    copy_file src/compiler builtin_types.py
+    copy_file src/compiler builtin_types_h.py
+    copy_file src/compiler builtin_types_c.py
     copy_file src/compiler/glsl ir_expression_operation.py
     copy_file src/compiler/nir nir_builder_opcodes_h.py
     copy_file src/compiler/nir nir_constant_expressions.py
@@ -65,7 +68,7 @@ copy_custom_steps_sources() {
     copy_file src/compiler/spirv vtn_generator_ids_h.py
     copy_file src/microsoft/compiler dxil_nir_algebraic.py
     copy_file src/util format_srgb.py
-    copy_file src/util/format u_format.csv
+    copy_file src/util/format u_format.yaml
     copy_file src/util/format u_format_pack.py
     copy_file src/util/format u_format_parse.py
     copy_file src/util/format u_format_table.py
@@ -89,7 +92,6 @@ copy_sources() {
     }
 
     # These are the first we know for sure we want to copy.
-    copy_file . VERSION
     copy_file . .editorconfig
     copy_subir_sources src/microsoft/compiler
     copy_subir_sources src/microsoft/spirv_to_dxil
@@ -106,6 +108,7 @@ copy_sources() {
     copy_subir_headers include/KHR
     copy_subir_headers src/c11
     copy_file src/c11/impl threads_win32.*
+    copy_file src/c11/impl time.c
     copy_subir_sources src/compiler
     copy_subir_headers src/compiler/glsl
     copy_subir_sources src/compiler/nir
@@ -126,34 +129,52 @@ copy_sources() {
     copy_file src/util blob.c
     copy_file src/util bitscan.c
     copy_file src/util double.c
+    copy_file src/util float8.c
     copy_file src/util half_float.c
     copy_file src/util hash_table.c
     copy_file src/util log.c
+    copy_file src/util mesa-blake3.c
     copy_file src/util mesa-sha1.c
     copy_file src/util memstream.c
     copy_file src/util os_misc.c
     copy_file src/util ralloc.c
+    copy_file src/util range_minimum_query.c
     copy_file src/util rb_tree.c
     copy_file src/util rgtc.c
     copy_file src/util set.c
     copy_file src/util simple_mtx.c
     copy_file src/util softfloat.c
     copy_file src/util string_buffer.c
+    copy_file src/util strndup.c
     copy_file src/util u_call_once.c
+    copy_file src/util u_cpu_detect.c
     copy_file src/util u_debug.c
+    copy_file src/util u_dynarray.c
     copy_file src/util u_printf.c
     copy_file src/util u_qsort.cpp
+    copy_file src/util u_thread.c
     copy_file src/util u_vector.c
     copy_file src/util u_worklist.c
+    copy_subir_headers src/util/blake3
+    copy_file src/util/blake3 blake3.c
+    copy_file src/util/blake3 blake3_dispatch.c
+    copy_file src/util/blake3 blake3_portable.c
+    copy_subir_headers src/util/perf
 
-    cp ./mesa/VERSION godot-mesa
+    cp ./mesa/VERSION godot-mesa/VERSION.info
     check_error
 }
 
 blacklist_sources() {
+    rm godot-mesa/src/compiler/nir/nir_stub.c
+    check_error
     # These are programs. Not needed and makes build hungrier for dependencies.
     rm godot-mesa/src/compiler/spirv/spirv2nir.c
     check_error
+    rm godot-mesa/src/compiler/spirv/vtn_bindgen2.c
+    check_error
+    rm godot-mesa/src/microsoft/compiler/dxil_buffer_test.c
+    check_error
     rm godot-mesa/src/microsoft/spirv_to_dxil/spirv2dxil.c
     check_error
 }
@@ -195,17 +216,20 @@ custom_source_gen() {
     run_step 'src/compiler' 'glsl/ir_expression_operation.py enum' 'ir_expression_operation.h'
     run_step 'src/compiler/nir' 'nir_builder_opcodes_h.py' 'nir_builder_opcodes.h'
     run_step 'src/compiler/nir' 'nir_constant_expressions.py' 'nir_constant_expressions.c'
-    run_step 'src/compiler/nir' 'nir_intrinsics_h.py --outdir $GENDIR'
-    run_step 'src/compiler/nir' 'nir_intrinsics_c.py --outdir $GENDIR'
-    run_step 'src/compiler/nir' 'nir_intrinsics_indices_h.py --outdir $GENDIR'
+    run_step 'src/compiler/nir' 'nir_intrinsics_h.py --out $GENDIR/nir_intrinsics.h'
+    run_step 'src/compiler/nir' 'nir_intrinsics_c.py --out $GENDIR/nir_intrinsics.c'
+    run_step 'src/compiler/nir' 'nir_intrinsics_indices_h.py --out $GENDIR/nir_intrinsics_indices.h'
     run_step 'src/compiler/nir' 'nir_opcodes_h.py' 'nir_opcodes.h'
     run_step 'src/compiler/nir' 'nir_opcodes_c.py' 'nir_opcodes.c'
-    run_step 'src/compiler/nir' 'nir_opt_algebraic.py' 'nir_opt_algebraic.c'
+    run_step 'src/compiler/nir' 'nir_opt_algebraic.py --out $GENDIR/nir_opt_algebraic.c'
     run_step 'src/compiler/spirv' 'vtn_generator_ids_h.py spir-v.xml $GENDIR/vtn_generator_ids.h'
     run_step 'src/microsoft/compiler' 'dxil_nir_algebraic.py -p ../../../src/compiler/nir' 'dxil_nir_algebraic.c'
     run_step 'src/util' 'format_srgb.py' 'format_srgb.c'
-    run_step 'src/util/format' 'u_format_table.py u_format.csv --header' 'u_format_pack.h'
-    run_step 'src/util/format' 'u_format_table.py u_format.csv' 'u_format_table.c'
+    run_step 'src/util/format' 'u_format_table.py u_format.yaml --enums' 'u_format_gen.h'
+    run_step 'src/util/format' 'u_format_table.py u_format.yaml --header' 'u_format_pack.h'
+    run_step 'src/util/format' 'u_format_table.py u_format.yaml' 'u_format_table.c'
+    run_step 'src/compiler' 'builtin_types_h.py $GENDIR/builtin_types.h'
+    run_step 'src/compiler' 'builtin_types_c.py $GENDIR/builtin_types.c'
 }