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

radeonsi: fold si_create_function into si_llvm_create_func

Browse Source 
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
This commit is contained in:
Marek Olšák
2020-01-10 20:22:47 -05:00
parent 42112010a3
commit 37916a66b1
4 changed files with 30 additions and 43 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/gallium/drivers/radeonsi/si_compute_prim_discard.c
View File

@@ -364,7 +364,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
/* Create the compute shader function. */
unsigned old_type = ctx->type;
ctx->type = PIPE_SHADER_COMPUTE;
si_create_function(ctx, "prim_discard_cs", NULL, 0, THREADGROUP_SIZE);
si_llvm_create_func(ctx, "prim_discard_cs", NULL, 0, THREADGROUP_SIZE);
ctx->type = old_type;
if (VERTEX_COUNTER_GDS_MODE == 1) {

42
src/gallium/drivers/radeonsi/si_shader.c
View File

@@ -3225,27 +3225,6 @@ static void si_llvm_emit_barrier(struct si_shader_context *ctx)
ac_build_s_barrier(&ctx->ac);
}
void si_create_function(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *returns, unsigned num_returns,
unsigned max_workgroup_size)
{
si_llvm_create_func(ctx, name, returns, num_returns);
ctx->return_value = LLVMGetUndef(ctx->return_type);
if (ctx->screen->info.address32_hi) {
ac_llvm_add_target_dep_function_attr(ctx->main_fn,
"amdgpu-32bit-address-high-bits",
ctx->screen->info.address32_hi);
}
LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
"no-signed-zeros-fp-math",
"true");
ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
}
static void declare_streamout_params(struct si_shader_context *ctx,
struct pipe_stream_output_info *so)
{
@@ -3811,8 +3790,8 @@ static void create_function(struct si_shader_context *ctx)
return;
}
si_create_function(ctx, "main", returns, num_returns,
si_get_max_workgroup_size(shader));
si_llvm_create_func(ctx, "main", returns, num_returns,
si_get_max_workgroup_size(shader));
/* Reserve register locations for VGPR inputs the PS prolog may need. */
if (ctx->type == PIPE_SHADER_FRAGMENT && !ctx->shader->is_monolithic) {
@@ -5350,8 +5329,7 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx,
}
/* Create the function. */
si_create_function(ctx, "gs_prolog", returns, num_sgprs + num_vgprs,
0);
si_llvm_create_func(ctx, "gs_prolog", returns, num_sgprs + num_vgprs, 0);
func = ctx->main_fn;
/* Set the full EXEC mask for the prolog, because we are only fiddling
@@ -5535,8 +5513,8 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
unreachable("unexpected type");
}
si_create_function(ctx, "wrapper", returns, num_returns,
si_get_max_workgroup_size(ctx->shader));
si_llvm_create_func(ctx, "wrapper", returns, num_returns,
si_get_max_workgroup_size(ctx->shader));
if (is_merged_shader(ctx))
ac_init_exec_full_mask(&ctx->ac);
@@ -6209,7 +6187,7 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
returns[num_returns++] = ctx->f32;
/* Create the function. */
si_create_function(ctx, "vs_prolog", returns, num_returns, 0);
si_llvm_create_func(ctx, "vs_prolog", returns, num_returns, 0);
func = ctx->main_fn;
for (i = 0; i < num_input_vgprs; i++) {
@@ -6440,8 +6418,8 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &tess_factors[i]);
/* Create the function. */
si_create_function(ctx, "tcs_epilog", NULL, 0,
ctx->screen->info.chip_class >= GFX7 ? 128 : 0);
si_llvm_create_func(ctx, "tcs_epilog", NULL, 0,
ctx->screen->info.chip_class >= GFX7 ? 128 : 0);
ac_declare_lds_as_pointer(&ctx->ac);
LLVMValueRef invoc0_tess_factors[6];
@@ -6585,7 +6563,7 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
return_types[num_returns++] = ctx->f32;
/* Create the function. */
si_create_function(ctx, "ps_prolog", return_types, num_returns, 0);
si_llvm_create_func(ctx, "ps_prolog", return_types, num_returns, 0);
func = ctx->main_fn;
/* Copy inputs to outputs. This should be no-op, as the registers match,
@@ -6861,7 +6839,7 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx,
ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_FLOAT, NULL);
/* Create the function. */
si_create_function(ctx, "ps_epilog", NULL, 0, 0);
si_llvm_create_func(ctx, "ps_epilog", NULL, 0, 0);
/* Disable elimination of unused inputs. */
ac_llvm_add_target_dep_function_attr(ctx->main_fn,
"InitialPSInputAddr", 0xffffff);

10
src/gallium/drivers/radeonsi/si_shader_internal.h
View File

@@ -199,10 +199,6 @@ si_shader_context_from_abi(struct ac_shader_abi *abi)
return container_of(abi, ctx, abi);
}
void si_create_function(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *returns, unsigned num_returns,
unsigned max_workgroup_size);
unsigned si_llvm_compile(LLVMModuleRef M, struct si_shader_binary *binary,
struct ac_llvm_compiler *compiler,
struct pipe_debug_callback *debug,
@@ -220,9 +216,9 @@ void si_llvm_context_init(struct si_shader_context *ctx,
void si_llvm_context_set_ir(struct si_shader_context *ctx,
struct si_shader *shader);
void si_llvm_create_func(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems);
void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems,
unsigned max_workgroup_size);
void si_llvm_dispose(struct si_shader_context *ctx);

19
src/gallium/drivers/radeonsi/si_shader_llvm.c
View File

@@ -167,9 +167,9 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
ctx->num_images = util_last_bit(info->images_declared);
}
void si_llvm_create_func(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems)
void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems,
unsigned max_workgroup_size)
{
LLVMTypeRef ret_type;
enum ac_llvm_calling_convention call_conv;
@@ -217,6 +217,19 @@ void si_llvm_create_func(struct si_shader_context *ctx,
ctx->return_type = ret_type;
ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
ret_type, ctx->ac.module);
ctx->return_value = LLVMGetUndef(ctx->return_type);
if (ctx->screen->info.address32_hi) {
ac_llvm_add_target_dep_function_attr(ctx->main_fn,
"amdgpu-32bit-address-high-bits",
ctx->screen->info.address32_hi);
}
LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
"no-signed-zeros-fp-math",
"true");
ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
}
void si_llvm_optimize_module(struct si_shader_context *ctx)