fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-27 12:40:09 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

microsoft/clc: Stop heap-allocating tiny fixed-size transparent structs

Browse source 
The caller can allocate these however they want. They don't need
independent allocations. Removes some unnecessary failure handling.

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:31:05 -07:00 • committed by Marge Bot
parent 27467700e9
commit c3666eec7e
3 changed files with 66 additions and 85 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -575,67 +575,44 @@ struct clc_libclc *
return ctx;
}
struct clc_object *
bool
clc_compile(const struct clc_compile_args *args,
const struct clc_logger *logger)
const struct clc_logger *logger,
struct clc_object *out_spirv)
{
struct clc_object *obj;
int ret;
obj = calloc(1, sizeof(*obj));
if (!obj) {
clc_error(logger, "D3D12: failed to allocate a clc_object");
return NULL;
}
ret = clc_to_spirv(args, &obj->spvbin, logger);
if (ret < 0) {
free(obj);
return NULL;
}
if (clc_to_spirv(args, &out_spirv->spvbin, logger) < 0)
return false;
if (debug_get_option_debug_clc() & CLC_DEBUG_DUMP_SPIRV)
clc_dump_spirv(&obj->spvbin, stdout);
clc_dump_spirv(&out_spirv->spvbin, stdout);
return obj;
return true;
}
struct clc_object *
bool
clc_link(const struct clc_linker_args *args,
const struct clc_logger *logger)
const struct clc_logger *logger,
struct clc_object *out_spirv)
{
struct clc_object *out_obj;
int ret;
out_obj = malloc(sizeof(*out_obj));
if (!out_obj) {
clc_error(logger, "failed to allocate a clc_object");
return NULL;
}
ret = clc_link_spirv_binaries(args, &out_obj->spvbin, logger);
if (ret < 0) {
free(out_obj);
return NULL;
}
if (clc_link_spirv_binaries(args, &out_spirv->spvbin, logger) < 0)
return false;
if (debug_get_option_debug_clc() & CLC_DEBUG_DUMP_SPIRV)
clc_dump_spirv(&out_obj->spvbin, stdout);
clc_dump_spirv(&out_spirv->spvbin, stdout);
out_obj->kernels = clc_spirv_get_kernels_info(&out_obj->spvbin,
&out_obj->num_kernels);
out_spirv->kernels = clc_spirv_get_kernels_info(&out_spirv->spvbin,
&out_spirv->num_kernels);
if (debug_get_option_debug_clc() & CLC_DEBUG_VERBOSE)
clc_print_kernels_info(out_obj);
clc_print_kernels_info(out_spirv);
return out_obj;
return true;
}
void clc_free_object(struct clc_object *obj)
{
clc_free_kernels_info(obj->kernels, obj->num_kernels);
clc_free_spirv_binary(&obj->spvbin);
free(obj);
}
static nir_variable *
@ -1011,32 +988,26 @@ scale_fdiv(nir_shader *nir)
return progress;
}
struct clc_dxil_object *
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)
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil)
{
struct clc_dxil_object *dxil;
struct nir_shader *nir;
dxil = calloc(1, sizeof(*dxil));
if (!dxil) {
clc_error(logger, "failed to allocate the dxil object");
return NULL;
}
for (unsigned i = 0; i < obj->num_kernels; i++) {
if (!strcmp(obj->kernels[i].name, entrypoint)) {
dxil->kernel = &obj->kernels[i];
out_dxil->kernel = &obj->kernels[i];
break;
}
}
if (!dxil->kernel) {
if (!out_dxil->kernel) {
clc_error(logger, "no '%s' kernel found", entrypoint);
goto err_free_dxil;
return false;
}
const struct spirv_to_nir_options spirv_options = {
@ -1088,9 +1059,9 @@ clc_to_dxil(struct clc_libclc *lib,
NIR_PASS_V(nir, nir_lower_goto_ifs);
NIR_PASS_V(nir, nir_opt_dead_cf);
struct clc_dxil_metadata *metadata = &dxil->metadata;
struct clc_dxil_metadata *metadata = &out_dxil->metadata;
metadata->args = calloc(dxil->kernel->num_args,
metadata->args = calloc(out_dxil->kernel->num_args,
sizeof(*metadata->args));
if (!metadata->args) {
clc_error(logger, "failed to allocate arg positions");
@ -1220,8 +1191,8 @@ clc_to_dxil(struct clc_libclc *lib,
metadata->args[i].size = size;
metadata->kernel_inputs_buf_size = MAX2(metadata->kernel_inputs_buf_size,
var->data.driver_location + size);
if ((dxil->kernel->args[i].address_qualifier == CLC_KERNEL_ARG_ADDRESS_GLOBAL ||
dxil->kernel->args[i].address_qualifier == CLC_KERNEL_ARG_ADDRESS_CONSTANT) &&
if ((out_dxil->kernel->args[i].address_qualifier == CLC_KERNEL_ARG_ADDRESS_GLOBAL ||
out_dxil->kernel->args[i].address_qualifier == CLC_KERNEL_ARG_ADDRESS_CONSTANT) &&
// Ignore images during this pass - global memory buffers need to have contiguous bindings
!glsl_type_is_image(var->type)) {
metadata->args[i].globconstptr.buf_id = uav_id++;
@ -1303,7 +1274,7 @@ clc_to_dxil(struct clc_libclc *lib,
glsl_get_cl_type_size_align);
NIR_PASS_V(nir, dxil_nir_lower_ubo_to_temp);
NIR_PASS_V(nir, clc_lower_constant_to_ssbo, dxil->kernel, &uav_id);
NIR_PASS_V(nir, clc_lower_constant_to_ssbo, out_dxil->kernel, &uav_id);
NIR_PASS_V(nir, clc_lower_global_to_ssbo);
bool has_printf = false;
@ -1337,9 +1308,9 @@ clc_to_dxil(struct clc_libclc *lib,
unsigned cbv_id = 0;
nir_variable *inputs_var =
add_kernel_inputs_var(dxil, nir, &cbv_id);
add_kernel_inputs_var(out_dxil, nir, &cbv_id);
nir_variable *work_properties_var =
add_work_properties_var(dxil, nir, &cbv_id);
add_work_properties_var(out_dxil, nir, &cbv_id);
memcpy(metadata->local_size, nir->info.workgroup_size,
sizeof(metadata->local_size));
@ -1398,8 +1369,8 @@ clc_to_dxil(struct clc_libclc *lib,
.num_kernel_globals = num_global_inputs,
};
for (unsigned i = 0; i < dxil->kernel->num_args; i++) {
if (dxil->kernel->args[i].address_qualifier != CLC_KERNEL_ARG_ADDRESS_LOCAL)
for (unsigned i = 0; i < out_dxil->kernel->num_args; i++) {
if (out_dxil->kernel->args[i].address_qualifier != CLC_KERNEL_ARG_ADDRESS_LOCAL)
continue;
/* If we don't have the runtime conf yet, we just create a dummy variable.
@ -1469,13 +1440,13 @@ clc_to_dxil(struct clc_libclc *lib,
ralloc_free(nir);
glsl_type_singleton_decref();
blob_finish_get_buffer(&tmp, &dxil->binary.data,
&dxil->binary.size);
return dxil;
blob_finish_get_buffer(&tmp, &out_dxil->binary.data,
&out_dxil->binary.size);
return true;
err_free_dxil:
clc_free_dxil_object(dxil);
return NULL;
clc_free_dxil_object(out_dxil);
return false;
}
void clc_free_dxil_object(struct clc_dxil_object *dxil)
@ -1490,7 +1461,6 @@ void clc_free_dxil_object(struct clc_dxil_object *dxil)
free(dxil->metadata.printf.infos);
free(dxil->binary.data);
free(dxil);
}
uint64_t clc_compiler_get_version()

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

@ -203,13 +203,15 @@ struct clc_libclc *clc_libclc_deserialize(void *serialized, size_t size);
struct clc_object *
bool
clc_compile(const struct clc_compile_args *args,
const struct clc_logger *logger);
const struct clc_logger *logger,
struct clc_object *out_spirv);
struct clc_object *
bool
clc_link(const struct clc_linker_args *args,
const struct clc_logger *logger);
const struct clc_logger *logger,
struct clc_object *out_spirv);
void clc_free_object(struct clc_object *obj);
@ -234,12 +236,13 @@ struct clc_runtime_kernel_conf {
unsigned support_workgroup_id_offsets;
};
struct clc_dxil_object *
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);
const struct clc_logger *logger,
struct clc_dxil_object *out_dxil);
void clc_free_dxil_object(struct clc_dxil_object *dxil);

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

@ -803,12 +803,16 @@ ComputeTest::compile(const std::vector<const char *> &sources,
for (unsigned i = 0; i < sources.size(); i++) {
args.source.value = sources[i];
auto obj = clc_compile(&args, &logger);
if (!obj)
clc_object spirv{};
if (!clc_compile(&args, &logger, &spirv))
throw runtime_error("failed to compile object!");
Shader shader;
shader.obj = std::shared_ptr<struct clc_object>(obj, clc_free_object);
shader.obj = std::shared_ptr<clc_object>(new clc_object(spirv), [](clc_object *spirv)
{
clc_free_object(spirv);
delete spirv;
});
shaders.push_back(shader);
}
@ -830,13 +834,16 @@ 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;
struct clc_object *obj = clc_link(&link_args,
&logger);
if (!obj)
clc_object spirv{};
if (!clc_link(&link_args, &logger, &spirv))
throw runtime_error("failed to link objects!");
ComputeTest::Shader shader;
shader.obj = std::shared_ptr<struct clc_object>(obj, clc_free_object);
shader.obj = std::shared_ptr<clc_object>(new clc_object(spirv), [](clc_object *spirv)
{
clc_free_object(spirv);
delete spirv;
});
if (!link_args.create_library)
configure(shader, NULL);
@ -847,13 +854,14 @@ void
ComputeTest::configure(Shader &shader,
const struct clc_runtime_kernel_conf *conf)
{
struct clc_dxil_object *dxil;
dxil = clc_to_dxil(compiler_ctx, shader.obj.get(), "main_test", conf, &logger);
if (!dxil)
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()))
throw runtime_error("failed to compile kernel!");
shader.dxil = std::shared_ptr<struct clc_dxil_object>(dxil, clc_free_dxil_object);
}
void