radeonsi: use shader_info::cs::local_size_variable to clean up some code

[mesa.git] / src / gallium / drivers / radeonsi / si_shader_llvm.c

diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c
index 51072c1efc97748d49ef0ee8986c46074d101c7a..0602593ba6e4b435dcc3d33ace4e1817079da731 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shader_llvm.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c
@@ -70,11 +70,11 @@ static void si_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
 bool si_compile_llvm(struct si_screen *sscreen, struct si_shader_binary *binary,
                      struct ac_shader_config *conf, struct ac_llvm_compiler *compiler,
                      struct ac_llvm_context *ac, struct pipe_debug_callback *debug,
-                     enum pipe_shader_type shader_type, const char *name, bool less_optimized)
+                     gl_shader_stage stage, const char *name, bool less_optimized)
 {
    unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
 
-   if (si_can_dump_shader(sscreen, shader_type)) {
+   if (si_can_dump_shader(sscreen, stage)) {
       fprintf(stderr, "radeonsi: Compiling shader %d\n", count);
 
       if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) {
@@ -114,7 +114,7 @@ bool si_compile_llvm(struct si_screen *sscreen, struct si_shader_binary *binary,
    struct ac_rtld_binary rtld;
    if (!ac_rtld_open(&rtld, (struct ac_rtld_open_info){
                                .info = &sscreen->info,
-                               .shader_type = tgsi_processor_to_shader_stage(shader_type),
+                               .shader_type = stage,
                                .wave_size = ac->wave_size,
                                .num_parts = 1,
                                .elf_ptrs = &binary->elf_buffer,
@@ -189,15 +189,13 @@ void si_llvm_create_func(struct si_shader_context *ctx, const char *name, LLVMTy
                                            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)
 {
    /* Dump LLVM IR before any optimization passes */
-   if (ctx->screen->debug_flags & DBG(PREOPT_IR) && si_can_dump_shader(ctx->screen, ctx->type))
+   if (ctx->screen->debug_flags & DBG(PREOPT_IR) && si_can_dump_shader(ctx->screen, ctx->stage))
       LLVMDumpModule(ctx->ac.module);
 
    /* Run the pass */
@@ -372,15 +370,12 @@ LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
    LLVMValueRef values[3];
    LLVMValueRef result;
    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]};
+   if (!ctx->shader->selector->info.base.cs.local_size_variable) {
+      uint16_t *local_size = ctx->shader->selector->info.base.cs.local_size;
 
       for (i = 0; i < 3; ++i)
-         values[i] = LLVMConstInt(ctx->ac.i32, sizes[i], 0);
+         values[i] = LLVMConstInt(ctx->ac.i32, local_size[i], 0);
 
       result = ac_build_gather_values(&ctx->ac, values, 3);
    } else {
@@ -393,7 +388,7 @@ LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
 void si_llvm_declare_compute_memory(struct si_shader_context *ctx)
 {
    struct si_shader_selector *sel = ctx->shader->selector;
-   unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
+   unsigned lds_size = sel->info.base.cs.shared_size;
 
    LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
    LLVMValueRef var;
@@ -453,6 +448,9 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
          ctx->abi.user_data = ac_build_expand_to_vec4(&ctx->ac, ctx->abi.user_data,
                                                       nir->info.cs.user_data_components_amd);
       }
+
+      if (ctx->shader->selector->info.base.cs.shared_size)
+         si_llvm_declare_compute_memory(ctx);
    }
 
    ctx->abi.inputs = &ctx->inputs[0];
@@ -461,11 +459,6 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
    ctx->abi.convert_undef_to_zero = true;
    ctx->abi.clamp_div_by_zero = ctx->screen->options.clamp_div_by_zero;
 
-   if (ctx->shader->selector->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE]) {
-      assert(gl_shader_stage_is_compute(nir->info.stage));
-      si_llvm_declare_compute_memory(ctx);
-   }
-
    const struct si_shader_info *info = &ctx->shader->selector->info;
    for (unsigned i = 0; i < info->num_outputs; i++) {
       for (unsigned j = 0; j < 4; j++)