ac: use amdgpu-flat-work-group-size

Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>

diff --git a/src/amd/common/ac_llvm_util.c b/src/amd/common/ac_llvm_util.c
index 5b701603ebbe254bda0242bcd5a740a943c0501e..c8a8bf146fe56ee3ba080556e418ed6eb4a516f2 100644 (file)
--- a/src/amd/common/ac_llvm_util.c
+++ b/src/amd/common/ac_llvm_util.c
@@ -269,6 +269,16 @@ ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
        LLVMAddTargetDependentFunctionAttr(F, name, str);
 }
 
+void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size)
+{
+       if (!size)
+               return;
+
+       char str[32];
+       snprintf(str, sizeof(str), "%u,%u", size, size);
+       LLVMAddTargetDependentFunctionAttr(F, "amdgpu-flat-work-group-size", str);
+}
+
 unsigned
 ac_count_scratch_private_memory(LLVMValueRef function)
 {

diff --git a/src/amd/common/ac_llvm_util.h b/src/amd/common/ac_llvm_util.h
index ca00540da8063cc9e4586beb5406fc14b3f3f924..18102be5207d6459c8ca8a1d74fb239e1b7af30c 100644 (file)
--- a/src/amd/common/ac_llvm_util.h
+++ b/src/amd/common/ac_llvm_util.h
@@ -109,6 +109,7 @@ LLVMBuilderRef ac_create_builder(LLVMContextRef ctx,
 void
 ac_llvm_add_target_dep_function_attr(LLVMValueRef F,
                                     const char *name, unsigned value);
+void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size);
 
 static inline unsigned
 ac_get_load_intr_attribs(bool can_speculate)

diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
index dca4bebcdd16cc87e854705617c691744ad78633..98a04e27b25ffb327d989045154bf5be64cf5254 100644 (file)
--- a/src/amd/vulkan/radv_nir_to_llvm.c
+++ b/src/amd/vulkan/radv_nir_to_llvm.c
@@ -518,11 +518,8 @@ create_llvm_function(LLVMContextRef ctx, LLVMModuleRef module,
                                                     options->address32_hi);
        }
 
-       if (max_workgroup_size) {
-               ac_llvm_add_target_dep_function_attr(main_function,
-                                                    "amdgpu-max-work-group-size",
-                                                    max_workgroup_size);
-       }
+       ac_llvm_set_workgroup_size(main_function, max_workgroup_size);
+
        if (options->unsafe_math) {
                /* These were copied from some LLVM test. */
                LLVMAddTargetDependentFunctionAttr(main_function,

diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 5bd65e0f65c3fe37fd68f9358eb7e1e603e7f3c6..e044e180778a683cfed81f531d6d32dab4bf1f07 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -4307,11 +4307,8 @@ void si_create_function(struct si_shader_context *ctx,
                                                     ctx->screen->info.address32_hi);
        }
 
-       if (max_workgroup_size) {
-               ac_llvm_add_target_dep_function_attr(ctx->main_fn,
-                                                    "amdgpu-max-work-group-size",
-                                                    max_workgroup_size);
-       }
+       ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
+
        LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
                                           "no-signed-zeros-fp-math",
                                           "true");