[Mesa-dev] [PATCH] radeonsi: store group_size_variable in struct si_compute

[Nicolai Hähnle]

From: Nicolai Hähnle <nicolai.haehnle at amd.com>

For compute shaders, we free the selector after the shader has been
compiled, so we need to save this bit somewhere else.  Also, make sure that
this type of bug cannot re-appear, by NULL-ing the selector pointer after
we're done with it.

This bug has been there since the feature was added, but was only exposed
in piglit arb_compute_variable_group_size-local-size by commit
9bfee7047b70cb0aa026ca9536465762f96cb2b1 (which is totally unrelated).

Cc: 13.0 <mesa-stable at lists.freedesktop.org>
---
 src/gallium/drivers/radeonsi/si_compute.c | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c
index f1887bb..69d57b9 100644
--- a/src/gallium/drivers/radeonsi/si_compute.c
+++ b/src/gallium/drivers/radeonsi/si_compute.c
@@ -35,21 +35,22 @@
 #define MAX_GLOBAL_BUFFERS 20
 
 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];
-	bool use_code_object_v2;
+	unsigned use_code_object_v2 : 1;
+	unsigned variable_group_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;
@@ -140,21 +141,25 @@ static void *si_create_compute_state(
 			   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_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_LDS_SIZE(shader->config.lds_size);
 
+		program->variable_group_size =
+			sel.info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0;
+
 		FREE(sel.tokens);
+		program->shader.selector = NULL;
 	} else {
 		const struct pipe_llvm_program_header *header;
 		const char *code;
 		header = cso->prog;
 		code = cso->prog + sizeof(struct pipe_llvm_program_header);
 
 		radeon_elf_read(code, header->num_bytes, &program->shader.binary);
 		if (program->use_code_object_v2) {
 			const amd_kernel_code_t *code_object =
 				si_compute_get_code_object(program, 0);
@@ -600,28 +605,26 @@ static void si_setup_tgsi_grid(struct si_context *sctx,
 			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);
 		}
 	} else {
 		struct si_compute *program = sctx->cs_shader_state.program;
-		bool variable_group_size =
-			program->shader.selector->info.properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0;
 
-		radeon_set_sh_reg_seq(cs, grid_size_reg, variable_group_size ? 6 : 3);
+		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 (variable_group_size) {
+		if (program->variable_group_size) {
 			radeon_emit(cs, info->block[0]);
 			radeon_emit(cs, info->block[1]);
 			radeon_emit(cs, info->block[2]);
 		}
 	}
 }
 
 static void si_emit_dispatch_packets(struct si_context *sctx,
                                      const struct pipe_grid_info *info)
 {
-- 
2.7.4