fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-02 05:48:07 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

microsoft/clc: Split clc_object and rename entrypoints

Browse source 
clc_object was overloaded, containing SPIR-V binary and metadata,
and it would only sometimes contain metadata (after linking). Split
it into a more generic clc_binary class which holds some type of data
(the kind depends on where it came from), and clc_metadata which can
be independently parsed on compiled or linked data.

Rename a couple entrypoints to be more explicit about what they're
actually transforming (c_to_spirv, link_spirv, spirv_to_dxil).

Add a logger to SPIR-V binary parsing so it can report errors on app-
provided SPIR-V.

Re-order helper function parameters to be more consistent (out params last).

Acked-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10322>
This commit is contained in:
Jesse Natalie 2021-04-19 06:54:13 -07:00 • committed by Marge Bot
parent c3666eec7e
commit 91e08312d8
8 changed files with 158 additions and 106 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

75
src/microsoft/clc/clc_compiler.c
View file

@ -54,7 +54,7 @@ static const struct debug_named_value clc_debug_options[] = {
DEBUG_GET_ONCE_FLAGS_OPTION(debug_clc, "CLC_DEBUG", clc_debug_options, 0)
static void
clc_print_kernels_info(const struct clc_object *obj)
clc_print_kernels_info(const struct clc_parsed_spirv *obj)
{
fprintf(stdout, "Kernels:\n");
for (unsigned i = 0; i < obj->num_kernels; i++) {
@ -575,44 +575,60 @@ struct clc_libclc *
return ctx;
}
bool
clc_compile(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_object *out_spirv)
void
clc_free_spirv(struct clc_binary *spirv)
{
if (clc_to_spirv(args, &out_spirv->spvbin, logger) < 0)
clc_free_spirv_binary(spirv);
}
bool
clc_compile_c_to_spirv(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv)
{
if (clc_c_to_spirv(args, logger, out_spirv) < 0)
return false;
if (debug_get_option_debug_clc() & CLC_DEBUG_DUMP_SPIRV)
clc_dump_spirv(&out_spirv->spvbin, stdout);
clc_dump_spirv(out_spirv, stdout);
return true;
}
bool
clc_link(const struct clc_linker_args *args,
const struct clc_logger *logger,
struct clc_object *out_spirv)
clc_link_spirv(const struct clc_linker_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv)
{
if (clc_link_spirv_binaries(args, &out_spirv->spvbin, logger) < 0)
if (clc_link_spirv_binaries(args, logger, out_spirv) < 0)
return false;
if (debug_get_option_debug_clc() & CLC_DEBUG_DUMP_SPIRV)
clc_dump_spirv(&out_spirv->spvbin, stdout);
clc_dump_spirv(out_spirv, stdout);
out_spirv->kernels = clc_spirv_get_kernels_info(&out_spirv->spvbin,
&out_spirv->num_kernels);
return true;
}
bool
clc_parse_spirv(const struct clc_binary *in_spirv,
const struct clc_logger *logger,
struct clc_parsed_spirv *out_data)
{
if (!clc_spirv_get_kernels_info(in_spirv,
&out_data->kernels,
&out_data->num_kernels,
logger))
return false;
if (debug_get_option_debug_clc() & CLC_DEBUG_VERBOSE)
clc_print_kernels_info(out_spirv);
clc_print_kernels_info(out_data);
return true;
}
void clc_free_object(struct clc_object *obj)
void clc_free_parsed_spirv(struct clc_parsed_spirv *data)
{
clc_free_kernels_info(obj->kernels, obj->num_kernels);
clc_free_spirv_binary(&obj->spvbin);
clc_free_kernels_info(data->kernels, data->num_kernels);
}
static nir_variable *
@ -989,18 +1005,19 @@ scale_fdiv(nir_shader *nir)
}
bool
clc_to_dxil(struct clc_libclc *lib,
const struct clc_object *obj,
const char *entrypoint,
const struct clc_runtime_kernel_conf *conf,
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil)
clc_spirv_to_dxil(struct clc_libclc *lib,
const struct clc_binary *linked_spirv,
const struct clc_parsed_spirv *parsed_data,
const char *entrypoint,
const struct clc_runtime_kernel_conf *conf,
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil)
{
struct nir_shader *nir;
for (unsigned i = 0; i < obj->num_kernels; i++) {
if (!strcmp(obj->kernels[i].name, entrypoint)) {
out_dxil->kernel = &obj->kernels[i];
for (unsigned i = 0; i < parsed_data->num_kernels; i++) {
if (!strcmp(parsed_data->kernels[i].name, entrypoint)) {
out_dxil->kernel = &parsed_data->kernels[i];
break;
}
}
@ -1045,7 +1062,7 @@ clc_to_dxil(struct clc_libclc *lib,
glsl_type_singleton_init_or_ref();
nir = spirv_to_nir(obj->spvbin.data, obj->spvbin.size / 4,
nir = spirv_to_nir(linked_spirv->data, linked_spirv->size / 4,
NULL, 0,
MESA_SHADER_KERNEL, entrypoint,
&spirv_options,
@ -1374,7 +1391,7 @@ clc_to_dxil(struct clc_libclc *lib,
continue;
/* If we don't have the runtime conf yet, we just create a dummy variable.
* This will be adjusted when clc_to_dxil() is called with a conf
* This will be adjusted when clc_spirv_to_dxil() is called with a conf
* argument.
*/
unsigned size = 4;

45
src/microsoft/clc/clc_compiler.h
View file

@ -45,7 +45,7 @@ struct clc_compile_args {
};
struct clc_linker_args {
const struct clc_object * const *in_objs;
const struct clc_binary * const *in_objs;
unsigned num_in_objs;
unsigned create_library;
};
@ -58,8 +58,8 @@ struct clc_logger {
clc_msg_callback warning;
};
struct spirv_binary {
uint32_t *data;
struct clc_binary {
void *data;
size_t size;
};
@ -108,8 +108,7 @@ struct clc_kernel_info {
enum clc_vec_hint_type vec_hint_type;
};
struct clc_object {
struct spirv_binary spvbin;
struct clc_parsed_spirv {
const struct clc_kernel_info *kernels;
unsigned num_kernels;
};
@ -201,19 +200,26 @@ void clc_libclc_serialize(struct clc_libclc *lib, void **serialized, size_t *siz
void clc_libclc_free_serialized(void *serialized);
struct clc_libclc *clc_libclc_deserialize(void *serialized, size_t size);
void
clc_free_spirv(struct clc_binary *spirv);
bool
clc_compile(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_object *out_spirv);
clc_compile_c_to_spirv(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv);
bool
clc_link(const struct clc_linker_args *args,
const struct clc_logger *logger,
struct clc_object *out_spirv);
clc_link_spirv(const struct clc_linker_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv);
void clc_free_object(struct clc_object *obj);
bool
clc_parse_spirv(const struct clc_binary *in_spirv,
const struct clc_logger *logger,
struct clc_parsed_spirv *out_data);
void
clc_free_parsed_spirv(struct clc_parsed_spirv *data);
struct clc_runtime_arg_info {
union {
@ -237,12 +243,13 @@ struct clc_runtime_kernel_conf {
};
bool
clc_to_dxil(struct clc_libclc *ctx,
const struct clc_object *obj,
const char *entrypoint,
const struct clc_runtime_kernel_conf *conf,
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil);
clc_spirv_to_dxil(struct clc_libclc *lib,
const struct clc_binary *linked_spirv,
const struct clc_parsed_spirv *parsed_data,
const char *entrypoint,
const struct clc_runtime_kernel_conf *conf,
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil);
void clc_free_dxil_object(struct clc_dxil_object *dxil);

10
src/microsoft/clc/clc_compiler_test.cpp
View file

@ -1688,8 +1688,8 @@ TEST_F(ComputeTest, vec_hint_float4)
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.obj->kernels[0].vec_hint_size, 4);
EXPECT_EQ(shader.obj->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_FLOAT);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 4);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_FLOAT);
}
TEST_F(ComputeTest, vec_hint_uchar2)
@ -1700,8 +1700,8 @@ TEST_F(ComputeTest, vec_hint_uchar2)
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.obj->kernels[0].vec_hint_size, 2);
EXPECT_EQ(shader.obj->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_CHAR);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 2);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_CHAR);
}
TEST_F(ComputeTest, vec_hint_none)
@ -1712,7 +1712,7 @@ TEST_F(ComputeTest, vec_hint_none)
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.obj->kernels[0].vec_hint_size, 0);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 0);
}
TEST_F(ComputeTest, DISABLED_debug_layer_failure)

70
src/microsoft/clc/clc_helpers.cpp
View file

@ -474,7 +474,7 @@ public:
return true;
}
void parseBinary(const struct spirv_binary &spvbin)
bool parseBinary(const struct clc_binary &spvbin, const struct clc_logger *logger)
{
/* 3 passes should be enough to retrieve all kernel information:
* 1st pass: all entry point name and number of args
@ -482,15 +482,23 @@ public:
* 3rd pass: pointer type names
*/
for (unsigned pass = 0; pass < 3; pass++) {
spvBinaryParse(ctx, reinterpret_cast<void *>(this),
spvbin.data, spvbin.size / 4,
NULL, parseInstruction, NULL);
spv_diagnostic diagnostic = NULL;
auto result = spvBinaryParse(ctx, reinterpret_cast<void *>(this),
static_cast<uint32_t*>(spvbin.data), spvbin.size / 4,
NULL, parseInstruction, &diagnostic);
if (result != SPV_SUCCESS) {
if (diagnostic && logger)
logger->error(logger->priv, diagnostic->error);
return false;
}
if (parsingComplete())
return;
return true;
}
assert(0);
return false;
}
std::vector<SPIRVKernelInfo> kernels;
@ -499,18 +507,22 @@ public:
spv_context ctx;
};
const struct clc_kernel_info *
clc_spirv_get_kernels_info(const struct spirv_binary *spvbin,
unsigned *num_kernels)
bool
clc_spirv_get_kernels_info(const struct clc_binary *spvbin,
const struct clc_kernel_info **out_kernels,
unsigned *num_kernels,
const struct clc_logger *logger)
{
struct clc_kernel_info *kernels;
SPIRVKernelParser parser;
parser.parseBinary(*spvbin);
if (!parser.parseBinary(*spvbin, logger))
return false;
*num_kernels = parser.kernels.size();
if (!*num_kernels)
return NULL;
return false;
kernels = reinterpret_cast<struct clc_kernel_info *>(calloc(*num_kernels,
sizeof(*kernels)));
@ -539,7 +551,9 @@ clc_spirv_get_kernels_info(const struct spirv_binary *spvbin,
}
}
return kernels;
*out_kernels = kernels;
return true;
}
void
@ -563,9 +577,9 @@ clc_free_kernels_info(const struct clc_kernel_info *kernels,
}
int
clc_to_spirv(const struct clc_compile_args *args,
struct spirv_binary *spvbin,
const struct clc_logger *logger)
clc_c_to_spirv(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv)
{
LLVMInitializeAllTargets();
LLVMInitializeAllTargetInfos();
@ -694,9 +708,9 @@ clc_to_spirv(const struct clc_compile_args *args,
}
const std::string spv_out = spv_stream.str();
spvbin->size = spv_out.size();
spvbin->data = static_cast<uint32_t *>(malloc(spvbin->size));
memcpy(spvbin->data, spv_out.data(), spvbin->size);
out_spirv->size = spv_out.size();
out_spirv->data = malloc(out_spirv->size);
memcpy(out_spirv->data, spv_out.data(), out_spirv->size);
return 0;
}
@ -762,15 +776,14 @@ private:
int
clc_link_spirv_binaries(const struct clc_linker_args *args,
struct spirv_binary *dst_bin,
const struct clc_logger *logger)
const struct clc_logger *logger,
struct clc_binary *out_spirv)
{
std::vector<std::vector<uint32_t>> binaries;
for (unsigned i = 0; i < args->num_in_objs; i++) {
std::vector<uint32_t> bin(args->in_objs[i]->spvbin.data,
args->in_objs[i]->spvbin.data +
(args->in_objs[i]->spvbin.size / 4));
const uint32_t *data = static_cast<const uint32_t *>(args->in_objs[i]->data);
std::vector<uint32_t> bin(data, data + (args->in_objs[i]->size / 4));
binaries.push_back(bin);
}
@ -786,18 +799,19 @@ clc_link_spirv_binaries(const struct clc_linker_args *args,
return -1;
}
dst_bin->size = linkingResult.size() * 4;
dst_bin->data = static_cast<uint32_t *>(malloc(dst_bin->size));
memcpy(dst_bin->data, linkingResult.data(), dst_bin->size);
out_spirv->size = linkingResult.size() * 4;
out_spirv->data = static_cast<uint32_t *>(malloc(out_spirv->size));
memcpy(out_spirv->data, linkingResult.data(), out_spirv->size);
return 0;
}
void
clc_dump_spirv(const struct spirv_binary *spvbin, FILE *f)
clc_dump_spirv(const struct clc_binary *spvbin, FILE *f)
{
spvtools::SpirvTools tools(SPV_ENV_UNIVERSAL_1_0);
std::vector<uint32_t> bin(spvbin->data, spvbin->data + (spvbin->size / 4));
const uint32_t *data = static_cast<const uint32_t *>(spvbin->data);
std::vector<uint32_t> bin(data, data + (spvbin->size / 4));
std::string out;
tools.Disassemble(bin, &out,
SPV_BINARY_TO_TEXT_OPTION_INDENT |
@ -806,7 +820,7 @@ clc_dump_spirv(const struct spirv_binary *spvbin, FILE *f)
}
void
clc_free_spirv_binary(struct spirv_binary *spvbin)
clc_free_spirv_binary(struct clc_binary *spvbin)
{
free(spvbin->data);
}

22
src/microsoft/clc/clc_helpers.h
View file

@ -38,29 +38,31 @@ extern "C" {
#include <stdio.h>
#include <stdint.h>
const struct clc_kernel_info *
clc_spirv_get_kernels_info(const struct spirv_binary *spvbin,
unsigned *num_kernels);
bool
clc_spirv_get_kernels_info(const struct clc_binary *spvbin,
const struct clc_kernel_info **kernels,
unsigned *num_kernels,
const struct clc_logger *logger);
void
clc_free_kernels_info(const struct clc_kernel_info *kernels,
unsigned num_kernels);
int
clc_to_spirv(const struct clc_compile_args *args,
struct spirv_binary *spvbin,
const struct clc_logger *logger);
clc_c_to_spirv(const struct clc_compile_args *args,
const struct clc_logger *logger,
struct clc_binary *out_spirv);
int
clc_link_spirv_binaries(const struct clc_linker_args *args,
struct spirv_binary *dst_bin,
const struct clc_logger *logger);
const struct clc_logger *logger,
struct clc_binary *out_spirv);
void
clc_dump_spirv(const struct spirv_binary *spvbin, FILE *f);
clc_dump_spirv(const struct clc_binary *spvbin, FILE *f);
void
clc_free_spirv_binary(struct spirv_binary *spvbin);
clc_free_spirv_binary(struct clc_binary *spvbin);
#define clc_log(logger, level, fmt, ...) do { \
if (!logger || !logger->level) break; \

10
src/microsoft/clc/clon12compiler.def
View file

@ -4,9 +4,11 @@ EXPORTS
clc_libclc_serialize
clc_libclc_free_serialized
clc_libclc_deserialize
clc_compile
clc_link
clc_free_object
clc_to_dxil
clc_free_spirv
clc_compile_c_to_spirv
clc_link_spirv
clc_parse_spirv
clc_free_parsed_spirv
clc_spirv_to_dxil
clc_free_dxil_object
clc_compiler_get_version

29
src/microsoft/clc/compute_test.cpp
View file

@ -803,14 +803,14 @@ ComputeTest::compile(const std::vector<const char *> &sources,
for (unsigned i = 0; i < sources.size(); i++) {
args.source.value = sources[i];
clc_object spirv{};
if (!clc_compile(&args, &logger, &spirv))
clc_binary spirv{};
if (!clc_compile_c_to_spirv(&args, &logger, &spirv))
throw runtime_error("failed to compile object!");
Shader shader;
shader.obj = std::shared_ptr<clc_object>(new clc_object(spirv), [](clc_object *spirv)
shader.obj = std::shared_ptr<clc_binary>(new clc_binary(spirv), [](clc_binary *spirv)
{
clc_free_object(spirv);
clc_free_spirv(spirv);
delete spirv;
});
shaders.push_back(shader);
@ -826,7 +826,7 @@ ComputeTest::Shader
ComputeTest::link(const std::vector<Shader> &sources,
bool create_library)
{
std::vector<const clc_object*> objs;
std::vector<const clc_binary*> objs;
for (auto& source : sources)
objs.push_back(&*source.obj);
@ -834,14 +834,14 @@ ComputeTest::link(const std::vector<Shader> &sources,
link_args.in_objs = objs.data();
link_args.num_in_objs = (unsigned)objs.size();
link_args.create_library = create_library;
clc_object spirv{};
if (!clc_link(&link_args, &logger, &spirv))
clc_binary spirv{};
if (!clc_link_spirv(&link_args, &logger, &spirv))
throw runtime_error("failed to link objects!");
ComputeTest::Shader shader;
shader.obj = std::shared_ptr<clc_object>(new clc_object(spirv), [](clc_object *spirv)
shader.obj = std::shared_ptr<clc_binary>(new clc_binary(spirv), [](clc_binary *spirv)
{
clc_free_object(spirv);
clc_free_spirv(spirv);
delete spirv;
});
if (!link_args.create_library)
@ -854,13 +854,22 @@ void
ComputeTest::configure(Shader &shader,
const struct clc_runtime_kernel_conf *conf)
{
if (!shader.metadata) {
shader.metadata = std::shared_ptr<clc_parsed_spirv>(new clc_parsed_spirv{}, [](clc_parsed_spirv *metadata)
{
clc_free_parsed_spirv(metadata);
delete metadata;
});
if (!clc_parse_spirv(shader.obj.get(), NULL, shader.metadata.get()))
throw runtime_error("failed to parse spirv!");
}
shader.dxil = std::shared_ptr<clc_dxil_object>(new clc_dxil_object{}, [](clc_dxil_object *dxil)
{
clc_free_dxil_object(dxil);
delete dxil;
});
if (!clc_to_dxil(compiler_ctx, shader.obj.get(), "main_test", conf, &logger, shader.dxil.get()))
if (!clc_spirv_to_dxil(compiler_ctx, shader.obj.get(), shader.metadata.get(), "main_test", conf, &logger, shader.dxil.get()))
throw runtime_error("failed to compile kernel!");
}

3
src/microsoft/clc/compute_test.h
View file

@ -52,7 +52,8 @@ align(size_t value, unsigned alignment)
class ComputeTest : public ::testing::Test {
protected:
struct Shader {
std::shared_ptr<struct clc_object> obj;
std::shared_ptr<struct clc_binary> obj;
std::shared_ptr<struct clc_parsed_spirv> metadata;
std::shared_ptr<struct clc_dxil_object> dxil;
};