OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

clover: drop support for nir drivers

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

Acked-by: David Heidelberg <david@ixit.cz>
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Daniel Stone <daniels@collabora.com>
Signed-off-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27783>
This commit is contained in:
Karol Herbst
2024-02-25 19:25:03 +01:00
committed by Marge Bot
parent 277925471e
commit 80c4ffb61a
14 changed files with 8 additions and 1847 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
.gitlab-ci/build/gitlab-ci.yml
View File

@@ -657,7 +657,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,r300,r600,freedreno,llvmpipe,softpipe,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
@@ -695,7 +694,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:

1
.gitlab-ci/windows/mesa_build.ps1
View File

@@ -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 `

7
meson.build
View File

@@ -802,8 +802,7 @@ if with_gallium_rusticl or with_nouveau_vk or with_tools.contains('etnaviv')
endif
endif
with_clover_spirv = with_gallium_clover and get_option('opencl-spirv')
with_clc = with_microsoft_clc or with_intel_clc or with_gallium_asahi or with_asahi_vk or with_gallium_rusticl or with_clover_spirv
with_clc = with_microsoft_clc or with_intel_clc or with_gallium_asahi or with_asahi_vk or with_gallium_rusticl
dep_clc = null_dep
if with_gallium_clover or with_clc
@@ -1817,7 +1816,7 @@ pre_args += '-DDRAW_LLVM_AVAILABLE=@0@'.format((with_llvm and draw_with_llvm).to
pre_args += '-DAMD_LLVM_AVAILABLE=@0@'.format(amd_with_llvm.to_int())
pre_args += '-DGALLIVM_USE_ORCJIT=@0@'.format((with_llvm and llvm_with_orcjit).to_int())
if with_clover_spirv or with_clc
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()
@@ -1847,7 +1846,7 @@ endif
dep_spirv_tools = dependency(
'SPIRV-Tools',
required : with_clover_spirv or with_clc,
required : with_clc,
version : '>= 2022.1'
)
if dep_spirv_tools.found()

8
meson_options.txt
View File

@@ -184,14 +184,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',

18
src/gallium/frontends/clover/api/program.cpp
View File

@@ -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;
}
}

25
src/gallium/frontends/clover/core/compiler.hpp
View File

@@ -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:

32
src/gallium/frontends/clover/core/device.cpp
View File

@@ -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, false);
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 *

76
src/gallium/frontends/clover/llvm/invocation.cpp
View File

@@ -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"
@@ -407,30 +400,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
@@ -574,48 +543,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

8
src/gallium/frontends/clover/llvm/invocation.hpp
View File

@@ -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
}
}

35
src/gallium/frontends/clover/meson.build
View File

@@ -13,7 +13,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
@@ -27,10 +26,6 @@ if with_opencl_icd
clover_cpp_args += '-DHAVE_CLOVER_ICD'
endif
if with_clover_spirv
clover_spirv_cpp_args += '-DHAVE_CLOVER_SPIRV'
endif
libclllvm = static_library(
'clllvm',
files(
@@ -48,7 +43,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',
@@ -58,32 +52,6 @@ libclllvm = static_library(
dependencies : [dep_llvm, dep_elf, dep_llvmspirvlib, idep_mesautil],
)
idep_opencl_spirv = null_dep
if with_clover_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',
@@ -150,10 +118,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],
)

485
src/gallium/frontends/clover/nir/invocation.cpp
View File

@@ -1,485 +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 <compiler/spirv/spirv_info.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_capabilities
create_spirv_caps(const device &dev)
{
struct spirv_capabilities caps = {};
caps.Addresses = true;
caps.Float64 = true;
caps.Int8 = true;
caps.Int16 = true;
caps.Int64 = true;
caps.Kernel = true;
caps.ImageBasic = dev.image_support();
caps.Int64Atomics = dev.has_int64_atomics();
return caps;
}
static spirv_to_nir_options
create_spirv_options(const device &dev,
spirv_capabilities &caps,
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.capabilities = &caps;
spirv_options.debug.func = &debug_function;
spirv_options.debug.private_data = &r_log;
spirv_options.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_capabilities caps = create_spirv_caps(dev);
spirv_to_nir_options spirv_options = create_spirv_options(dev, caps, 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_capabilities caps = create_spirv_caps(dev);
spirv_to_nir_options spirv_options = create_spirv_options(dev, caps, 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

46
src/gallium/frontends/clover/nir/invocation.hpp
View File

@@ -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

1031
src/gallium/frontends/clover/spirv/invocation.cpp
View File

File diff suppressed because it is too large Load Diff

81
src/gallium/frontends/clover/spirv/invocation.hpp
View File

@@ -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