File Mesa-MR27783_drop-nir-in-clover.patch of Package Mesa

Overview Repositories Revisions Requests Users Attributes Meta

File Mesa-MR27783_drop-nir-in-clover.patch of Package Mesa (Revision 03d9a31ac43184a49cbd3b74c5e56ac7)

Currently displaying revision 03d9a31ac43184a49cbd3b74c5e56ac7 , Show latest

From 481ca6704379413b0bfee333ed96c71fa7928c64 Mon Sep 17 00:00:00 2001
From: Karol Herbst <kherbst@redhat.com>
Date: Mon, 26 Feb 2024 15:51:46 +0100
Subject: [PATCH 1/4] nvc0: return NULL instead of asserting in
 nvc0_resource_from_user_memory

Fixes: 212f1ab40ea ("nvc0: support PIPE_CAP_RESOURCE_FROM_USER_MEMORY_COMPUTE_ONLY")
Signed-off-by: Karol Herbst <kherbst@redhat.com>
---
 src/gallium/drivers/nouveau/nvc0/nvc0_resource.c | 8 +++-----
 1 file changed, 3 insertions(+), 5 deletions(-)

diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_resource.c b/src/gallium/drivers/nouveau/nvc0/nvc0_resource.c
index c51990a557601..385cfdd19c07d 100644
--- a/src/gallium/drivers/nouveau/nvc0/nvc0_resource.c
+++ b/src/gallium/drivers/nouveau/nvc0/nvc0_resource.c
@@ -141,11 +141,9 @@ nvc0_resource_from_user_memory(struct pipe_screen *pipe,
                                const struct pipe_resource *templ,
                                void *user_memory)
 {
-   ASSERTED struct nouveau_screen *screen = nouveau_screen(pipe);
-
-   assert(screen->has_svm);
-   assert(templ->target == PIPE_BUFFER);
-
+   struct nouveau_screen *screen = nouveau_screen(pipe);
+   if (!screen->has_svm || templ->target != PIPE_BUFFER)
+      return NULL;
    return nouveau_buffer_create_from_user(pipe, templ, user_memory);
 }
 
-- 
GitLab

From a976f2c9f0c07f9e06cc9affd9124b45bc60c2bd Mon Sep 17 00:00:00 2001
From: Karol Herbst <kherbst@redhat.com>
Date: Sun, 25 Feb 2024 19:25:03 +0100
Subject: [PATCH 2/4] clover: drop support for nir drivers

People had enough time to migrate to rusticl, also nobody would support
this anyway anymore.

Fixes: 815a6647eb1 ("meson: do not pull in clc for clover")
Signed-off-by: Karol Herbst <kherbst@redhat.com>
---
 .gitlab-ci/build/gitlab-ci.yml                |    2 -
 .gitlab-ci/windows/mesa_build.ps1             |    1 -
 meson.build                                   |    3 +-
 meson_options.txt                             |    8 -
 src/gallium/frontends/clover/api/program.cpp  |   18 +-
 .../frontends/clover/core/compiler.hpp        |   25 -
 src/gallium/frontends/clover/core/device.cpp  |   32 +-
 .../frontends/clover/llvm/invocation.cpp      |   76 --
 .../frontends/clover/llvm/invocation.hpp      |    8 -
 src/gallium/frontends/clover/meson.build      |   35 +-
 .../frontends/clover/nir/invocation.cpp       |  472 --------
 .../frontends/clover/nir/invocation.hpp       |   46 -
 .../frontends/clover/spirv/invocation.cpp     | 1031 -----------------
 .../frontends/clover/spirv/invocation.hpp     |   81 --
 14 files changed, 6 insertions(+), 1832 deletions(-)
 delete mode 100644 src/gallium/frontends/clover/nir/invocation.cpp
 delete mode 100644 src/gallium/frontends/clover/nir/invocation.hpp
 delete mode 100644 src/gallium/frontends/clover/spirv/invocation.cpp
 delete mode 100644 src/gallium/frontends/clover/spirv/invocation.hpp

diff --git a/.gitlab-ci/build/gitlab-ci.yml b/.gitlab-ci/build/gitlab-ci.yml
index 2a12acb39075b..ed249d8d8edfb 100644
--- a/.gitlab-ci/build/gitlab-ci.yml
+++ b/.gitlab-ci/build/gitlab-ci.yml
@@ -536,7 +536,6 @@ debian-clang:
       -D llvm=enabled
       -D microsoft-clc=disabled
       -D shared-llvm=enabled
-      -D opencl-spirv=true
       -D shared-glapi=enabled
     GALLIUM_DRIVERS: "iris,nouveau,kmsro,r300,r600,freedreno,swrast,svga,v3d,vc4,virgl,etnaviv,panfrost,lima,zink,radeonsi,tegra,d3d12,crocus,i915,asahi"
     VULKAN_DRIVERS: intel,amd,freedreno,broadcom,virtio,swrast,panfrost,imagination-experimental,microsoft-experimental,nouveau-experimental
@@ -574,7 +573,6 @@ debian-clang-release:
       -D llvm=enabled
       -D microsoft-clc=disabled
       -D shared-llvm=enabled
-      -D opencl-spirv=true
       -D shared-glapi=disabled
 
 windows-msvc:
diff --git a/.gitlab-ci/windows/mesa_build.ps1 b/.gitlab-ci/windows/mesa_build.ps1
index 9b768e6281835..01625dd7fdd6d 100644
--- a/.gitlab-ci/windows/mesa_build.ps1
+++ b/.gitlab-ci/windows/mesa_build.ps1
@@ -56,7 +56,6 @@ meson setup `
 -Dgles2=enabled `
 -Dgallium-opencl=icd `
 -Dgallium-rusticl=false `
--Dopencl-spirv=true `
 -Dmicrosoft-clc=enabled `
 -Dstatic-libclc=all `
 -Dspirv-to-dxil=true `
diff --git a/meson.build b/meson.build
index 90d2f3644facb..f38b9e14d3504 100644
--- a/meson.build
+++ b/meson.build
@@ -1824,8 +1824,7 @@ endif
 pre_args += '-DLLVM_AVAILABLE=' + (with_llvm ? '1' : '0')
 pre_args += '-DDRAW_LLVM_AVAILABLE=' + (with_llvm and draw_with_llvm ? '1' : '0')
 
-with_opencl_spirv = (_opencl != 'disabled' and get_option('opencl-spirv')) or with_clc
-if with_opencl_spirv
+if with_clc
   chosen_llvm_version_array = dep_llvm.version().split('.')
   chosen_llvm_version_major = chosen_llvm_version_array[0].to_int()
   chosen_llvm_version_minor = chosen_llvm_version_array[1].to_int()
diff --git a/meson_options.txt b/meson_options.txt
index 3da8e6a3350bd..7efe72f86ab35 100644
--- a/meson_options.txt
+++ b/meson_options.txt
@@ -192,14 +192,6 @@ option(
                 'defaults to libgallium_d3d10.dll to match DRI',
 )
 
-option(
-  'opencl-spirv',
-  type : 'boolean',
-  value : false,
-  description : 'build gallium "clover" OpenCL frontend with SPIR-V ' +
-                'binary support.',
-)
-
 option(
   'static-libclc',
   type : 'array',
diff --git a/src/gallium/frontends/clover/api/program.cpp b/src/gallium/frontends/clover/api/program.cpp
index bd7daa6dc0caa..b01ce08d01051 100644
--- a/src/gallium/frontends/clover/api/program.cpp
+++ b/src/gallium/frontends/clover/api/program.cpp
@@ -23,7 +23,6 @@
 #include "api/util.hpp"
 #include "core/program.hpp"
 #include "core/platform.hpp"
-#include "spirv/invocation.hpp"
 #include "util/u_debug.h"
 
 #include <limits>
@@ -80,22 +79,7 @@ namespace {
                             const cl_version opencl_version,
                             const context::notify_action &notify) {
 
-      enum program::il_type il_type = program::il_type::none;
-
-#ifdef HAVE_CLOVER_SPIRV
-      if (spirv::is_binary_spirv(il)) {
-         std::string log;
-         if (!spirv::is_valid_spirv(il, opencl_version, log)) {
-            if (notify) {
-               notify(log.c_str());
-            }
-            throw error(CL_INVALID_VALUE);
-         }
-         il_type = program::il_type::spirv;
-      }
-#endif
-
-      return il_type;
+      return program::il_type::none;
    }
 }
 
diff --git a/src/gallium/frontends/clover/core/compiler.hpp b/src/gallium/frontends/clover/core/compiler.hpp
index 265b280f356cc..c598a9ea879db 100644
--- a/src/gallium/frontends/clover/core/compiler.hpp
+++ b/src/gallium/frontends/clover/core/compiler.hpp
@@ -26,8 +26,6 @@
 #include "core/device.hpp"
 #include "core/binary.hpp"
 #include "llvm/invocation.hpp"
-#include "nir/invocation.hpp"
-#include "spirv/invocation.hpp"
 
 namespace clover {
    namespace compiler {
@@ -36,18 +34,6 @@ namespace clover {
                       const device &dev, const std::string &opts,
                       std::string &log) {
          switch (dev.ir_format()) {
-#ifdef HAVE_CLOVER_SPIRV
-         case PIPE_SHADER_IR_NIR_SERIALIZED:
-            switch (prog.il_type()) {
-            case program::il_type::source:
-               return llvm::compile_to_spirv(prog.source(), headers, dev, opts, log);
-            case program::il_type::spirv:
-               return spirv::compile_program(prog.source(), dev, log);
-            default:
-               unreachable("device with unsupported IL");
-               throw error(CL_INVALID_VALUE);
-            }
-#endif
          case PIPE_SHADER_IR_NATIVE:
             if (prog.il_type() == program::il_type::source)
                return llvm::compile_program(prog.source(), headers, dev, opts, log);
@@ -63,17 +49,6 @@ namespace clover {
       link_program(const std::vector<binary> &bs, const device &dev,
                    const std::string &opts, std::string &log) {
          switch (dev.ir_format()) {
-#ifdef HAVE_CLOVER_SPIRV
-         case PIPE_SHADER_IR_NIR_SERIALIZED: {
-            const bool create_library =
-               opts.find("-create-library") != std::string::npos;
-            auto spirv_linked_module = spirv::link_program(bs, dev, opts, log);
-            if (create_library)
-               return spirv_linked_module;
-            return nir::spirv_to_nir(spirv_linked_module,
-                                     dev, log);
-         }
-#endif
          case PIPE_SHADER_IR_NATIVE:
             return llvm::link_program(bs, dev, opts, log);
          default:
diff --git a/src/gallium/frontends/clover/core/device.cpp b/src/gallium/frontends/clover/core/device.cpp
index 848c19b7660a1..dd4b5ebb7bf1d 100644
--- a/src/gallium/frontends/clover/core/device.cpp
+++ b/src/gallium/frontends/clover/core/device.cpp
@@ -31,11 +31,6 @@
 #include "nir.h"
 #include <fstream>
 
-#ifdef HAVE_CLOVER_SPIRV
-#include "spirv/invocation.hpp"
-#include "nir/invocation.hpp"
-#endif
-
 using namespace clover;
 
 namespace {
@@ -165,8 +160,7 @@ device::device(clover::platform &platform, pipe_loader_device *ldev) :
    platform(platform), clc_cache(NULL), ldev(ldev) {
    pipe = pipe_loader_create_screen(ldev);
    if (pipe && pipe->get_param(pipe, PIPE_CAP_COMPUTE)) {
-      const bool has_supported_ir = supports_ir(PIPE_SHADER_IR_NATIVE) ||
-                                    supports_ir(PIPE_SHADER_IR_NIR_SERIALIZED);
+      const bool has_supported_ir = supports_ir(PIPE_SHADER_IR_NATIVE);
       if (has_supported_ir) {
          unsigned major = 1, minor = 1;
          debug_get_version_option("CLOVER_DEVICE_CLC_VERSION_OVERRIDE",
@@ -184,14 +178,6 @@ device::device(clover::platform &platform, pipe_loader_device *ldev) :
 
       if (supports_ir(PIPE_SHADER_IR_NATIVE))
          return;
-#ifdef HAVE_CLOVER_SPIRV
-      if (supports_ir(PIPE_SHADER_IR_NIR_SERIALIZED)) {
-         nir::check_for_libclc(*this);
-         clc_cache = nir::create_clc_disk_cache();
-         clc_nir = lazy<std::shared_ptr<nir_shader>>([&] () { std::string log; return std::shared_ptr<nir_shader>(nir::load_libclc_nir(*this, log), ralloc_free); });
-         return;
-      }
-#endif
    }
    if (pipe)
       pipe->destroy(pipe);
@@ -460,11 +446,8 @@ device::vendor_name() const {
 
 enum pipe_shader_ir
 device::ir_format() const {
-   if (supports_ir(PIPE_SHADER_IR_NATIVE))
-      return PIPE_SHADER_IR_NATIVE;
-
-   assert(supports_ir(PIPE_SHADER_IR_NIR_SERIALIZED));
-   return PIPE_SHADER_IR_NIR_SERIALIZED;
+   assert(supports_ir(PIPE_SHADER_IR_NATIVE));
+   return PIPE_SHADER_IR_NATIVE;
 }
 
 std::string
@@ -528,11 +511,6 @@ device::supported_extensions() const {
       vec.push_back( cl_name_version{ CL_MAKE_VERSION(1, 0, 0), "cl_khr_fp16" } );
    if (svm_support())
       vec.push_back( cl_name_version{ CL_MAKE_VERSION(1, 0, 0), "cl_arm_shared_virtual_memory" } );
-#ifdef HAVE_CLOVER_SPIRV
-   if (!clover::spirv::supported_versions().empty() &&
-       supports_ir(PIPE_SHADER_IR_NIR_SERIALIZED))
-      vec.push_back( cl_name_version{ CL_MAKE_VERSION(1, 0, 0), "cl_khr_il_program" } );
-#endif
    vec.push_back( cl_name_version{ CL_MAKE_VERSION(1, 0, 0), "cl_khr_extended_versioning" } );
    return vec;
 }
@@ -555,11 +533,7 @@ device::supported_extensions_as_string() const {
 
 std::vector<cl_name_version>
 device::supported_il_versions() const {
-#ifdef HAVE_CLOVER_SPIRV
-   return clover::spirv::supported_versions();
-#else
    return {};
-#endif
 }
 
 const void *
diff --git a/src/gallium/frontends/clover/llvm/invocation.cpp b/src/gallium/frontends/clover/llvm/invocation.cpp
index 6ab32befbcd38..d1a1fd570e22e 100644
--- a/src/gallium/frontends/clover/llvm/invocation.cpp
+++ b/src/gallium/frontends/clover/llvm/invocation.cpp
@@ -31,10 +31,6 @@
 #include <llvm/Support/raw_ostream.h>
 #include <llvm/Transforms/IPO/Internalize.h>
 #include <llvm-c/Target.h>
-#ifdef HAVE_CLOVER_SPIRV
-#include <LLVMSPIRVLib/LLVMSPIRVLib.h>
-#endif
-
 #include <llvm-c/TargetMachine.h>
 #include <llvm-c/Transforms/PassBuilder.h>
 #include <llvm/Support/CBindingWrapping.h>
@@ -58,9 +54,6 @@
 #include "llvm/invocation.hpp"
 #include "llvm/metadata.hpp"
 #include "llvm/util.hpp"
-#ifdef HAVE_CLOVER_SPIRV
-#include "spirv/invocation.hpp"
-#endif
 #include "util/algorithm.hpp"
 
 
@@ -398,30 +391,6 @@ namespace {
 
       return act.takeModule();
    }
-
-#ifdef HAVE_CLOVER_SPIRV
-   SPIRV::TranslatorOpts
-   get_spirv_translator_options(const device &dev) {
-      const auto supported_versions = clover::spirv::supported_versions();
-      const auto max_supported = clover::spirv::to_spirv_version_encoding(supported_versions.back().version);
-      const auto maximum_spirv_version =
-         std::min(static_cast<SPIRV::VersionNumber>(max_supported),
-                  SPIRV::VersionNumber::MaximumVersion);
-
-      SPIRV::TranslatorOpts::ExtensionsStatusMap spirv_extensions;
-      for (auto &ext : clover::spirv::supported_extensions()) {
-         #define EXT(X) if (ext == #X) spirv_extensions.insert({ SPIRV::ExtensionID::X, true });
-         #include <LLVMSPIRVLib/LLVMSPIRVExtensions.inc>
-         #undef EXT
-      }
-
-      auto translator_opts = SPIRV::TranslatorOpts(maximum_spirv_version, spirv_extensions);
-#if LLVM_VERSION_MAJOR >= 13
-      translator_opts.setPreserveOCLKernelArgTypeMetadataThroughString(true);
-#endif
-      return translator_opts;
-   }
-#endif
 }
 
 binary
@@ -564,48 +533,3 @@ clover::llvm::link_program(const std::vector<binary> &binaries,
       unreachable("Unsupported IR.");
    }
 }
-
-#ifdef HAVE_CLOVER_SPIRV
-binary
-clover::llvm::compile_to_spirv(const std::string &source,
-                               const header_map &headers,
-                               const device &dev,
-                               const std::string &opts,
-                               std::string &r_log) {
-   if (has_flag(debug::clc))
-      debug::log(".cl", "// Options: " + opts + '\n' + source);
-
-   auto ctx = create_context(r_log);
-   const std::string target = dev.address_bits() == 32u ?
-      "-spir-unknown-unknown" :
-      "-spir64-unknown-unknown";
-   auto c = create_compiler_instance(dev, target,
-                                     tokenize(opts + " -O0 -fgnu89-inline input.cl"), r_log);
-   auto mod = compile(*ctx, *c, "input.cl", source, headers, dev, opts, false,
-                      r_log);
-
-   if (has_flag(debug::llvm))
-      debug::log(".ll", print_module_bitcode(*mod));
-
-   const auto spirv_options = get_spirv_translator_options(dev);
-
-   std::string error_msg;
-   std::ostringstream os;
-   if (!::llvm::writeSpirv(mod.get(), spirv_options, os, error_msg)) {
-      r_log += "Translation from LLVM IR to SPIR-V failed: " + error_msg + ".\n";
-      throw error(CL_INVALID_VALUE);
-   }
-
-   const std::string osContent = os.str();
-   std::string binary(osContent.begin(), osContent.end());
-   if (binary.empty()) {
-      r_log += "Failed to retrieve SPIR-V binary.\n";
-      throw error(CL_INVALID_VALUE);
-   }
-
-   if (has_flag(debug::spirv))
-      debug::log(".spvasm", spirv::print_module(binary, dev.device_version()));
-
-   return spirv::compile_program(binary, dev, r_log);
-}
-#endif
diff --git a/src/gallium/frontends/clover/llvm/invocation.hpp b/src/gallium/frontends/clover/llvm/invocation.hpp
index f6304978f6b33..3ef51c77175eb 100644
--- a/src/gallium/frontends/clover/llvm/invocation.hpp
+++ b/src/gallium/frontends/clover/llvm/invocation.hpp
@@ -40,14 +40,6 @@ namespace clover {
                           const device &device,
                           const std::string &opts,
                           std::string &r_log);
-
-#ifdef HAVE_CLOVER_SPIRV
-      binary compile_to_spirv(const std::string &source,
-                              const header_map &headers,
-                              const device &dev,
-                              const std::string &opts,
-                              std::string &r_log);
-#endif
    }
 }
 
diff --git a/src/gallium/frontends/clover/meson.build b/src/gallium/frontends/clover/meson.build
index 0dffda0d4b4a3..c43d5a049bcec 100644
--- a/src/gallium/frontends/clover/meson.build
+++ b/src/gallium/frontends/clover/meson.build
@@ -30,7 +30,6 @@ clover_opencl_cpp_args = [
   '-DLIBCLC_INCLUDEDIR="@0@/"'.format(dep_clc.get_variable(pkgconfig : 'includedir')),
   '-DLIBCLC_LIBEXECDIR="@0@/"'.format(dep_clc.get_variable(pkgconfig : 'libexecdir'))
 ]
-clover_spirv_cpp_args = []
 clover_incs = [inc_include, inc_src, inc_gallium, inc_gallium_aux]
 
 # the CL header files declare attributes on the CL types. Compilers warn if
@@ -44,10 +43,6 @@ if with_opencl_icd
   clover_cpp_args += '-DHAVE_CLOVER_ICD'
 endif
 
-if with_opencl_spirv
-  clover_spirv_cpp_args += '-DHAVE_CLOVER_SPIRV'
-endif
-
 libclllvm = static_library(
   'clllvm',
   files(
@@ -65,7 +60,6 @@ libclllvm = static_library(
   cpp_args : [
     clover_cpp_args,
     clover_opencl_cpp_args,
-    clover_spirv_cpp_args,
     '-DCLANG_RESOURCE_DIR="@0@"'.format(join_paths(
       dep_llvm.get_variable(cmake : 'LLVM_LIBRARY_DIR', configtool: 'libdir'), 'clang',
       dep_llvm.version(), 'include',
@@ -75,32 +69,6 @@ libclllvm = static_library(
   dependencies : [dep_llvm, dep_elf, dep_llvmspirvlib, idep_mesautil],
 )
 
-idep_opencl_spirv = null_dep
-if with_opencl_spirv
-  libclspirv = static_library(
-    'clspirv',
-    files('spirv/invocation.cpp', 'spirv/invocation.hpp'),
-    include_directories : clover_incs,
-    cpp_args : [clover_opencl_cpp_args, clover_spirv_cpp_args],
-    gnu_symbol_visibility : 'hidden',
-    dependencies : [dep_spirv_tools, idep_mesautil],
-  )
-
-  libclnir = static_library(
-    'clnir',
-    files('nir/invocation.cpp', 'nir/invocation.hpp'),
-    include_directories : [clover_incs, inc_mesa],
-    dependencies : [idep_nir, idep_vtn, idep_mesaclc],
-    cpp_args : [clover_opencl_cpp_args, clover_spirv_cpp_args],
-    gnu_symbol_visibility : 'hidden',
-  )
-
-  idep_opencl_spirv = declare_dependency(
-    dependencies : [idep_nir],
-    link_with : [libclspirv, libclnir],
-  )
-endif
-
 clover_files = files(
   'api/context.cpp',
   'api/device.cpp',
@@ -167,10 +135,9 @@ libclover = static_library(
   include_directories : clover_incs,
   cpp_args : [
     clover_opencl_cpp_args,
-    clover_spirv_cpp_args,
     clover_cpp_args,
   ],
   gnu_symbol_visibility : 'hidden',
   link_with : [libclllvm],
-  dependencies : [idep_mesautil, idep_nir, idep_opencl_spirv],
+  dependencies : [idep_mesautil, idep_nir],
 )
diff --git a/src/gallium/frontends/clover/nir/invocation.cpp b/src/gallium/frontends/clover/nir/invocation.cpp
deleted file mode 100644
index b5e18fbe21c7f..0000000000000
--- a/src/gallium/frontends/clover/nir/invocation.cpp
+++ /dev/null
@@ -1,472 +0,0 @@
-//
-// Copyright 2019 Karol Herbst
-//
-// 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.
-//
-
-#include "invocation.hpp"
-
-#include <tuple>
-
-#include "core/device.hpp"
-#include "core/error.hpp"
-#include "core/binary.hpp"
-#include "pipe/p_state.h"
-#include "util/algorithm.hpp"
-#include "util/functional.hpp"
-
-#include <compiler/glsl_types.h>
-#include <compiler/clc/nir_clc_helpers.h>
-#include <compiler/nir/nir_builder.h>
-#include <compiler/nir/nir_serialize.h>
-#include <compiler/spirv/nir_spirv.h>
-#include <util/u_math.h>
-#include <util/hex.h>
-
-using namespace clover;
-
-#ifdef HAVE_CLOVER_SPIRV
-
-// Refs and unrefs the glsl_type_singleton.
-static class glsl_type_ref {
-public:
-   glsl_type_ref() {
-      glsl_type_singleton_init_or_ref();
-   }
-
-   ~glsl_type_ref() {
-      glsl_type_singleton_decref();
-   }
-} glsl_type_ref;
-
-static const nir_shader_compiler_options *
-dev_get_nir_compiler_options(const device &dev)
-{
-   const void *co = dev.get_compiler_options(PIPE_SHADER_IR_NIR);
-   return static_cast<const nir_shader_compiler_options*>(co);
-}
-
-static void debug_function(void *private_data,
-                   enum nir_spirv_debug_level level, size_t spirv_offset,
-                   const char *message)
-{
-   assert(private_data);
-   auto r_log = reinterpret_cast<std::string *>(private_data);
-   *r_log += message;
-}
-
-static void
-clover_arg_size_align(const glsl_type *type, unsigned *size, unsigned *align)
-{
-   if (glsl_type_is_sampler(type) || glsl_type_is_image(type)) {
-      *size = 0;
-      *align = 1;
-   } else {
-      *size = glsl_get_cl_size(type);
-      *align = glsl_get_cl_alignment(type);
-   }
-}
-
-static void
-clover_nir_add_image_uniforms(nir_shader *shader)
-{
-   /* Clover expects each image variable to take up a cl_mem worth of space in
-    * the arguments data.  Add uniforms as needed to match this expectation.
-    */
-   nir_foreach_image_variable_safe(var, shader) {
-      nir_variable *uniform = rzalloc(shader, nir_variable);
-      uniform->name = ralloc_strdup(uniform, var->name);
-      uniform->type = glsl_uintN_t_type(sizeof(cl_mem) * 8);
-      uniform->data.mode = nir_var_uniform;
-      uniform->data.read_only = true;
-      uniform->data.location = var->data.location;
-
-      exec_node_insert_node_before(&var->node, &uniform->node);
-   }
-}
-
-struct clover_lower_nir_state {
-   std::vector<binary::argument> &args;
-   uint32_t global_dims;
-   nir_variable *constant_var;
-   nir_variable *printf_buffer;
-   nir_variable *offset_vars[3];
-};
-
-static bool
-clover_lower_nir_filter(const nir_instr *instr, const void *)
-{
-   return instr->type == nir_instr_type_intrinsic;
-}
-
-static nir_def *
-clover_lower_nir_instr(nir_builder *b, nir_instr *instr, void *_state)
-{
-   clover_lower_nir_state *state = reinterpret_cast<clover_lower_nir_state*>(_state);
-   nir_intrinsic_instr *intrinsic = nir_instr_as_intrinsic(instr);
-
-   switch (intrinsic->intrinsic) {
-   case nir_intrinsic_load_printf_buffer_address: {
-      if (!state->printf_buffer) {
-         unsigned location = state->args.size();
-         state->args.emplace_back(binary::argument::global, sizeof(size_t),
-                                  8, 8, binary::argument::zero_ext,
-                                  binary::argument::printf_buffer);
-
-         const glsl_type *type = glsl_uint64_t_type();
-         state->printf_buffer = nir_variable_create(b->shader, nir_var_uniform,
-                                                    type, "global_printf_buffer");
-         state->printf_buffer->data.location = location;
-      }
-      return nir_load_var(b, state->printf_buffer);
-   }
-   case nir_intrinsic_load_base_global_invocation_id: {
-      nir_def *loads[3];
-
-      /* create variables if we didn't do so alrady */
-      if (!state->offset_vars[0]) {
-         /* TODO: fix for 64 bit */
-         /* Even though we only place one scalar argument, clover will bind up to
-          * three 32 bit values
-         */
-         unsigned location = state->args.size();
-         state->args.emplace_back(binary::argument::scalar, 4, 4, 4,
-                                  binary::argument::zero_ext,
-                                  binary::argument::grid_offset);
-
-         const glsl_type *type = glsl_uint_type();
-         for (uint32_t i = 0; i < 3; i++) {
-            state->offset_vars[i] =
-               nir_variable_create(b->shader, nir_var_uniform, type,
-                                   "global_invocation_id_offsets");
-            state->offset_vars[i]->data.location = location + i;
-         }
-      }
-
-      for (int i = 0; i < 3; i++) {
-         nir_variable *var = state->offset_vars[i];
-         loads[i] = var ? nir_load_var(b, var) : nir_imm_int(b, 0);
-      }
-
-      return nir_u2uN(b, nir_vec(b, loads, state->global_dims),
-                     intrinsic->def.bit_size);
-   }
-   case nir_intrinsic_load_constant_base_ptr: {
-      return nir_load_var(b, state->constant_var);
-   }
-
-   default:
-      return NULL;
-   }
-}
-
-static bool
-clover_lower_nir(nir_shader *nir, std::vector<binary::argument> &args,
-                 uint32_t dims, uint32_t pointer_bit_size)
-{
-   nir_variable *constant_var = NULL;
-   if (nir->constant_data_size) {
-      const glsl_type *type = pointer_bit_size == 64 ? glsl_uint64_t_type() : glsl_uint_type();
-
-      constant_var = nir_variable_create(nir, nir_var_uniform, type,
-                                         "constant_buffer_addr");
-      constant_var->data.location = args.size();
-
-      args.emplace_back(binary::argument::global, sizeof(cl_mem),
-                        pointer_bit_size / 8, pointer_bit_size / 8,
-                        binary::argument::zero_ext,
-                        binary::argument::constant_buffer);
-   }
-
-   clover_lower_nir_state state = { args, dims, constant_var };
-   return nir_shader_lower_instructions(nir,
-      clover_lower_nir_filter, clover_lower_nir_instr, &state);
-}
-
-static spirv_to_nir_options
-create_spirv_options(const device &dev, std::string &r_log)
-{
-   struct spirv_to_nir_options spirv_options = {};
-   spirv_options.environment = NIR_SPIRV_OPENCL;
-   if (dev.address_bits() == 32u) {
-      spirv_options.shared_addr_format = nir_address_format_32bit_offset;
-      spirv_options.global_addr_format = nir_address_format_32bit_global;
-      spirv_options.temp_addr_format = nir_address_format_32bit_offset;
-      spirv_options.constant_addr_format = nir_address_format_32bit_global;
-   } else {
-      spirv_options.shared_addr_format = nir_address_format_32bit_offset_as_64bit;
-      spirv_options.global_addr_format = nir_address_format_64bit_global;
-      spirv_options.temp_addr_format = nir_address_format_32bit_offset_as_64bit;
-      spirv_options.constant_addr_format = nir_address_format_64bit_global;
-   }
-   spirv_options.caps.address = true;
-   spirv_options.caps.float64 = true;
-   spirv_options.caps.int8 = true;
-   spirv_options.caps.int16 = true;
-   spirv_options.caps.int64 = true;
-   spirv_options.caps.kernel = true;
-   spirv_options.caps.kernel_image = dev.image_support();
-   spirv_options.caps.int64_atomics = dev.has_int64_atomics();
-   spirv_options.debug.func = &debug_function;
-   spirv_options.debug.private_data = &r_log;
-   spirv_options.caps.printf = true;
-   return spirv_options;
-}
-
-struct disk_cache *clover::nir::create_clc_disk_cache(void)
-{
-   struct mesa_sha1 ctx;
-   unsigned char sha1[20];
-   char cache_id[20 * 2 + 1];
-   _mesa_sha1_init(&ctx);
-
-   if (!disk_cache_get_function_identifier((void *)clover::nir::create_clc_disk_cache, &ctx))
-      return NULL;
-
-   _mesa_sha1_final(&ctx, sha1);
-
-   mesa_bytes_to_hex(cache_id, sha1, 20);
-   return disk_cache_create("clover-clc", cache_id, 0);
-}
-
-void clover::nir::check_for_libclc(const device &dev)
-{
-   if (!nir_can_find_libclc(dev.address_bits()))
-      throw error(CL_COMPILER_NOT_AVAILABLE);
-}
-
-nir_shader *clover::nir::load_libclc_nir(const device &dev, std::string &r_log)
-{
-   spirv_to_nir_options spirv_options = create_spirv_options(dev, r_log);
-   auto *compiler_options = dev_get_nir_compiler_options(dev);
-
-   return nir_load_libclc_shader(dev.address_bits(), dev.clc_cache,
-                                 &spirv_options, compiler_options,
-                                 dev.clc_cache != nullptr);
-}
-
-static bool
-can_remove_var(nir_variable *var, void *data)
-{
-   return !(glsl_type_is_sampler(var->type) ||
-            glsl_type_is_texture(var->type) ||
-            glsl_type_is_image(var->type));
-}
-
-binary clover::nir::spirv_to_nir(const binary &mod, const device &dev,
-                                 std::string &r_log)
-{
-   spirv_to_nir_options spirv_options = create_spirv_options(dev, r_log);
-   std::shared_ptr<nir_shader> nir = dev.clc_nir;
-   spirv_options.clc_shader = nir.get();
-
-   binary b;
-   // We only insert one section.
-   assert(mod.secs.size() == 1);
-   auto &section = mod.secs[0];
-
-   binary::resource_id section_id = 0;
-   for (const auto &sym : mod.syms) {
-      assert(sym.section == 0);
-
-      const auto *binary =
-         reinterpret_cast<const pipe_binary_program_header *>(section.data.data());
-      const uint32_t *data = reinterpret_cast<const uint32_t *>(binary->blob);
-      const size_t num_words = binary->num_bytes / 4;
-      const char *name = sym.name.c_str();
-      auto *compiler_options = dev_get_nir_compiler_options(dev);
-
-      nir_shader *nir = spirv_to_nir(data, num_words, nullptr, 0,
-                                     MESA_SHADER_KERNEL, name,
-                                     &spirv_options, compiler_options);
-      if (!nir) {
-         r_log += "Translation from SPIR-V to NIR for kernel \"" + sym.name +
-                  "\" failed.\n";
-         throw build_error();
-      }
-
-      nir->info.workgroup_size_variable = sym.reqd_work_group_size[0] == 0;
-      nir->info.workgroup_size[0] = sym.reqd_work_group_size[0];
-      nir->info.workgroup_size[1] = sym.reqd_work_group_size[1];
-      nir->info.workgroup_size[2] = sym.reqd_work_group_size[2];
-      nir_validate_shader(nir, "clover");
-
-      // Inline all functions first.
-      // according to the comment on nir_inline_functions
-      NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp);
-      NIR_PASS_V(nir, nir_lower_returns);
-      NIR_PASS_V(nir, nir_link_shader_functions, spirv_options.clc_shader);
-
-      NIR_PASS_V(nir, nir_inline_functions);
-      NIR_PASS_V(nir, nir_copy_prop);
-      NIR_PASS_V(nir, nir_opt_deref);
-
-      // Pick off the single entrypoint that we want.
-      nir_remove_non_entrypoints(nir);
-
-      nir_validate_shader(nir, "clover after function inlining");
-
-      NIR_PASS_V(nir, nir_lower_variable_initializers, ~nir_var_function_temp);
-
-      struct nir_lower_printf_options printf_options;
-      printf_options.max_buffer_size = dev.max_printf_buffer_size();
-
-      NIR_PASS_V(nir, nir_lower_printf, &printf_options);
-
-      NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
-
-      // copy propagate to prepare for lower_explicit_io
-      NIR_PASS_V(nir, nir_split_var_copies);
-      NIR_PASS_V(nir, nir_opt_copy_prop_vars);
-      NIR_PASS_V(nir, nir_lower_var_copies);
-      NIR_PASS_V(nir, nir_lower_vars_to_ssa);
-      NIR_PASS_V(nir, nir_opt_dce);
-      NIR_PASS_V(nir, nir_lower_convert_alu_types, NULL);
-
-      if (compiler_options->lower_to_scalar) {
-         NIR_PASS_V(nir, nir_lower_alu_to_scalar,
-                    compiler_options->lower_to_scalar_filter, NULL);
-      }
-      NIR_PASS_V(nir, nir_lower_system_values);
-      nir_lower_compute_system_values_options sysval_options = { 0 };
-      sysval_options.has_base_global_invocation_id = true;
-      NIR_PASS_V(nir, nir_lower_compute_system_values, &sysval_options);
-
-      // constant fold before lowering mem constants
-      NIR_PASS_V(nir, nir_opt_constant_folding);
-
-      NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_mem_constant, NULL);
-      NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_mem_constant,
-                 glsl_get_cl_type_size_align);
-      if (nir->constant_data_size > 0) {
-         assert(nir->constant_data == NULL);
-         nir->constant_data = rzalloc_size(nir, nir->constant_data_size);
-         nir_gather_explicit_io_initializers(nir, nir->constant_data,
-                                             nir->constant_data_size,
-                                             nir_var_mem_constant);
-      }
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant,
-                 spirv_options.constant_addr_format);
-
-      auto args = sym.args;
-      NIR_PASS_V(nir, clover_lower_nir, args, dev.max_block_size().size(),
-                 dev.address_bits());
-
-      NIR_PASS_V(nir, clover_nir_add_image_uniforms);
-      NIR_PASS_V(nir, nir_lower_vars_to_explicit_types,
-                 nir_var_uniform, clover_arg_size_align);
-      NIR_PASS_V(nir, nir_lower_vars_to_explicit_types,
-                 nir_var_mem_shared | nir_var_mem_global |
-                 nir_var_function_temp,
-                 glsl_get_cl_type_size_align);
-
-      NIR_PASS_V(nir, nir_opt_deref);
-      NIR_PASS_V(nir, nir_lower_readonly_images_to_tex, false);
-      NIR_PASS_V(nir, nir_lower_cl_images, true, true);
-      NIR_PASS_V(nir, nir_lower_memcpy);
-
-      /* use offsets for kernel inputs (uniform) */
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_uniform,
-                 nir->info.cs.ptr_size == 64 ?
-                 nir_address_format_32bit_offset_as_64bit :
-                 nir_address_format_32bit_offset);
-
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant,
-                 spirv_options.constant_addr_format);
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_shared,
-                 spirv_options.shared_addr_format);
-
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_function_temp,
-                 spirv_options.temp_addr_format);
-
-      NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_global,
-                 spirv_options.global_addr_format);
-
-      struct nir_remove_dead_variables_options remove_dead_variables_options = {};
-      remove_dead_variables_options.can_remove_var = can_remove_var;
-      NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_all, &remove_dead_variables_options);
-
-      if (compiler_options->lower_int64_options)
-         NIR_PASS_V(nir, nir_lower_int64);
-
-      NIR_PASS_V(nir, nir_opt_dce);
-
-      if (nir->constant_data_size) {
-         const char *ptr = reinterpret_cast<const char *>(nir->constant_data);
-         const binary::section constants {
-            section_id,
-            binary::section::data_constant,
-            nir->constant_data_size,
-            { ptr, ptr + nir->constant_data_size }
-         };
-         nir->constant_data = NULL;
-         nir->constant_data_size = 0;
-         b.secs.push_back(constants);
-      }
-
-      void *mem_ctx = ralloc_context(NULL);
-      unsigned printf_info_count = nir->printf_info_count;
-      u_printf_info *printf_infos = nir->printf_info;
-
-      ralloc_steal(mem_ctx, printf_infos);
-
-      struct blob blob;
-      blob_init(&blob);
-      nir_serialize(&blob, nir, false);
-
-      ralloc_free(nir);
-
-      const pipe_binary_program_header header { uint32_t(blob.size) };
-      binary::section text { section_id, binary::section::text_executable, header.num_bytes, {} };
-      text.data.insert(text.data.end(), reinterpret_cast<const char *>(&header),
-                       reinterpret_cast<const char *>(&header) + sizeof(header));
-      text.data.insert(text.data.end(), blob.data, blob.data + blob.size);
-
-      free(blob.data);
-
-      b.printf_strings_in_buffer = false;
-      b.printf_infos.reserve(printf_info_count);
-      for (unsigned i = 0; i < printf_info_count; i++) {
-         binary::printf_info info;
-
-         info.arg_sizes.reserve(printf_infos[i].num_args);
-         for (unsigned j = 0; j < printf_infos[i].num_args; j++)
-            info.arg_sizes.push_back(printf_infos[i].arg_sizes[j]);
-
-         info.strings.resize(printf_infos[i].string_size);
-         memcpy(info.strings.data(), printf_infos[i].strings, printf_infos[i].string_size);
-         b.printf_infos.push_back(info);
-      }
-
-      ralloc_free(mem_ctx);
-
-      b.syms.emplace_back(sym.name, sym.attributes,
-                          sym.reqd_work_group_size, section_id, 0, args);
-      b.secs.push_back(text);
-      section_id++;
-   }
-   return b;
-}
-#else
-binary clover::nir::spirv_to_nir(const binary &mod, const device &dev, std::string &r_log)
-{
-   r_log += "SPIR-V support in clover is not enabled.\n";
-   throw error(CL_LINKER_NOT_AVAILABLE);
-}
-#endif
diff --git a/src/gallium/frontends/clover/nir/invocation.hpp b/src/gallium/frontends/clover/nir/invocation.hpp
deleted file mode 100644
index 873f625ef98f0..0000000000000
--- a/src/gallium/frontends/clover/nir/invocation.hpp
+++ /dev/null
@@ -1,46 +0,0 @@
-//
-// Copyright 2019 Karol Herbst
-//
-// 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.
-//
-
-#ifndef CLOVER_NIR_INVOCATION_HPP
-#define CLOVER_NIR_INVOCATION_HPP
-
-#include "core/binary.hpp"
-#include <util/disk_cache.h>
-
-struct nir_shader;
-
-namespace clover {
-   class device;
-   namespace nir {
-      void check_for_libclc(const device &dev);
-
-      // converts libclc spirv into nir
-      nir_shader *load_libclc_nir(const device &dev, std::string &r_log);
-
-      struct disk_cache *create_clc_disk_cache(void);
-
-      // converts a given spirv binary to nir
-      binary spirv_to_nir(const binary &bin, const device &dev, std::string &r_log);
-   }
-}
-
-#endif
diff --git a/src/gallium/frontends/clover/spirv/invocation.cpp b/src/gallium/frontends/clover/spirv/invocation.cpp
deleted file mode 100644
index 5fa2779ca6359..0000000000000
--- a/src/gallium/frontends/clover/spirv/invocation.cpp
+++ /dev/null
@@ -1,1031 +0,0 @@
-//
-// Copyright 2018 Pierre Moreau
-//
-// 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.
-//
-
-#include "invocation.hpp"
-
-#include <limits>
-#include <unordered_map>
-#include <unordered_set>
-#include <vector>
-
-#ifdef HAVE_CLOVER_SPIRV
-#include <spirv-tools/libspirv.hpp>
-#include <spirv-tools/linker.hpp>
-#endif
-
-#include "core/error.hpp"
-#include "core/platform.hpp"
-#include "invocation.hpp"
-#include "llvm/util.hpp"
-#include "pipe/p_state.h"
-#include "util/algorithm.hpp"
-#include "util/functional.hpp"
-#include "util/u_math.h"
-
-#include "compiler/spirv/spirv.h"
-
-#define SPIRV_HEADER_WORD_SIZE 5
-
-using namespace clover;
-
-using clover::detokenize;
-
-#ifdef HAVE_CLOVER_SPIRV
-namespace {
-
-   static const std::array<std::string,7> type_strs = {
-      "uchar", "ushort", "uint", "ulong", "half", "float", "double"
-   };
-
-   template<typename T>
-   T get(const char *source, size_t index) {
-      const uint32_t *word_ptr = reinterpret_cast<const uint32_t *>(source);
-      return static_cast<T>(word_ptr[index]);
-   }
-
-   enum binary::argument::type
-   convert_storage_class(SpvStorageClass storage_class, std::string &err) {
-      switch (storage_class) {
-      case SpvStorageClassFunction:
-         return binary::argument::scalar;
-      case SpvStorageClassUniformConstant:
-         return binary::argument::global;
-      case SpvStorageClassWorkgroup:
-         return binary::argument::local;
-      case SpvStorageClassCrossWorkgroup:
-         return binary::argument::global;
-      default:
-         err += "Invalid storage type " + std::to_string(storage_class) + "\n";
-         throw build_error();
-      }
-   }
-
-   cl_kernel_arg_address_qualifier
-   convert_storage_class_to_cl(SpvStorageClass storage_class) {
-      switch (storage_class) {
-      case SpvStorageClassUniformConstant:
-         return CL_KERNEL_ARG_ADDRESS_CONSTANT;
-      case SpvStorageClassWorkgroup:
-         return CL_KERNEL_ARG_ADDRESS_LOCAL;
-      case SpvStorageClassCrossWorkgroup:
-         return CL_KERNEL_ARG_ADDRESS_GLOBAL;
-      case SpvStorageClassFunction:
-      default:
-         return CL_KERNEL_ARG_ADDRESS_PRIVATE;
-      }
-   }
-
-   enum binary::argument::type
-   convert_image_type(SpvId id, SpvDim dim, SpvAccessQualifier access,
-                      std::string &err) {
-      switch (dim) {
-      case SpvDim1D:
-      case SpvDim2D:
-      case SpvDim3D:
-      case SpvDimBuffer:
-         switch (access) {
-         case SpvAccessQualifierReadOnly:
-            return binary::argument::image_rd;
-         case SpvAccessQualifierWriteOnly:
-            return binary::argument::image_wr;
-         default:
-            err += "Unknown access qualifier " + std::to_string(access) + " for image "
-                +  std::to_string(id) + ".\n";
-            throw build_error();
-         }
-      default:
-         err += "Unknown dimension " + std::to_string(dim) + " for image "
-             +  std::to_string(id) + ".\n";
-         throw build_error();
-      }
-   }
-
-   binary::section
-   make_text_section(const std::string &code,
-                     enum binary::section::type section_type) {
-      const pipe_binary_program_header header { uint32_t(code.size()) };
-      binary::section text { 0, section_type, header.num_bytes, {} };
-
-      text.data.reserve(sizeof(header) + header.num_bytes);
-      text.data.insert(text.data.end(), reinterpret_cast<const char *>(&header),
-                       reinterpret_cast<const char *>(&header) + sizeof(header));
-      text.data.insert(text.data.end(), code.begin(), code.end());
-
-      return text;
-   }
-
-   binary
-   create_binary_from_spirv(const std::string &source,
-                            size_t pointer_byte_size,
-                            std::string &err) {
-      const size_t length = source.size() / sizeof(uint32_t);
-      size_t i = SPIRV_HEADER_WORD_SIZE; // Skip header
-
-      std::string kernel_name;
-      size_t kernel_nb = 0u;
-      std::vector<binary::argument> args;
-      std::vector<size_t> req_local_size;
-
-      binary b;
-
-      std::vector<std::string> attributes;
-      std::unordered_map<SpvId, std::vector<size_t> > req_local_sizes;
-      std::unordered_map<SpvId, std::string> kernels;
-      std::unordered_map<SpvId, binary::argument> types;
-      std::unordered_map<SpvId, SpvId> pointer_types;
-      std::unordered_map<SpvId, unsigned int> constants;
-      std::unordered_set<SpvId> packed_structures;
-      std::unordered_map<SpvId, std::vector<SpvFunctionParameterAttribute>>
-         func_param_attr_map;
-      std::unordered_map<SpvId, std::string> names;
-      std::unordered_map<SpvId, cl_kernel_arg_type_qualifier> qualifiers;
-      std::unordered_map<std::string, std::vector<std::string> > param_type_names;
-      std::unordered_map<std::string, std::vector<std::string> > param_qual_names;
-
-      while (i < length) {
-         const auto inst = &source[i * sizeof(uint32_t)];
-         const auto desc_word = get<uint32_t>(inst, 0);
-         const auto opcode = static_cast<SpvOp>(desc_word & SpvOpCodeMask);
-         const unsigned int num_operands = desc_word >> SpvWordCountShift;
-
-         switch (opcode) {
-         case SpvOpName: {
-            names.emplace(get<SpvId>(inst, 1),
-                          source.data() + (i + 2u) * sizeof(uint32_t));
-            break;
-         }
-
-         case SpvOpString: {
-            // SPIRV-LLVM-Translator stores param type names as OpStrings
-            std::string str(source.data() + (i + 2u) * sizeof(uint32_t));
-            if (str.find("kernel_arg_type.") == 0) {
-               std::string line;
-               std::istringstream istream(str.substr(16));
-
-               std::getline(istream, line, '.');
-
-               std::string k = line;
-               while (std::getline(istream, line, ','))
-                  param_type_names[k].push_back(line);
-            } else if (str.find("kernel_arg_type_qual.") == 0) {
-               std::string line;
-               std::istringstream istream(str.substr(21));
-
-               std::getline(istream, line, '.');
-               std::string k = line;
-               while (std::getline(istream, line, ','))
-                  param_qual_names[k].push_back(line);
-            } else
-               continue;
-            break;
-         }
-
-         case SpvOpEntryPoint:
-            if (get<SpvExecutionModel>(inst, 1) == SpvExecutionModelKernel)
-               kernels.emplace(get<SpvId>(inst, 2),
-                               source.data() + (i + 3u) * sizeof(uint32_t));
-            break;
-
-         case SpvOpExecutionMode:
-            switch (get<SpvExecutionMode>(inst, 2)) {
-            case SpvExecutionModeLocalSize: {
-               req_local_sizes[get<SpvId>(inst, 1)] = {
-                  get<uint32_t>(inst, 3),
-                  get<uint32_t>(inst, 4),
-                  get<uint32_t>(inst, 5)
-               };
-               std::string s = "reqd_work_group_size(";
-               s += std::to_string(get<uint32_t>(inst, 3));
-               s += ",";
-               s += std::to_string(get<uint32_t>(inst, 4));
-               s += ",";
-               s += std::to_string(get<uint32_t>(inst, 5));
-               s += ")";
-               attributes.emplace_back(s);
-               break;
-            }
-            case SpvExecutionModeLocalSizeHint: {
-               std::string s = "work_group_size_hint(";
-               s += std::to_string(get<uint32_t>(inst, 3));
-               s += ",";
-               s += std::to_string(get<uint32_t>(inst, 4));
-               s += ",";
-               s += std::to_string(get<uint32_t>(inst, 5));
-               s += ")";
-               attributes.emplace_back(s);
-               break;
-            }
-	    case SpvExecutionModeVecTypeHint: {
-               uint32_t val = get<uint32_t>(inst, 3);
-               uint32_t size = val >> 16;
-
-               val &= 0xf;
-               if (val > 6)
-                  val = 0;
-               std::string s = "vec_type_hint(";
-               s += type_strs[val];
-               s += std::to_string(size);
-               s += ")";
-               attributes.emplace_back(s);
-	       break;
-            }
-            default:
-               break;
-            }
-            break;
-
-         case SpvOpDecorate: {
-            const auto id = get<SpvId>(inst, 1);
-            const auto decoration = get<SpvDecoration>(inst, 2);
-            switch (decoration) {
-            case SpvDecorationCPacked:
-               packed_structures.emplace(id);
-               break;
-            case SpvDecorationFuncParamAttr: {
-               const auto attribute =
-                  get<SpvFunctionParameterAttribute>(inst, 3u);
-               func_param_attr_map[id].push_back(attribute);
-               break;
-            }
-            case SpvDecorationVolatile:
-               qualifiers[id] |= CL_KERNEL_ARG_TYPE_VOLATILE;
-               break;
-            default:
-               break;
-            }
-            break;
-         }
-
-         case SpvOpGroupDecorate: {
-            const auto group_id = get<SpvId>(inst, 1);
-            if (packed_structures.count(group_id)) {
-               for (unsigned int i = 2u; i < num_operands; ++i)
-                  packed_structures.emplace(get<SpvId>(inst, i));
-            }
-            const auto func_param_attr_iter =
-               func_param_attr_map.find(group_id);
-            if (func_param_attr_iter != func_param_attr_map.end()) {
-               for (unsigned int i = 2u; i < num_operands; ++i) {
-                  auto &attrs = func_param_attr_map[get<SpvId>(inst, i)];
-                  attrs.insert(attrs.begin(),
-                               func_param_attr_iter->second.begin(),
-                               func_param_attr_iter->second.end());
-               }
-            }
-            if (qualifiers.count(group_id)) {
-               for (unsigned int i = 2u; i < num_operands; ++i)
-                  qualifiers[get<SpvId>(inst, i)] |= qualifiers[group_id];
-            }
-            break;
-         }
-
-         case SpvOpConstant:
-            // We only care about constants that represent the size of arrays.
-            // If they are passed as argument, they will never be more than
-            // 4GB-wide, and even if they did, a clover::binary::argument size
-            // is represented by an int.
-            constants[get<SpvId>(inst, 2)] = get<unsigned int>(inst, 3u);
-            break;
-
-         case SpvOpTypeInt:
-         case SpvOpTypeFloat: {
-            const auto size = get<uint32_t>(inst, 2) / 8u;
-            const auto id = get<SpvId>(inst, 1);
-            types[id] = { binary::argument::scalar, size, size, size,
-                          binary::argument::zero_ext };
-            types[id].info.address_qualifier = CL_KERNEL_ARG_ADDRESS_PRIVATE;
-            break;
-         }
-
-         case SpvOpTypeArray: {
-            const auto id = get<SpvId>(inst, 1);
-            const auto type_id = get<SpvId>(inst, 2);
-            const auto types_iter = types.find(type_id);
-            if (types_iter == types.end())
-               break;
-
-            const auto constant_id = get<SpvId>(inst, 3);
-            const auto constants_iter = constants.find(constant_id);
-            if (constants_iter == constants.end()) {
-               err += "Constant " + std::to_string(constant_id) +
-                  " is missing\n";
-               throw build_error();
-            }
-            const auto elem_size = types_iter->second.size;
-            const auto elem_nbs = constants_iter->second;
-            const auto size = elem_size * elem_nbs;
-            types[id] = { binary::argument::scalar, size, size,
-                          types_iter->second.target_align,
-                          binary::argument::zero_ext };
-            break;
-         }
-
-         case SpvOpTypeStruct: {
-            const auto id = get<SpvId>(inst, 1);
-            const bool is_packed = packed_structures.count(id);
-
-            unsigned struct_size = 0u;
-            unsigned struct_align = 1u;
-            for (unsigned j = 2u; j < num_operands; ++j) {
-               const auto type_id = get<SpvId>(inst, j);
-               const auto types_iter = types.find(type_id);
-
-               // If a type was not found, that means it is not one of the
-               // types allowed as kernel arguments. And since the binary has
-               // been validated, this means this type is not used for kernel
-               // arguments, and therefore can be ignored.
-               if (types_iter == types.end())
-                  break;
-
-               const auto alignment = is_packed ? 1u
-                                                : types_iter->second.target_align;
-               const auto padding = (-struct_size) & (alignment - 1u);
-               struct_size += padding + types_iter->second.target_size;
-               struct_align = std::max(struct_align, alignment);
-            }
-            struct_size += (-struct_size) & (struct_align - 1u);
-            types[id] = { binary::argument::scalar, struct_size, struct_size,
-                          struct_align, binary::argument::zero_ext };
-            break;
-         }
-
-         case SpvOpTypeVector: {
-            const auto id = get<SpvId>(inst, 1);
-            const auto type_id = get<SpvId>(inst, 2);
-            const auto types_iter = types.find(type_id);
-
-            // If a type was not found, that means it is not one of the
-            // types allowed as kernel arguments. And since the binary has
-            // been validated, this means this type is not used for kernel
-            // arguments, and therefore can be ignored.
-            if (types_iter == types.end())
-               break;
-
-            const auto elem_size = types_iter->second.size;
-            const auto elem_nbs = get<uint32_t>(inst, 3);
-            const auto size = elem_size * (elem_nbs != 3 ? elem_nbs : 4);
-            types[id] = { binary::argument::scalar, size, size, size,
-                          binary::argument::zero_ext };
-            types[id].info.address_qualifier = CL_KERNEL_ARG_ADDRESS_PRIVATE;
-            break;
-         }
-
-         case SpvOpTypeForwardPointer: // FALLTHROUGH
-         case SpvOpTypePointer: {
-            const auto id = get<SpvId>(inst, 1);
-            const auto storage_class = get<SpvStorageClass>(inst, 2);
-            // Input means this is for a builtin variable, which can not be
-            // passed as an argument to a kernel.
-            if (storage_class == SpvStorageClassInput)
-               break;
-
-            if (opcode == SpvOpTypePointer)
-               pointer_types[id] = get<SpvId>(inst, 3);
-
-            binary::size_t alignment;
-            if (storage_class == SpvStorageClassWorkgroup)
-               alignment = opcode == SpvOpTypePointer ? types[pointer_types[id]].target_align : 0;
-            else
-               alignment = pointer_byte_size;
-
-            types[id] = { convert_storage_class(storage_class, err),
-                          sizeof(cl_mem),
-                          static_cast<binary::size_t>(pointer_byte_size),
-                          alignment,
-                          binary::argument::zero_ext };
-            types[id].info.address_qualifier = convert_storage_class_to_cl(storage_class);
-            break;
-         }
-
-         case SpvOpTypeSampler:
-            types[get<SpvId>(inst, 1)] = { binary::argument::sampler,
-                                             sizeof(cl_sampler) };
-            break;
-
-         case SpvOpTypeImage: {
-            const auto id = get<SpvId>(inst, 1);
-            const auto dim = get<SpvDim>(inst, 3);
-            const auto access = get<SpvAccessQualifier>(inst, 9);
-            types[id] = { convert_image_type(id, dim, access, err),
-                          sizeof(cl_mem), sizeof(cl_mem), sizeof(cl_mem),
-                          binary::argument::zero_ext };
-            break;
-         }
-
-         case SpvOpTypePipe: // FALLTHROUGH
-         case SpvOpTypeQueue: {
-            err += "TypePipe and TypeQueue are valid SPIR-V 1.0 types, but are "
-                   "not available in the currently supported OpenCL C version."
-                   "\n";
-            throw build_error();
-         }
-
-         case SpvOpFunction: {
-            auto id = get<SpvId>(inst, 2);
-            const auto kernels_iter = kernels.find(id);
-            if (kernels_iter != kernels.end())
-               kernel_name = kernels_iter->second;
-
-            const auto req_local_size_iter = req_local_sizes.find(id);
-            if (req_local_size_iter != req_local_sizes.end())
-               req_local_size =  (*req_local_size_iter).second;
-            else
-               req_local_size = { 0, 0, 0 };
-
-            break;
-         }
-
-         case SpvOpFunctionParameter: {
-            if (kernel_name.empty())
-               break;
-
-            const auto id = get<SpvId>(inst, 2);
-            const auto type_id = get<SpvId>(inst, 1);
-            auto arg = types.find(type_id)->second;
-            const auto &func_param_attr_iter =
-               func_param_attr_map.find(get<SpvId>(inst, 2));
-            if (func_param_attr_iter != func_param_attr_map.end()) {
-               for (auto &i : func_param_attr_iter->second) {
-                  switch (i) {
-                  case SpvFunctionParameterAttributeSext:
-                     arg.ext_type = binary::argument::sign_ext;
-                     break;
-                  case SpvFunctionParameterAttributeZext:
-                     arg.ext_type = binary::argument::zero_ext;
-                     break;
-                  case SpvFunctionParameterAttributeByVal: {
-                     const SpvId ptr_type_id =
-                        pointer_types.find(type_id)->second;
-                     arg = types.find(ptr_type_id)->second;
-                     break;
-                  }
-                  case SpvFunctionParameterAttributeNoAlias:
-                     arg.info.type_qualifier |= CL_KERNEL_ARG_TYPE_RESTRICT;
-                     break;
-                  case SpvFunctionParameterAttributeNoWrite:
-                     arg.info.type_qualifier |= CL_KERNEL_ARG_TYPE_CONST;
-                     break;
-                  default:
-                     break;
-                  }
-               }
-            }
-
-            auto name_it = names.find(id);
-            if (name_it != names.end())
-               arg.info.arg_name = (*name_it).second;
-
-            arg.info.type_qualifier |= qualifiers[id];
-            arg.info.address_qualifier = types[type_id].info.address_qualifier;
-            arg.info.access_qualifier = CL_KERNEL_ARG_ACCESS_NONE;
-            args.emplace_back(arg);
-            break;
-         }
-
-         case SpvOpFunctionEnd: {
-            if (kernel_name.empty())
-               break;
-
-            for (size_t i = 0; i < param_type_names[kernel_name].size(); i++)
-               args[i].info.type_name = param_type_names[kernel_name][i];
-
-            for (size_t i = 0; i < param_qual_names[kernel_name].size(); i++)
-               if (param_qual_names[kernel_name][i].find("const") != std::string::npos)
-                  args[i].info.type_qualifier |= CL_KERNEL_ARG_TYPE_CONST;
-            b.syms.emplace_back(kernel_name, detokenize(attributes, " "),
-                                req_local_size, 0, kernel_nb, args);
-            ++kernel_nb;
-            kernel_name.clear();
-            args.clear();
-            attributes.clear();
-            break;
-         }
-         default:
-            break;
-         }
-
-         i += num_operands;
-      }
-
-      b.secs.push_back(make_text_section(source,
-                                         binary::section::text_intermediate));
-      return b;
-   }
-
-   bool
-   check_spirv_version(const device &dev, const char *binary,
-                       std::string &r_log) {
-      const auto spirv_version = get<uint32_t>(binary, 1u);
-      const auto supported_spirv_versions = clover::spirv::supported_versions();
-      const auto compare_versions =
-         [module_version =
-            clover::spirv::to_opencl_version_encoding(spirv_version)](const cl_name_version &supported){
-         return supported.version == module_version;
-      };
-
-      if (std::find_if(supported_spirv_versions.cbegin(),
-                       supported_spirv_versions.cend(),
-                       compare_versions) != supported_spirv_versions.cend())
-         return true;
-
-      r_log += "SPIR-V version " +
-               clover::spirv::version_to_string(spirv_version) +
-               " is not supported; supported versions:";
-      for (const auto &version : supported_spirv_versions) {
-         r_log += " " + clover::spirv::version_to_string(version.version);
-      }
-      r_log += "\n";
-      return false;
-   }
-
-   bool
-   check_capabilities(const device &dev, const std::string &source,
-                      std::string &r_log) {
-      const size_t length = source.size() / sizeof(uint32_t);
-      size_t i = SPIRV_HEADER_WORD_SIZE; // Skip header
-
-      while (i < length) {
-         const auto desc_word = get<uint32_t>(source.data(), i);
-         const auto opcode = static_cast<SpvOp>(desc_word & SpvOpCodeMask);
-         const unsigned int num_operands = desc_word >> SpvWordCountShift;
-
-         if (opcode != SpvOpCapability)
-            break;
-
-         const auto capability = get<SpvCapability>(source.data(), i + 1u);
-         switch (capability) {
-         // Mandatory capabilities
-         case SpvCapabilityAddresses:
-         case SpvCapabilityFloat16Buffer:
-         case SpvCapabilityGroups:
-         case SpvCapabilityInt64:
-         case SpvCapabilityInt16:
-         case SpvCapabilityInt8:
-         case SpvCapabilityKernel:
-         case SpvCapabilityLinkage:
-         case SpvCapabilityVector16:
-            break;
-         // Optional capabilities
-         case SpvCapabilityImageBasic:
-         case SpvCapabilityLiteralSampler:
-         case SpvCapabilitySampled1D:
-         case SpvCapabilityImage1D:
-         case SpvCapabilitySampledBuffer:
-         case SpvCapabilityImageBuffer:
-            if (!dev.image_support()) {
-               r_log += "Capability 'ImageBasic' is not supported.\n";
-               return false;
-            }
-            break;
-         case SpvCapabilityFloat64:
-            if (!dev.has_doubles()) {
-               r_log += "Capability 'Float64' is not supported.\n";
-               return false;
-            }
-            break;
-         // Enabled through extensions
-         case SpvCapabilityFloat16:
-            if (!dev.has_halves()) {
-               r_log += "Capability 'Float16' is not supported.\n";
-               return false;
-            }
-            break;
-         case SpvCapabilityInt64Atomics:
-            if (!dev.has_int64_atomics()) {
-               r_log += "Capability 'Int64Atomics' is not supported.\n";
-               return false;
-            }
-            break;
-         default:
-            r_log += "Capability '" + std::to_string(capability) +
-                     "' is not supported.\n";
-            return false;
-         }
-
-         i += num_operands;
-      }
-
-      return true;
-   }
-
-   bool
-   check_extensions(const device &dev, const std::string &source,
-                    std::string &r_log) {
-      const size_t length = source.size() / sizeof(uint32_t);
-      size_t i = SPIRV_HEADER_WORD_SIZE; // Skip header
-      const auto spirv_extensions = spirv::supported_extensions();
-
-      while (i < length) {
-         const auto desc_word = get<uint32_t>(source.data(), i);
-         const auto opcode = static_cast<SpvOp>(desc_word & SpvOpCodeMask);
-         const unsigned int num_operands = desc_word >> SpvWordCountShift;
-
-         if (opcode == SpvOpCapability) {
-            i += num_operands;
-            continue;
-         }
-         if (opcode != SpvOpExtension)
-            break;
-
-         const std::string extension = source.data() + (i + 1u) * sizeof(uint32_t);
-         if (spirv_extensions.count(extension) == 0) {
-            r_log += "Extension '" + extension + "' is not supported.\n";
-            return false;
-         }
-
-         i += num_operands;
-      }
-
-      return true;
-   }
-
-   bool
-   check_memory_model(const device &dev, const std::string &source,
-                      std::string &r_log) {
-      const size_t length = source.size() / sizeof(uint32_t);
-      size_t i = SPIRV_HEADER_WORD_SIZE; // Skip header
-
-      while (i < length) {
-         const auto desc_word = get<uint32_t>(source.data(), i);
-         const auto opcode = static_cast<SpvOp>(desc_word & SpvOpCodeMask);
-         const unsigned int num_operands = desc_word >> SpvWordCountShift;
-
-         switch (opcode) {
-         case SpvOpMemoryModel:
-            switch (get<SpvAddressingModel>(source.data(), i + 1u)) {
-            case SpvAddressingModelPhysical32:
-               return dev.address_bits() == 32;
-            case SpvAddressingModelPhysical64:
-               return dev.address_bits() == 64;
-            default:
-               unreachable("Only Physical32 and Physical64 are valid for OpenCL, and the binary was already validated");
-               return false;
-            }
-            break;
-         default:
-            break;
-         }
-
-         i += num_operands;
-      }
-
-      return false;
-   }
-
-   // Copies the input binary and convert it to the endianness of the host CPU.
-   std::string
-   spirv_to_cpu(const std::string &binary)
-   {
-      const uint32_t first_word = get<uint32_t>(binary.data(), 0u);
-      if (first_word == SpvMagicNumber)
-         return binary;
-
-      std::vector<char> cpu_endianness_binary(binary.size());
-      for (size_t i = 0; i < (binary.size() / 4u); ++i) {
-         const uint32_t word = get<uint32_t>(binary.data(), i);
-         reinterpret_cast<uint32_t *>(cpu_endianness_binary.data())[i] =
-            util_bswap32(word);
-      }
-
-      return std::string(cpu_endianness_binary.begin(),
-                         cpu_endianness_binary.end());
-   }
-
-#ifdef HAVE_CLOVER_SPIRV
-   std::string
-   format_validator_msg(spv_message_level_t level, const char * /* source */,
-                        const spv_position_t &position, const char *message) {
-      std::string level_str;
-      switch (level) {
-      case SPV_MSG_FATAL:
-         level_str = "Fatal";
-         break;
-      case SPV_MSG_INTERNAL_ERROR:
-         level_str = "Internal error";
-         break;
-      case SPV_MSG_ERROR:
-         level_str = "Error";
-         break;
-      case SPV_MSG_WARNING:
-         level_str = "Warning";
-         break;
-      case SPV_MSG_INFO:
-         level_str = "Info";
-         break;
-      case SPV_MSG_DEBUG:
-         level_str = "Debug";
-         break;
-      }
-      return "[" + level_str + "] At word No." +
-             std::to_string(position.index) + ": \"" + message + "\"\n";
-   }
-
-   spv_target_env
-   convert_opencl_version_to_target_env(const cl_version opencl_version) {
-      // Pick 1.2 for 3.0 for now
-      if (opencl_version == CL_MAKE_VERSION(3, 0, 0)) {
-         return SPV_ENV_OPENCL_1_2;
-      } else if (opencl_version == CL_MAKE_VERSION(2, 2, 0)) {
-         return SPV_ENV_OPENCL_2_2;
-      } else if (opencl_version == CL_MAKE_VERSION(2, 1, 0)) {
-         return SPV_ENV_OPENCL_2_1;
-      } else if (opencl_version == CL_MAKE_VERSION(2, 0, 0)) {
-         return SPV_ENV_OPENCL_2_0;
-      } else if (opencl_version == CL_MAKE_VERSION(1, 2, 0) ||
-                 opencl_version == CL_MAKE_VERSION(1, 1, 0) ||
-                 opencl_version == CL_MAKE_VERSION(1, 0, 0)) {
-         // SPIR-V is only defined for OpenCL >= 1.2, however some drivers
-         // might use it with OpenCL 1.0 and 1.1.
-         return SPV_ENV_OPENCL_1_2;
-      } else {
-         throw build_error("Invalid OpenCL version");
-      }
-   }
-#endif
-
-}
-
-bool
-clover::spirv::is_binary_spirv(const std::string &binary)
-{
-   // A SPIR-V binary is at the very least 5 32-bit words, which represent the
-   // SPIR-V header.
-   if (binary.size() < 20u)
-      return false;
-
-   const uint32_t first_word =
-      reinterpret_cast<const uint32_t *>(binary.data())[0u];
-   return (first_word == SpvMagicNumber) ||
-          (util_bswap32(first_word) == SpvMagicNumber);
-}
-
-std::string
-clover::spirv::version_to_string(uint32_t version) {
-   const uint32_t major_version = (version >> 16) & 0xff;
-   const uint32_t minor_version = (version >> 8) & 0xff;
-   return std::to_string(major_version) + '.' +
-      std::to_string(minor_version);
-}
-
-binary
-clover::spirv::compile_program(const std::string &binary,
-                               const device &dev, std::string &r_log,
-                               bool validate) {
-   std::string source = spirv_to_cpu(binary);
-
-   if (validate && !is_valid_spirv(source, dev.device_version(), r_log))
-      throw build_error();
-
-   if (!check_spirv_version(dev, source.data(), r_log))
-      throw build_error();
-   if (!check_capabilities(dev, source, r_log))
-      throw build_error();
-   if (!check_extensions(dev, source, r_log))
-      throw build_error();
-   if (!check_memory_model(dev, source, r_log))
-      throw build_error();
-
-   return create_binary_from_spirv(source,
-                                   dev.address_bits() == 32 ? 4u : 8u, r_log);
-}
-
-binary
-clover::spirv::link_program(const std::vector<binary> &binaries,
-                            const device &dev, const std::string &opts,
-                            std::string &r_log) {
-   std::vector<std::string> options = tokenize(opts);
-
-   bool create_library = false;
-
-   std::string ignored_options;
-   for (const std::string &option : options) {
-      if (option == "-create-library") {
-         create_library = true;
-      } else {
-         ignored_options += "'" + option + "' ";
-      }
-   }
-   if (!ignored_options.empty()) {
-      r_log += "Ignoring the following link options: " + ignored_options
-            + "\n";
-   }
-
-   spvtools::LinkerOptions linker_options;
-   linker_options.SetCreateLibrary(create_library);
-
-   binary b;
-
-   const auto section_type = create_library ? binary::section::text_library :
-                                              binary::section::text_executable;
-
-   std::vector<const uint32_t *> sections;
-   sections.reserve(binaries.size());
-   std::vector<size_t> lengths;
-   lengths.reserve(binaries.size());
-
-   auto const validator_consumer = [&r_log](spv_message_level_t level,
-                                            const char *source,
-                                            const spv_position_t &position,
-                                            const char *message) {
-      r_log += format_validator_msg(level, source, position, message);
-   };
-
-   for (const auto &bin : binaries) {
-      const auto &bsec = find([](const binary::section &sec) {
-                  return sec.type == binary::section::text_intermediate ||
-                         sec.type == binary::section::text_library;
-               }, bin.secs);
-
-      const auto c_il = ((struct pipe_binary_program_header*)bsec.data.data())->blob;
-      const auto length = bsec.size;
-
-      if (!check_spirv_version(dev, c_il, r_log))
-         throw error(CL_LINK_PROGRAM_FAILURE);
-
-      sections.push_back(reinterpret_cast<const uint32_t *>(c_il));
-      lengths.push_back(length / sizeof(uint32_t));
-   }
-
-   std::vector<uint32_t> linked_binary;
-
-   const cl_version opencl_version = dev.device_version();
-   const spv_target_env target_env =
-      convert_opencl_version_to_target_env(opencl_version);
-
-   const spvtools::MessageConsumer consumer = validator_consumer;
-   spvtools::Context context(target_env);
-   context.SetMessageConsumer(std::move(consumer));
-
-   if (Link(context, sections.data(), lengths.data(), sections.size(),
-            &linked_binary, linker_options) != SPV_SUCCESS)
-      throw error(CL_LINK_PROGRAM_FAILURE);
-
-   std::string final_binary{
-         reinterpret_cast<char *>(linked_binary.data()),
-         reinterpret_cast<char *>(linked_binary.data() +
-               linked_binary.size()) };
-   if (!is_valid_spirv(final_binary, opencl_version, r_log))
-      throw error(CL_LINK_PROGRAM_FAILURE);
-
-   if (has_flag(llvm::debug::spirv))
-      llvm::debug::log(".spvasm", spirv::print_module(final_binary, dev.device_version()));
-
-   for (const auto &bin : binaries)
-      b.syms.insert(b.syms.end(), bin.syms.begin(), bin.syms.end());
-
-   b.secs.emplace_back(make_text_section(final_binary, section_type));
-
-   return b;
-}
-
-bool
-clover::spirv::is_valid_spirv(const std::string &binary,
-                              const cl_version opencl_version,
-                              std::string &r_log) {
-   auto const validator_consumer =
-      [&r_log](spv_message_level_t level, const char *source,
-               const spv_position_t &position, const char *message) {
-      r_log += format_validator_msg(level, source, position, message);
-   };
-
-   const spv_target_env target_env =
-      convert_opencl_version_to_target_env(opencl_version);
-   spvtools::SpirvTools spvTool(target_env);
-   spvTool.SetMessageConsumer(validator_consumer);
-
-   spvtools::ValidatorOptions validator_options;
-   validator_options.SetUniversalLimit(spv_validator_limit_max_function_args,
-                                       std::numeric_limits<uint32_t>::max());
-
-   return spvTool.Validate(reinterpret_cast<const uint32_t *>(binary.data()),
-                           binary.size() / 4u, validator_options);
-}
-
-std::string
-clover::spirv::print_module(const std::string &binary,
-                            const cl_version opencl_version) {
-   const spv_target_env target_env =
-      convert_opencl_version_to_target_env(opencl_version);
-   spvtools::SpirvTools spvTool(target_env);
-   spv_context spvContext = spvContextCreate(target_env);
-   if (!spvContext)
-      return "Failed to create an spv_context for disassembling the binary.";
-
-   spv_text disassembly;
-   spvBinaryToText(spvContext,
-                   reinterpret_cast<const uint32_t *>(binary.data()),
-                   binary.size() / 4u, SPV_BINARY_TO_TEXT_OPTION_NONE,
-                   &disassembly, nullptr);
-   spvContextDestroy(spvContext);
-
-   const std::string disassemblyStr = disassembly->str;
-   spvTextDestroy(disassembly);
-
-   return disassemblyStr;
-}
-
-std::unordered_set<std::string>
-clover::spirv::supported_extensions() {
-   return {
-      /* this is only a hint so all devices support that */
-      "SPV_KHR_no_integer_wrap_decoration"
-   };
-}
-
-std::vector<cl_name_version>
-clover::spirv::supported_versions() {
-   return { cl_name_version { CL_MAKE_VERSION(1u, 0u, 0u), "SPIR-V" } };
-}
-
-cl_version
-clover::spirv::to_opencl_version_encoding(uint32_t version) {
-      return CL_MAKE_VERSION((version >> 16u) & 0xff,
-                             (version >> 8u) & 0xff, 0u);
-}
-
-uint32_t
-clover::spirv::to_spirv_version_encoding(cl_version version) {
-   return ((CL_VERSION_MAJOR(version) & 0xff) << 16u) |
-          ((CL_VERSION_MINOR(version) & 0xff) << 8u);
-}
-
-#else
-bool
-clover::spirv::is_binary_spirv(const std::string &binary)
-{
-   return false;
-}
-
-bool
-clover::spirv::is_valid_spirv(const std::string &/*binary*/,
-                              const cl_version opencl_version,
-                              std::string &/*r_log*/) {
-   return false;
-}
-
-std::string
-clover::spirv::version_to_string(uint32_t version) {
-   return "";
-}
-
-binary
-clover::spirv::compile_program(const std::string &binary,
-                               const device &dev, std::string &r_log,
-                               bool validate) {
-   r_log += "SPIR-V support in clover is not enabled.\n";
-   throw build_error();
-}
-
-binary
-clover::spirv::link_program(const std::vector<binary> &/*binaries*/,
-                            const device &/*dev*/, const std::string &/*opts*/,
-                            std::string &r_log) {
-   r_log += "SPIR-V support in clover is not enabled.\n";
-   throw error(CL_LINKER_NOT_AVAILABLE);
-}
-
-std::string
-clover::spirv::print_module(const std::string &binary,
-                            const cl_version opencl_version) {
-   return std::string();
-}
-
-std::unordered_set<std::string>
-clover::spirv::supported_extensions() {
-   return {};
-}
-
-std::vector<cl_name_version>
-clover::spirv::supported_versions() {
-   return {};
-}
-
-cl_version
-clover::spirv::to_opencl_version_encoding(uint32_t version) {
-   return CL_MAKE_VERSION(0u, 0u, 0u);
-}
-
-uint32_t
-clover::spirv::to_spirv_version_encoding(cl_version version) {
-   return 0u;
-}
-#endif
diff --git a/src/gallium/frontends/clover/spirv/invocation.hpp b/src/gallium/frontends/clover/spirv/invocation.hpp
deleted file mode 100644
index 50b0c085b4b45..0000000000000
--- a/src/gallium/frontends/clover/spirv/invocation.hpp
+++ /dev/null
@@ -1,81 +0,0 @@
-//
-// Copyright 2018 Pierre Moreau
-//
-// 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.
-//
-
-#ifndef CLOVER_SPIRV_INVOCATION_HPP
-#define CLOVER_SPIRV_INVOCATION_HPP
-
-#include <unordered_set>
-
-#include "core/context.hpp"
-#include "core/binary.hpp"
-#include "core/program.hpp"
-
-namespace clover {
-   namespace spirv {
-      // Returns whether the binary starts with the SPIR-V magic word.
-      //
-      // The first word is interpreted as little endian and big endian, but
-      // only one of them has to match.
-      bool is_binary_spirv(const std::string &binary);
-
-      // Returns whether the given binary is considered valid for the given
-      // OpenCL version.
-      //
-      // It uses SPIRV-Tools validator to do the validation, and potential
-      // warnings and errors are appended to |r_log|.
-      bool is_valid_spirv(const std::string &binary,
-                          const cl_version opencl_version,
-                          std::string &r_log);
-
-      // Converts an integer SPIR-V version into its textual representation.
-      std::string version_to_string(uint32_t version);
-
-      // Creates a clover binary out of the given SPIR-V binary.
-      binary compile_program(const std::string &binary,
-                             const device &dev, std::string &r_log,
-                             bool validate = true);
-
-      // Combines multiple clover objects into a single one, resolving
-      // link dependencies between them.
-      binary link_program(const std::vector<binary> &objects, const device &dev,
-                          const std::string &opts, std::string &r_log);
-
-      // Returns a textual representation of the given binary.
-      std::string print_module(const std::string &binary,
-                               const cl_version opencl_version);
-
-      // Returns a set of supported SPIR-V extensions.
-      std::unordered_set<std::string> supported_extensions();
-
-      // Returns a vector (sorted in increasing order) of supported SPIR-V
-      // versions.
-      std::vector<cl_name_version> supported_versions();
-
-      // Converts a version number from SPIR-V's encoding to OpenCL's one.
-      cl_version to_opencl_version_encoding(uint32_t version);
-
-      // Converts a version number from OpenCL's encoding to SPIR-V's one.
-      uint32_t to_spirv_version_encoding(cl_version version);
-   }
-}
-
-#endif
-- 
GitLab

From 43754900661ac7eae816d707cb942c87f8188d8a Mon Sep 17 00:00:00 2001
From: Karol Herbst <kherbst@redhat.com>
Date: Sun, 25 Feb 2024 19:38:42 +0100
Subject: [PATCH 3/4] gallium: drop non clover drivers from dynamic pipe-loader

Every other frontend moved to the static pipe-loader, so none of this is
needed anymore. Shaves of around 50MiB from a mesa install.

Signed-off-by: Karol Herbst <kherbst@redhat.com>
---
 src/gallium/targets/pipe-loader/meson.build   | 23 +--------
 src/gallium/targets/pipe-loader/pipe_crocus.c |  5 --
 src/gallium/targets/pipe-loader/pipe_i915.c   |  5 --
 src/gallium/targets/pipe-loader/pipe_iris.c   |  5 --
 src/gallium/targets/pipe-loader/pipe_kmsro.c  |  6 ---
 src/gallium/targets/pipe-loader/pipe_msm.c    |  5 --
 .../targets/pipe-loader/pipe_nouveau.c        |  5 --
 src/gallium/targets/pipe-loader/pipe_r300.c   |  5 --
 src/gallium/targets/pipe-loader/pipe_swrast.c | 51 -------------------
 src/gallium/targets/pipe-loader/pipe_vmwgfx.c |  6 ---
 10 files changed, 2 insertions(+), 114 deletions(-)
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_crocus.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_i915.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_iris.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_kmsro.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_msm.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_nouveau.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_r300.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_swrast.c
 delete mode 100644 src/gallium/targets/pipe-loader/pipe_vmwgfx.c

diff --git a/src/gallium/targets/pipe-loader/meson.build b/src/gallium/targets/pipe-loader/meson.build
index 943faec469d0f..ca270cff5e2e7 100644
--- a/src/gallium/targets/pipe-loader/meson.build
+++ b/src/gallium/targets/pipe-loader/meson.build
@@ -39,27 +39,9 @@ endif
 
 pipe_loader_install_dir = join_paths(get_option('libdir'), 'gallium-pipe')
 
-_kmsro_targets = [
-   driver_kmsro, driver_v3d, driver_vc4, driver_freedreno, driver_etnaviv,
-   driver_panfrost, driver_lima, driver_asahi,
-]
-
-if with_gallium_v3d
-   _kmsro_targets += [idep_xmlconfig, dep_expat]
-endif
-
 pipe_loaders = [
-  [with_gallium_i915, 'i915', driver_i915, []],
-  [with_gallium_crocus, 'crocus', [driver_crocus, idep_xmlconfig], []],
-  [with_gallium_iris, 'iris', [driver_iris, idep_xmlconfig], []],
-  [with_gallium_nouveau, 'nouveau', driver_nouveau, []],
-  [with_gallium_r300, 'r300', driver_r300, []],
   [with_gallium_r600, 'r600', driver_r600, []],
   [with_gallium_radeonsi, 'radeonsi', [driver_radeonsi, idep_xmlconfig], []],
-  [with_gallium_freedreno, 'msm', driver_freedreno, []],
-  [with_gallium_kmsro, 'kmsro', _kmsro_targets, [libpipe_loader_dynamic]],
-  [with_gallium_svga, 'vmwgfx', driver_svga, []],
-  [with_gallium_softpipe, 'swrast', driver_swrast, [libwsw, libws_null, libswdri, libswkmsdri]],
 ]
 
 foreach x : pipe_loaders
@@ -69,11 +51,10 @@ foreach x : pipe_loaders
 
   pipe_sym_config = configuration_data()
 
-  foreach d : [[x[1] in ['r300', 'r600', 'radeonsi'], 'radeon_drm_winsys_create'],
+  foreach d : [[x[1] in ['r600', 'radeonsi'], 'radeon_drm_winsys_create'],
                [x[1] == 'radeonsi', 'amdgpu_winsys_create'],
                [x[1] == 'radeonsi' and with_llvm, 'ac_init_shared_llvm_once'],
-               [x[1] != 'swrast', 'driver_descriptor'],
-               [x[1] == 'swrast', 'swrast_driver_descriptor']]
+               [true, 'driver_descriptor']]
     if d[0]
       pipe_sym_config.set(d[1], d[1] + ';')
     else
diff --git a/src/gallium/targets/pipe-loader/pipe_crocus.c b/src/gallium/targets/pipe-loader/pipe_crocus.c
deleted file mode 100644
index 1664939421e0d..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_crocus.c
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "frontend/drm_driver.h"
-#include "target-helpers/drm_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "crocus/drm/crocus_drm_public.h"
-#include "util/driconf.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_i915.c b/src/gallium/targets/pipe-loader/pipe_i915.c
deleted file mode 100644
index 82daf7083e5ea..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_i915.c
+++ /dev/null
@@ -1,5 +0,0 @@
-
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/drm_driver.h"
-#include "i915/drm/i915_drm_public.h"
-#include "i915/i915_public.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_iris.c b/src/gallium/targets/pipe-loader/pipe_iris.c
deleted file mode 100644
index d7e85a0b6c28f..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_iris.c
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "frontend/drm_driver.h"
-#include "target-helpers/drm_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "iris/drm/iris_drm_public.h"
-#include "util/driconf.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_kmsro.c b/src/gallium/targets/pipe-loader/pipe_kmsro.c
deleted file mode 100644
index 4110517d923d7..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_kmsro.c
+++ /dev/null
@@ -1,6 +0,0 @@
-
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/drm_driver.h"
-#include "kmsro/drm/kmsro_drm_public.h"
-#define GALLIUM_KMSRO_ONLY
-#include "target-helpers/drm_helper.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_msm.c b/src/gallium/targets/pipe-loader/pipe_msm.c
deleted file mode 100644
index 807ac109bb715..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_msm.c
+++ /dev/null
@@ -1,5 +0,0 @@
-
-#include "target-helpers/drm_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/drm_driver.h"
-#include "freedreno/drm/freedreno_drm_public.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_nouveau.c b/src/gallium/targets/pipe-loader/pipe_nouveau.c
deleted file mode 100644
index 9ff6b82329f43..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_nouveau.c
+++ /dev/null
@@ -1,5 +0,0 @@
-
-#include "target-helpers/drm_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/drm_driver.h"
-#include "nouveau/drm/nouveau_drm_public.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_r300.c b/src/gallium/targets/pipe-loader/pipe_r300.c
deleted file mode 100644
index a348c6744615e..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_r300.c
+++ /dev/null
@@ -1,5 +0,0 @@
-#include "target-helpers/drm_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/drm_driver.h"
-#include "winsys/radeon_winsys.h"
-#include "r300/r300_public.h"
diff --git a/src/gallium/targets/pipe-loader/pipe_swrast.c b/src/gallium/targets/pipe-loader/pipe_swrast.c
deleted file mode 100644
index 583dc3ad705de..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_swrast.c
+++ /dev/null
@@ -1,51 +0,0 @@
-
-#include "target-helpers/inline_sw_helper.h"
-#include "target-helpers/inline_debug_helper.h"
-#include "frontend/sw_driver.h"
-#include "sw/dri/dri_sw_winsys.h"
-#include "sw/kms-dri/kms_dri_sw_winsys.h"
-#include "sw/null/null_sw_winsys.h"
-#include "sw/wrapper/wrapper_sw_winsys.h"
-
-PUBLIC struct pipe_screen *
-swrast_create_screen(struct sw_winsys *ws, const struct pipe_screen_config *config, bool sw_vk);
-
-struct pipe_screen *
-swrast_create_screen(struct sw_winsys *ws, const struct pipe_screen_config *config, bool sw_vk)
-{
-   struct pipe_screen *screen;
-
-   screen = sw_screen_create(ws);
-   if (screen)
-      screen = debug_screen_wrap(screen);
-
-   return screen;
-}
-
-PUBLIC
-const struct sw_driver_descriptor swrast_driver_descriptor = {
-   .create_screen = swrast_create_screen,
-   .winsys = {
-#ifdef HAVE_DRI
-      {
-         .name = "dri",
-         .create_winsys_dri = dri_create_sw_winsys,
-      },
-#endif
-#ifdef HAVE_DRISW_KMS
-      {
-         .name = "kms_dri",
-         .create_winsys_kms_dri = kms_dri_create_winsys,
-      },
-#endif
-      {
-         .name = "null",
-         .create_winsys = null_sw_create,
-      },
-      {
-         .name = "wrapped",
-         .create_winsys_wrapped = wrapper_sw_winsys_wrap_pipe_screen,
-      },
-      { 0 },
-   }
-};
diff --git a/src/gallium/targets/pipe-loader/pipe_vmwgfx.c b/src/gallium/targets/pipe-loader/pipe_vmwgfx.c
deleted file mode 100644
index 33d08071def3b..0000000000000
--- a/src/gallium/targets/pipe-loader/pipe_vmwgfx.c
+++ /dev/null
@@ -1,6 +0,0 @@
-
-#include "target-helpers/inline_debug_helper.h"
-#include "target-helpers/drm_helper.h"
-#include "frontend/drm_driver.h"
-#include "svga/drm/svga_drm_public.h"
-#include "svga/svga_public.h"
-- 
GitLab

From 91f4f82ebe525723c0bc5fc7e06f1f9e86df70ef Mon Sep 17 00:00:00 2001
From: Karol Herbst <kherbst@redhat.com>
Date: Sun, 25 Feb 2024 19:48:03 +0100
Subject: [PATCH 4/4] gallium: drop PIPE_SHADER_IR_NIR_SERIALIZED

It's not used anymore

Signed-off-by: Karol Herbst <kherbst@redhat.com>
---
 .../drivers/freedreno/freedreno_screen.c        |  2 --
 src/gallium/drivers/freedreno/ir3/ir3_gallium.c | 10 ----------
 src/gallium/drivers/iris/iris_program.c         |  8 --------
 src/gallium/drivers/iris/iris_screen.c          | 17 ++---------------
 src/gallium/drivers/llvmpipe/lp_screen.c        |  3 +--
 src/gallium/drivers/llvmpipe/lp_state_cs.c      |  8 --------
 src/gallium/drivers/nouveau/nouveau_screen.c    |  1 -
 src/gallium/drivers/nouveau/nouveau_screen.h    |  1 -
 src/gallium/drivers/nouveau/nvc0/nvc0_screen.c  |  8 ++------
 src/gallium/drivers/nouveau/nvc0/nvc0_state.c   |  8 --------
 src/gallium/include/pipe/p_defines.h            |  1 -
 11 files changed, 5 insertions(+), 62 deletions(-)

diff --git a/src/gallium/drivers/freedreno/freedreno_screen.c b/src/gallium/drivers/freedreno/freedreno_screen.c
index d04ab107eefb4..3f3eb91bb4938 100644
--- a/src/gallium/drivers/freedreno/freedreno_screen.c
+++ b/src/gallium/drivers/freedreno/freedreno_screen.c
@@ -762,8 +762,6 @@ fd_screen_get_shader_param(struct pipe_screen *pscreen,
       return 16;
    case PIPE_SHADER_CAP_SUPPORTED_IRS:
       return (1 << PIPE_SHADER_IR_NIR) |
-             COND(has_compute(screen) && (shader == PIPE_SHADER_COMPUTE),
-                  (1 << PIPE_SHADER_IR_NIR_SERIALIZED)) |
              /* tgsi_to_nir doesn't support all stages: */
              COND((shader == PIPE_SHADER_VERTEX) ||
                   (shader == PIPE_SHADER_FRAGMENT) ||
diff --git a/src/gallium/drivers/freedreno/ir3/ir3_gallium.c b/src/gallium/drivers/freedreno/ir3/ir3_gallium.c
index cf1bf5c6f2f20..509ef96e96f3a 100644
--- a/src/gallium/drivers/freedreno/ir3/ir3_gallium.c
+++ b/src/gallium/drivers/freedreno/ir3/ir3_gallium.c
@@ -299,16 +299,6 @@ ir3_shader_compute_state_create(struct pipe_context *pctx,
    if (cso->ir_type == PIPE_SHADER_IR_NIR) {
       /* we take ownership of the reference: */
       nir = (nir_shader *)cso->prog;
-   } else if (cso->ir_type == PIPE_SHADER_IR_NIR_SERIALIZED) {
-      const nir_shader_compiler_options *options =
-            ir3_get_compiler_options(compiler);
-      const struct pipe_binary_program_header *hdr = cso->prog;
-      struct blob_reader reader;
-
-      blob_reader_init(&reader, hdr->blob, hdr->num_bytes);
-      nir = nir_deserialize(NULL, options, &reader);
-
-      ir3_finalize_nir(compiler, nir);
    } else {
       assert(cso->ir_type == PIPE_SHADER_IR_TGSI);
       if (ir3_shader_debug & IR3_DBG_DISASM) {
diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c
index ad7b927f13d2c..5634aad7bb7e9 100644
--- a/src/gallium/drivers/iris/iris_program.c
+++ b/src/gallium/drivers/iris/iris_program.c
@@ -3242,14 +3242,6 @@ iris_create_compute_state(struct pipe_context *ctx,
       nir = (void *)state->prog;
       break;
 
-   case PIPE_SHADER_IR_NIR_SERIALIZED: {
-      struct blob_reader reader;
-      const struct pipe_binary_program_header *hdr = state->prog;
-      blob_reader_init(&reader, hdr->blob, hdr->num_bytes);
-      nir = nir_deserialize(NULL, options, &reader);
-      break;
-   }
-
    default:
       unreachable("Unsupported IR");
    }
diff --git a/src/gallium/drivers/iris/iris_screen.c b/src/gallium/drivers/iris/iris_screen.c
index 25163d460a665..abe3c40dc4711 100644
--- a/src/gallium/drivers/iris/iris_screen.c
+++ b/src/gallium/drivers/iris/iris_screen.c
@@ -111,15 +111,6 @@ iris_get_driver_uuid(struct pipe_screen *pscreen, char *uuid)
    intel_uuid_compute_driver_id((uint8_t *)uuid, devinfo, PIPE_UUID_SIZE);
 }
 
-static bool
-iris_enable_clover()
-{
-   static int enable = -1;
-   if (enable < 0)
-      enable = debug_get_bool_option("IRIS_ENABLE_CLOVER", false);
-   return enable;
-}
-
 static void
 iris_warn_cl()
 {
@@ -546,12 +537,8 @@ iris_get_shader_param(struct pipe_screen *pscreen,
    case PIPE_SHADER_CAP_MAX_HW_ATOMIC_COUNTERS:
    case PIPE_SHADER_CAP_MAX_HW_ATOMIC_COUNTER_BUFFERS:
       return 0;
-   case PIPE_SHADER_CAP_SUPPORTED_IRS: {
-      int irs = 1 << PIPE_SHADER_IR_NIR;
-      if (iris_enable_clover())
-         irs |= 1 << PIPE_SHADER_IR_NIR_SERIALIZED;
-      return irs;
-   }
+   case PIPE_SHADER_CAP_SUPPORTED_IRS:
+      return 1 << PIPE_SHADER_IR_NIR;
    case PIPE_SHADER_CAP_TGSI_ANY_INOUT_DECL_RANGE:
    case PIPE_SHADER_CAP_TGSI_SQRT_SUPPORTED:
       return 0;
diff --git a/src/gallium/drivers/llvmpipe/lp_screen.c b/src/gallium/drivers/llvmpipe/lp_screen.c
index ad3d66424e19d..bbd11dd1754f0 100644
--- a/src/gallium/drivers/llvmpipe/lp_screen.c
+++ b/src/gallium/drivers/llvmpipe/lp_screen.c
@@ -380,8 +380,7 @@ llvmpipe_get_shader_param(struct pipe_screen *screen,
    case PIPE_SHADER_COMPUTE:
       if ((lscreen->allow_cl) && param == PIPE_SHADER_CAP_SUPPORTED_IRS)
          return ((1 << PIPE_SHADER_IR_TGSI) |
-                 (1 << PIPE_SHADER_IR_NIR) |
-                 (1 << PIPE_SHADER_IR_NIR_SERIALIZED));
+                 (1 << PIPE_SHADER_IR_NIR));
       FALLTHROUGH;
    case PIPE_SHADER_MESH:
    case PIPE_SHADER_TASK:
diff --git a/src/gallium/drivers/llvmpipe/lp_state_cs.c b/src/gallium/drivers/llvmpipe/lp_state_cs.c
index 62fa37c153739..cf9b5d0e226d8 100644
--- a/src/gallium/drivers/llvmpipe/lp_state_cs.c
+++ b/src/gallium/drivers/llvmpipe/lp_state_cs.c
@@ -921,14 +921,6 @@ llvmpipe_create_compute_state(struct pipe_context *pipe,
 
    if (templ->ir_type == PIPE_SHADER_IR_TGSI) {
       shader->base.ir.nir = tgsi_to_nir(templ->prog, pipe->screen, false);
-   } else if (templ->ir_type == PIPE_SHADER_IR_NIR_SERIALIZED) {
-      struct blob_reader reader;
-      const struct pipe_binary_program_header *hdr = templ->prog;
-
-      blob_reader_init(&reader, hdr->blob, hdr->num_bytes);
-      shader->base.ir.nir = nir_deserialize(NULL, pipe->screen->get_compiler_options(pipe->screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE), &reader);
-
-      pipe->screen->finalize_nir(pipe->screen, shader->base.ir.nir);
    } else if (templ->ir_type == PIPE_SHADER_IR_NIR) {
       shader->base.ir.nir = (struct nir_shader *)templ->prog;
    }
diff --git a/src/gallium/drivers/nouveau/nouveau_screen.c b/src/gallium/drivers/nouveau/nouveau_screen.c
index 2292420451417..d0c8e60ac55bb 100644
--- a/src/gallium/drivers/nouveau/nouveau_screen.c
+++ b/src/gallium/drivers/nouveau/nouveau_screen.c
@@ -295,7 +295,6 @@ nouveau_screen_init(struct nouveau_screen *screen, struct nouveau_device *dev)
    if (nv_dbg)
       nouveau_mesa_debug = atoi(nv_dbg);
 
-   screen->force_enable_cl = debug_get_bool_option("NOUVEAU_ENABLE_CL", false);
    screen->disable_fences = debug_get_bool_option("NOUVEAU_DISABLE_FENCES", false);
 
    /* These must be set before any failure is possible, as the cleanup
diff --git a/src/gallium/drivers/nouveau/nouveau_screen.h b/src/gallium/drivers/nouveau/nouveau_screen.h
index 87a54c2c2f05a..ca0fb382f8923 100644
--- a/src/gallium/drivers/nouveau/nouveau_screen.h
+++ b/src/gallium/drivers/nouveau/nouveau_screen.h
@@ -65,7 +65,6 @@ struct nouveau_screen {
 
    struct disk_cache *disk_shader_cache;
 
-   bool force_enable_cl;
    bool has_svm;
    bool is_uma;
    bool disable_fences;
diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c
index 28a616a621fce..58e425ac34e2b 100644
--- a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c
+++ b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c
@@ -396,12 +396,8 @@ nvc0_screen_get_shader_param(struct pipe_screen *pscreen,
    }
 
    switch (param) {
-   case PIPE_SHADER_CAP_SUPPORTED_IRS: {
-      uint32_t irs = 1 << PIPE_SHADER_IR_NIR;
-      if (screen->force_enable_cl)
-         irs |= 1 << PIPE_SHADER_IR_NIR_SERIALIZED;
-      return irs;
-   }
+   case PIPE_SHADER_CAP_SUPPORTED_IRS:
+      return 1 << PIPE_SHADER_IR_NIR;
    case PIPE_SHADER_CAP_MAX_INSTRUCTIONS:
    case PIPE_SHADER_CAP_MAX_ALU_INSTRUCTIONS:
    case PIPE_SHADER_CAP_MAX_TEX_INSTRUCTIONS:
diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_state.c b/src/gallium/drivers/nouveau/nvc0/nvc0_state.c
index c7a22742fee3d..b7f6d6f1d28d1 100644
--- a/src/gallium/drivers/nouveau/nvc0/nvc0_state.c
+++ b/src/gallium/drivers/nouveau/nvc0/nvc0_state.c
@@ -748,14 +748,6 @@ nvc0_cp_state_create(struct pipe_context *pipe,
    case PIPE_SHADER_IR_NIR:
       prog->nir = (nir_shader *)cso->prog;
       break;
-   case PIPE_SHADER_IR_NIR_SERIALIZED: {
-      struct blob_reader reader;
-      const struct pipe_binary_program_header *hdr = cso->prog;
-
-      blob_reader_init(&reader, hdr->blob, hdr->num_bytes);
-      prog->nir = nir_deserialize(NULL, pipe->screen->get_compiler_options(pipe->screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE), &reader);
-      break;
-   }
    default:
       assert(!"unsupported IR!");
       free(prog);
diff --git a/src/gallium/include/pipe/p_defines.h b/src/gallium/include/pipe/p_defines.h
index ee6d894a65a88..99e53154df7bd 100644
--- a/src/gallium/include/pipe/p_defines.h
+++ b/src/gallium/include/pipe/p_defines.h
@@ -1054,7 +1054,6 @@ enum pipe_shader_ir
    PIPE_SHADER_IR_TGSI = 0,
    PIPE_SHADER_IR_NATIVE,
    PIPE_SHADER_IR_NIR,
-   PIPE_SHADER_IR_NIR_SERIALIZED,
 };
 
 /**
-- 
GitLab

Places

File Mesa-MR27783_drop-nir-in-clover.patch of Package Mesa (Revision 03d9a31ac43184a49cbd3b74c5e56ac7)

Places