Sign Up
Log In
Log In
or
Sign Up
Places
All Projects
Status Monitor
Collapse sidebar
home:X0F:HSF
Mesa
Mesa-MR27783_drop-nir-in-clover.patch
Overview
Repositories
Revisions
Requests
Users
Attributes
Meta
File Mesa-MR27783_drop-nir-in-clover.patch of Package Mesa
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 ¬ify) { - 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 §ion = 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
Locations
Projects
Search
Status Monitor
Help
OpenBuildService.org
Documentation
API Documentation
Code of Conduct
Contact
Support
@OBShq
Terms
openSUSE Build Service is sponsored by
The Open Build Service is an
openSUSE project
.
Sign Up
Log In
Places
Places
All Projects
Status Monitor
