Module: Mesa Branch: main Commit: 91e08312d87dd303f7306d2eaa0f48ed008b8844 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=91e08312d87dd303f7306d2eaa0f48ed008b8844
Author: Jesse Natalie <[email protected]> Date: Mon Apr 19 06:54:13 2021 -0700 microsoft/clc: Split clc_object and rename entrypoints 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 <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10322> --- src/microsoft/clc/clc_compiler.c | 73 ++++++++++++++++++++------------- src/microsoft/clc/clc_compiler.h | 43 +++++++++++-------- src/microsoft/clc/clc_compiler_test.cpp | 10 ++--- src/microsoft/clc/clc_helpers.cpp | 70 ++++++++++++++++++------------- src/microsoft/clc/clc_helpers.h | 22 +++++----- src/microsoft/clc/clon12compiler.def | 10 +++-- src/microsoft/clc/compute_test.cpp | 29 ++++++++----- src/microsoft/clc/compute_test.h | 3 +- 8 files changed, 156 insertions(+), 104 deletions(-) diff --git a/src/microsoft/clc/clc_compiler.c b/src/microsoft/clc/clc_compiler.c index 4d7b8ddb4bb..3af1493efaa 100644 --- a/src/microsoft/clc/clc_compiler.c +++ b/src/microsoft/clc/clc_compiler.c @@ -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; } +void +clc_free_spirv(struct clc_binary *spirv) +{ + clc_free_spirv_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) { - if (clc_to_spirv(args, &out_spirv->spvbin, logger) < 0) + 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; diff --git a/src/microsoft/clc/clc_compiler.h b/src/microsoft/clc/clc_compiler.h index 58c7ceded12..3f2cfea2eeb 100644 --- a/src/microsoft/clc/clc_compiler.h +++ b/src/microsoft/clc/clc_compiler.h @@ -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_c_to_spirv(const struct clc_compile_args *args, + const struct clc_logger *logger, + struct clc_binary *out_spirv); bool -clc_compile(const struct clc_compile_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); bool -clc_link(const struct clc_linker_args *args, - const struct clc_logger *logger, - struct clc_object *out_spirv); +clc_parse_spirv(const struct clc_binary *in_spirv, + const struct clc_logger *logger, + struct clc_parsed_spirv *out_data); -void clc_free_object(struct clc_object *obj); +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); diff --git a/src/microsoft/clc/clc_compiler_test.cpp b/src/microsoft/clc/clc_compiler_test.cpp index 4d3182f17c3..2b3a8da5828 100644 --- a/src/microsoft/clc/clc_compiler_test.cpp +++ b/src/microsoft/clc/clc_compiler_test.cpp @@ -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) diff --git a/src/microsoft/clc/clc_helpers.cpp b/src/microsoft/clc/clc_helpers.cpp index 20bad26112b..18d056ea0d2 100644 --- a/src/microsoft/clc/clc_helpers.cpp +++ b/src/microsoft/clc/clc_helpers.cpp @@ -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); } diff --git a/src/microsoft/clc/clc_helpers.h b/src/microsoft/clc/clc_helpers.h index 653e99a8a27..c85caac0b04 100644 --- a/src/microsoft/clc/clc_helpers.h +++ b/src/microsoft/clc/clc_helpers.h @@ -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; \ diff --git a/src/microsoft/clc/clon12compiler.def b/src/microsoft/clc/clon12compiler.def index faa5da32f89..8a76cf87d83 100644 --- a/src/microsoft/clc/clon12compiler.def +++ b/src/microsoft/clc/clon12compiler.def @@ -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 diff --git a/src/microsoft/clc/compute_test.cpp b/src/microsoft/clc/compute_test.cpp index a0ed48814fd..34482e3ea49 100644 --- a/src/microsoft/clc/compute_test.cpp +++ b/src/microsoft/clc/compute_test.cpp @@ -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!"); } diff --git a/src/microsoft/clc/compute_test.h b/src/microsoft/clc/compute_test.h index 5fa666a6dfe..c915e74d554 100644 --- a/src/microsoft/clc/compute_test.h +++ b/src/microsoft/clc/compute_test.h @@ -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; };
