[Mesa-dev] [PATCH 2/2] radeonsi: don't load unused compute shader input SGPRs and VGPRs
Nicolai Hähnle
nhaehnle at gmail.com
Tue Apr 25 06:24:52 UTC 2017
On 24.04.2017 18:22, Marek Olšák wrote:
> From: Marek Olšák <marek.olsak at amd.com>
>
> Basically, don't load GRID_SIZE or BLOCK_SIZE if they are unused, determine
> whether to load BLOCK_ID for each component separately, and set the number
> of THREAD_ID VGPRs to load. Now we should get the maximum CS launch wave
> rate in most cases.
> ---
> src/gallium/drivers/radeonsi/si_compute.c | 71 ++++++++++++++---------
> src/gallium/drivers/radeonsi/si_shader.c | 37 ++++++++----
> src/gallium/drivers/radeonsi/si_shader.h | 11 ----
> src/gallium/drivers/radeonsi/si_shader_internal.h | 5 ++
> 4 files changed, 76 insertions(+), 48 deletions(-)
>
> diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c
> index 2b2efae..b3399d1 100644
> --- a/src/gallium/drivers/radeonsi/si_compute.c
> +++ b/src/gallium/drivers/radeonsi/si_compute.c
> @@ -41,20 +41,22 @@ struct si_compute {
>
> unsigned ir_type;
> unsigned local_size;
> unsigned private_size;
> unsigned input_size;
> struct si_shader shader;
>
> struct pipe_resource *global_buffers[MAX_GLOBAL_BUFFERS];
> unsigned use_code_object_v2 : 1;
> unsigned variable_group_size : 1;
> + unsigned uses_grid_size:1;
> + unsigned uses_block_size:1;
> };
>
> struct dispatch_packet {
> uint16_t header;
> uint16_t setup;
> uint16_t workgroup_size_x;
> uint16_t workgroup_size_y;
> uint16_t workgroup_size_z;
> uint16_t reserved0;
> uint32_t grid_size_x;
> @@ -114,37 +116,45 @@ static void si_create_compute_state_async(void *job, int thread_index)
> memset(&sel, 0, sizeof(sel));
>
> sel.screen = program->screen;
> tgsi_scan_shader(program->tokens, &sel.info);
> sel.tokens = program->tokens;
> sel.type = PIPE_SHADER_COMPUTE;
> sel.local_size = program->local_size;
>
> program->shader.selector = &sel;
> program->shader.is_monolithic = true;
> + program->uses_grid_size = sel.info.uses_grid_size;
> + program->uses_block_size = sel.info.uses_block_size;
>
> if (si_shader_create(program->screen, tm, &program->shader, debug)) {
> program->shader.compilation_failed = true;
> } else {
> bool scratch_enabled = shader->config.scratch_bytes_per_wave > 0;
> + unsigned user_sgprs = SI_NUM_RESOURCE_SGPRS +
> + (sel.info.uses_grid_size ? 3 : 0) +
> + (sel.info.uses_block_size ? 3 : 0);
>
> shader->config.rsrc1 =
> S_00B848_VGPRS((shader->config.num_vgprs - 1) / 4) |
> S_00B848_SGPRS((shader->config.num_sgprs - 1) / 8) |
> S_00B848_DX10_CLAMP(1) |
> S_00B848_FLOAT_MODE(shader->config.float_mode);
>
> shader->config.rsrc2 =
> - S_00B84C_USER_SGPR(SI_CS_NUM_USER_SGPR) |
> + S_00B84C_USER_SGPR(user_sgprs) |
> S_00B84C_SCRATCH_EN(scratch_enabled) |
> - S_00B84C_TGID_X_EN(1) | S_00B84C_TGID_Y_EN(1) |
> - S_00B84C_TGID_Z_EN(1) | S_00B84C_TIDIG_COMP_CNT(2) |
> + S_00B84C_TGID_X_EN(sel.info.uses_block_id[0]) |
> + S_00B84C_TGID_Y_EN(sel.info.uses_block_id[1]) |
> + S_00B84C_TGID_Z_EN(sel.info.uses_block_id[2]) |
> + S_00B84C_TIDIG_COMP_CNT(sel.info.uses_thread_id[2] ? 2 :
> + sel.info.uses_thread_id[1] ? 1 : 0) |
> S_00B84C_LDS_SIZE(shader->config.lds_size);
>
> program->variable_group_size =
> sel.info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0;
> }
>
> FREE(program->tokens);
> program->shader.selector = NULL;
> }
>
> @@ -644,50 +654,57 @@ static bool si_upload_compute_input(struct si_context *sctx,
> }
>
> r600_resource_reference(&input_buffer, NULL);
>
> return true;
> }
>
> static void si_setup_tgsi_grid(struct si_context *sctx,
> const struct pipe_grid_info *info)
> {
> + struct si_compute *program = sctx->cs_shader_state.program;
> struct radeon_winsys_cs *cs = sctx->b.gfx.cs;
> unsigned grid_size_reg = R_00B900_COMPUTE_USER_DATA_0 +
> - 4 * SI_SGPR_GRID_SIZE;
> + 4 * SI_NUM_RESOURCE_SGPRS;
> + unsigned block_size_reg = grid_size_reg +
> + /* 12 bytes = 3 dwords. */
> + 12 * program->uses_grid_size;
>
> if (info->indirect) {
> - uint64_t base_va = r600_resource(info->indirect)->gpu_address;
> - uint64_t va = base_va + info->indirect_offset;
> - int i;
> -
> - radeon_add_to_buffer_list(&sctx->b, &sctx->b.gfx,
> - (struct r600_resource *)info->indirect,
> - RADEON_USAGE_READ, RADEON_PRIO_DRAW_INDIRECT);
> -
> - for (i = 0; i < 3; ++i) {
> - radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
> - radeon_emit(cs, COPY_DATA_SRC_SEL(COPY_DATA_MEM) |
> - COPY_DATA_DST_SEL(COPY_DATA_REG));
> - radeon_emit(cs, (va + 4 * i));
> - radeon_emit(cs, (va + 4 * i) >> 32);
> - radeon_emit(cs, (grid_size_reg >> 2) + i);
> - radeon_emit(cs, 0);
> + if (program->uses_grid_size) {
> + uint64_t base_va = r600_resource(info->indirect)->gpu_address;
> + uint64_t va = base_va + info->indirect_offset;
> + int i;
> +
> + radeon_add_to_buffer_list(&sctx->b, &sctx->b.gfx,
> + (struct r600_resource *)info->indirect,
> + RADEON_USAGE_READ, RADEON_PRIO_DRAW_INDIRECT);
> +
> + for (i = 0; i < 3; ++i) {
> + radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
> + radeon_emit(cs, COPY_DATA_SRC_SEL(COPY_DATA_MEM) |
> + COPY_DATA_DST_SEL(COPY_DATA_REG));
> + radeon_emit(cs, (va + 4 * i));
Extra space (was in the original as well, but may as well fix it now).
> + radeon_emit(cs, (va + 4 * i) >> 32);
> + radeon_emit(cs, (grid_size_reg >> 2) + i);
> + radeon_emit(cs, 0);
> + }
> }
> } else {
> - struct si_compute *program = sctx->cs_shader_state.program;
> -
> - radeon_set_sh_reg_seq(cs, grid_size_reg, program->variable_group_size ? 6 : 3);
> - radeon_emit(cs, info->grid[0]);
> - radeon_emit(cs, info->grid[1]);
> - radeon_emit(cs, info->grid[2]);
> - if (program->variable_group_size) {
> + if (program->uses_grid_size) {
> + radeon_set_sh_reg_seq(cs, grid_size_reg, 3);
> + radeon_emit(cs, info->grid[0]);
> + radeon_emit(cs, info->grid[1]);
> + radeon_emit(cs, info->grid[2]);
> + }
> + if (program->variable_group_size && program->uses_block_size) {
> + radeon_set_sh_reg_seq(cs, block_size_reg, 3);
> radeon_emit(cs, info->block[0]);
> radeon_emit(cs, info->block[1]);
> radeon_emit(cs, info->block[2]);
This is a slight regression if both are used, though I guess variable
group size is rare enough that it's not a big deal.
Either way,
Reviewed-by: Nicolai Hähnle <nicolai.haehnle at amd.com>
By the way: Do we have a test case that exercises the corner case of a
compute shader that reads only the Y or only the Z component of
BLOCK_ID? That might be a good idea.
> }
> }
> }
>
> static void si_emit_dispatch_packets(struct si_context *sctx,
> const struct pipe_grid_info *info)
> {
> diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
> index 086b279..0fef036 100644
> --- a/src/gallium/drivers/radeonsi/si_shader.c
> +++ b/src/gallium/drivers/radeonsi/si_shader.c
> @@ -1583,52 +1583,63 @@ static void declare_system_value(struct si_shader_context *ctx,
> LLVMConstInt(ctx->i32, (offset + i) * 4, 0));
> value = lp_build_gather_values(gallivm, val, 4);
> break;
> }
>
> case TGSI_SEMANTIC_PRIMID:
> value = get_primitive_id(&ctx->bld_base, 0);
> break;
>
> case TGSI_SEMANTIC_GRID_SIZE:
> - value = LLVMGetParam(ctx->main_fn, SI_PARAM_GRID_SIZE);
> + value = LLVMGetParam(ctx->main_fn, ctx->param_grid_size);
> break;
>
> case TGSI_SEMANTIC_BLOCK_SIZE:
> {
> LLVMValueRef values[3];
> unsigned i;
> unsigned *properties = ctx->shader->selector->info.properties;
>
> if (properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] != 0) {
> unsigned sizes[3] = {
> properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH],
> properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT],
> properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH]
> };
>
> for (i = 0; i < 3; ++i)
> values[i] = LLVMConstInt(ctx->i32, sizes[i], 0);
>
> value = lp_build_gather_values(gallivm, values, 3);
> } else {
> - value = LLVMGetParam(ctx->main_fn, SI_PARAM_BLOCK_SIZE);
> + value = LLVMGetParam(ctx->main_fn, ctx->param_block_size);
> }
> break;
> }
>
> case TGSI_SEMANTIC_BLOCK_ID:
> - value = LLVMGetParam(ctx->main_fn, SI_PARAM_BLOCK_ID);
> + {
> + LLVMValueRef values[3];
> +
> + for (int i = 0; i < 3; i++) {
> + values[i] = ctx->i32_0;
> + if (ctx->param_block_id[i] >= 0) {
> + values[i] = LLVMGetParam(ctx->main_fn,
> + ctx->param_block_id[i]);
> + }
> + }
> + value = lp_build_gather_values(gallivm, values, 3);
> break;
> + }
>
> case TGSI_SEMANTIC_THREAD_ID:
> - value = LLVMGetParam(ctx->main_fn, SI_PARAM_THREAD_ID);
> + value = LLVMGetParam(ctx->main_fn, ctx->param_thread_id);
> break;
>
> case TGSI_SEMANTIC_HELPER_INVOCATION:
> if (HAVE_LLVM >= 0x0309) {
> value = lp_build_intrinsic(gallivm->builder,
> "llvm.amdgcn.ps.live",
> ctx->i1, NULL, 0,
> LP_FUNC_ATTR_READNONE);
> value = LLVMBuildNot(gallivm->builder, value, "");
> value = LLVMBuildSExt(gallivm->builder, value, ctx->i32, "");
> @@ -6176,27 +6187,33 @@ static void create_function(struct si_shader_context *ctx)
> PS_EPILOG_SAMPLEMASK_MIN_LOC + 1);
>
> for (i = 0; i < num_return_sgprs; i++)
> returns[i] = ctx->i32;
> for (; i < num_returns; i++)
> returns[i] = ctx->f32;
> break;
>
> case PIPE_SHADER_COMPUTE:
> declare_default_desc_pointers(ctx, params, &num_params);
> - params[SI_PARAM_GRID_SIZE] = v3i32;
> - params[SI_PARAM_BLOCK_SIZE] = v3i32;
> - params[SI_PARAM_BLOCK_ID] = v3i32;
> - last_sgpr = SI_PARAM_BLOCK_ID;
> + if (shader->selector->info.uses_grid_size)
> + params[ctx->param_grid_size = num_params++] = v3i32;
> + if (shader->selector->info.uses_block_size)
> + params[ctx->param_block_size = num_params++] = v3i32;
> +
> + for (i = 0; i < 3; i++) {
> + ctx->param_block_id[i] = -1;
> + if (shader->selector->info.uses_block_id[i])
> + params[ctx->param_block_id[i] = num_params++] = ctx->i32;
> + }
> + last_sgpr = num_params - 1;
>
> - params[SI_PARAM_THREAD_ID] = v3i32;
> - num_params = SI_PARAM_THREAD_ID + 1;
> + params[ctx->param_thread_id = num_params++] = v3i32;
> break;
> default:
> assert(0 && "unimplemented shader");
> return;
> }
>
> assert(num_params <= ARRAY_SIZE(params));
>
> si_create_function(ctx, "main", returns, num_returns, params,
> num_params, last_sgpr,
> diff --git a/src/gallium/drivers/radeonsi/si_shader.h b/src/gallium/drivers/radeonsi/si_shader.h
> index c9bd904..4248b80 100644
> --- a/src/gallium/drivers/radeonsi/si_shader.h
> +++ b/src/gallium/drivers/radeonsi/si_shader.h
> @@ -213,25 +213,20 @@ enum {
> GFX9_SGPR_GS_SHADER_BUFFERS_HI,
> GFX9_GS_NUM_USER_SGPR,
>
> /* GS limits */
> GFX6_GS_NUM_USER_SGPR = SI_NUM_RESOURCE_SGPRS,
> SI_GSCOPY_NUM_USER_SGPR = SI_SGPR_RW_BUFFERS_HI + 1,
>
> /* PS only */
> SI_SGPR_ALPHA_REF = SI_NUM_RESOURCE_SGPRS,
> SI_PS_NUM_USER_SGPR,
> -
> - /* CS only */
> - SI_SGPR_GRID_SIZE = SI_NUM_RESOURCE_SGPRS,
> - SI_SGPR_BLOCK_SIZE = SI_SGPR_GRID_SIZE + 3,
> - SI_CS_NUM_USER_SGPR = SI_SGPR_BLOCK_SIZE + 3
> };
>
> /* LLVM function parameter indices */
> enum {
> SI_NUM_RESOURCE_PARAMS = 5,
>
> /* PS only parameters */
> SI_PARAM_ALPHA_REF = SI_NUM_RESOURCE_PARAMS,
> SI_PARAM_PRIM_MASK,
> SI_PARAM_PERSP_SAMPLE,
> @@ -244,26 +239,20 @@ enum {
> SI_PARAM_LINE_STIPPLE_TEX,
> SI_PARAM_POS_X_FLOAT,
> SI_PARAM_POS_Y_FLOAT,
> SI_PARAM_POS_Z_FLOAT,
> SI_PARAM_POS_W_FLOAT,
> SI_PARAM_FRONT_FACE,
> SI_PARAM_ANCILLARY,
> SI_PARAM_SAMPLE_COVERAGE,
> SI_PARAM_POS_FIXED_PT,
>
> - /* CS only parameters */
> - SI_PARAM_GRID_SIZE = SI_NUM_RESOURCE_PARAMS,
> - SI_PARAM_BLOCK_SIZE,
> - SI_PARAM_BLOCK_ID,
> - SI_PARAM_THREAD_ID,
> -
> SI_NUM_PARAMS = SI_PARAM_POS_FIXED_PT + 9, /* +8 for COLOR[0..1] */
> };
>
> /* Fields of driver-defined VS state SGPR. */
> /* Clamp vertex color output (only used in VS as VS). */
> #define S_VS_STATE_CLAMP_VERTEX_COLOR(x) (((unsigned)(x) & 0x1) << 0)
> #define C_VS_STATE_CLAMP_VERTEX_COLOR 0xFFFFFFFE
> #define S_VS_STATE_INDEXED(x) (((unsigned)(x) & 0x1) << 1)
> #define C_VS_STATE_INDEXED 0xFFFFFFFD
> #define S_VS_STATE_LS_OUT_PATCH_SIZE(x) (((unsigned)(x) & 0x1FFF) << 8)
> diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h
> index 954b83d..cad2db3 100644
> --- a/src/gallium/drivers/radeonsi/si_shader_internal.h
> +++ b/src/gallium/drivers/radeonsi/si_shader_internal.h
> @@ -188,20 +188,25 @@ struct si_shader_context {
> int param_gs_vtx1_offset; /* in dwords (GFX6) */
> int param_gs_prim_id;
> int param_gs_vtx2_offset; /* in dwords (GFX6) */
> int param_gs_vtx3_offset; /* in dwords (GFX6) */
> int param_gs_vtx4_offset; /* in dwords (GFX6) */
> int param_gs_vtx5_offset; /* in dwords (GFX6) */
> int param_gs_instance_id;
> int param_gs_vtx01_offset; /* in dwords (GFX9) */
> int param_gs_vtx23_offset; /* in dwords (GFX9) */
> int param_gs_vtx45_offset; /* in dwords (GFX9) */
> + /* CS */
> + int param_grid_size;
> + int param_block_size;
> + int param_block_id[3];
> + int param_thread_id;
>
> LLVMTargetMachineRef tm;
>
> unsigned range_md_kind;
> unsigned fpmath_md_kind;
> LLVMValueRef fpmath_md_2p5_ulp;
>
> /* Preloaded descriptors. */
> LLVMValueRef esgs_ring;
> LLVMValueRef gsvs_ring[4];
>
--
Lerne, wie die Welt wirklich ist,
Aber vergiss niemals, wie sie sein sollte.
More information about the mesa-dev
mailing list