[Mesa-dev] [PATCH 3/9] radeonsi: use is_merged shader in si_prolog_get_rw_buffers
Marek Olšák
maraeo at gmail.com
Tue Aug 21 03:23:33 UTC 2018
From: Marek Olšák <marek.olsak at amd.com>
needed to change the input type to si_shader_context
---
src/gallium/drivers/radeonsi/si_shader.c | 32 +++++++++++-------------
1 file changed, 14 insertions(+), 18 deletions(-)
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 551671f4021..354c05e3d9d 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -94,29 +94,29 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx,
static bool llvm_type_is_64bit(struct si_shader_context *ctx,
LLVMTypeRef type)
{
if (type == ctx->ac.i64 || type == ctx->ac.f64)
return true;
return false;
}
-static bool is_merged_shader(struct si_shader *shader)
+static bool is_merged_shader(struct si_shader_context *ctx)
{
- if (shader->selector->screen->info.chip_class <= VI)
+ if (ctx->screen->info.chip_class <= VI)
return false;
- return shader->key.as_ls ||
- shader->key.as_es ||
- shader->selector->type == PIPE_SHADER_TESS_CTRL ||
- shader->selector->type == PIPE_SHADER_GEOMETRY;
+ return ctx->shader->key.as_ls ||
+ ctx->shader->key.as_es ||
+ ctx->type == PIPE_SHADER_TESS_CTRL ||
+ ctx->type == PIPE_SHADER_GEOMETRY;
}
static void si_init_function_info(struct si_function_info *fninfo)
{
fninfo->num_params = 0;
fninfo->num_sgpr_params = 0;
}
static unsigned add_arg_assign(struct si_function_info *fninfo,
enum si_arg_regfile regfile, LLVMTypeRef type,
@@ -6573,21 +6573,21 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
assert(ac_is_sgpr_param(param) == (gprs < num_sgprs));
assert(gprs + size <= num_sgprs + num_vgprs &&
(gprs >= num_sgprs || gprs + size <= num_sgprs));
gprs += size;
}
si_create_function(ctx, "wrapper", NULL, 0, &fninfo,
si_get_max_workgroup_size(ctx->shader));
- if (is_merged_shader(ctx->shader))
+ if (is_merged_shader(ctx))
ac_init_exec_full_mask(&ctx->ac);
/* Record the arguments of the function as if they were an output of
* a previous part.
*/
num_out = 0;
num_out_sgpr = 0;
for (unsigned i = 0; i < fninfo.num_params; ++i) {
LLVMValueRef param = LLVMGetParam(ctx->main_fn, i);
@@ -6631,21 +6631,21 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
/* Now chain the parts. */
for (unsigned part = 0; part < num_parts; ++part) {
LLVMValueRef in[48];
LLVMValueRef ret;
LLVMTypeRef ret_type;
unsigned out_idx = 0;
unsigned num_params = LLVMCountParams(parts[part]);
/* Merged shaders are executed conditionally depending
* on the number of enabled threads passed in the input SGPRs. */
- if (is_merged_shader(ctx->shader) && part == 0) {
+ if (is_merged_shader(ctx) && part == 0) {
LLVMValueRef ena, count = initial[3];
count = LLVMBuildAnd(builder, count,
LLVMConstInt(ctx->i32, 0x7f, 0), "");
ena = LLVMBuildICmp(builder, LLVMIntULT,
ac_get_thread_id(&ctx->ac), count, "");
lp_build_if(&if_state, &ctx->gallivm, ena);
}
/* Derive arguments for the next part from outputs of the
@@ -6693,21 +6693,21 @@ static void si_build_wrapper_function(struct si_shader_context *ctx,
arg = LLVMBuildBitCast(builder, arg, param_type, "");
}
}
in[param_idx] = arg;
out_idx += param_size;
}
ret = LLVMBuildCall(builder, parts[part], in, num_params, "");
- if (is_merged_shader(ctx->shader) &&
+ if (is_merged_shader(ctx) &&
part + 1 == next_shader_first_part) {
lp_build_endif(&if_state);
/* The second half of the merged shader should use
* the inputs from the toplevel (wrapper) function,
* not the return value from the last call.
*
* That's because the last call was executed condi-
* tionally, so we can't consume it in the main
* block.
@@ -7027,21 +7027,21 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
/* Just terminate the process, because dependent
* shaders can hang due to bad input data, but use
* the env var to allow shader-db to work.
*/
if (!debug_get_bool_option("SI_PASS_BAD_SHADERS", false))
abort();
}
}
/* Add the scratch offset to input SGPRs. */
- if (shader->config.scratch_bytes_per_wave && !is_merged_shader(shader))
+ if (shader->config.scratch_bytes_per_wave && !is_merged_shader(&ctx))
shader->info.num_input_sgprs += 1; /* scratch byte offset */
/* Calculate the number of fragment input VGPRs. */
if (ctx.type == PIPE_SHADER_FRAGMENT) {
shader->info.num_input_vgprs = 0;
shader->info.face_vgpr_index = -1;
shader->info.ancillary_vgpr_index = -1;
if (G_0286CC_PERSP_SAMPLE_ENA(shader->config.spi_ps_input_addr))
shader->info.num_input_vgprs += 2;
@@ -7173,36 +7173,32 @@ si_get_shader_part(struct si_screen *sscreen,
out:
si_llvm_dispose(&ctx);
mtx_unlock(&sscreen->shader_parts_mutex);
return result;
}
static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
{
LLVMValueRef ptr[2], list;
- bool is_merged_shader =
- ctx->screen->info.chip_class >= GFX9 &&
- (ctx->type == PIPE_SHADER_TESS_CTRL ||
- ctx->type == PIPE_SHADER_GEOMETRY ||
- ctx->shader->key.as_ls || ctx->shader->key.as_es);
+ bool merged_shader = is_merged_shader(ctx);
if (HAVE_32BIT_POINTERS) {
- ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
+ ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0],
ac_array_in_const32_addr_space(ctx->v4i32), "");
return list;
}
/* Get the pointer to rw buffers. */
- ptr[0] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
- ptr[1] = LLVMGetParam(ctx->main_fn, (is_merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS + 1);
+ ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
+ ptr[1] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS + 1);
list = ac_build_gather_values(&ctx->ac, ptr, 2);
list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, "");
list = LLVMBuildIntToPtr(ctx->ac.builder, list,
ac_array_in_const_addr_space(ctx->v4i32), "");
return list;
}
/**
* Build the vertex shader prolog function.
*
--
2.17.1
More information about the mesa-dev
mailing list