diff options
| author | Marek Olšák <[email protected]> | 2019-05-31 15:38:39 -0400 |
|---|---|---|
| committer | Marek Olšák <[email protected]> | 2019-06-03 14:32:47 -0400 |
| commit | 486bc1e17ecc975f98fb495bd2f8ae580eebbf6e (patch) | |
| tree | d2d58d3caf48f05aa55a82145923980c0f8e6edb /src/amd/common | |
| parent | 4b11ed443b85e4fcddc5d0ef60dec096ecdb951e (diff) | |
ac: use amdgpu-flat-work-group-size
Reviewed-by: Bas Nieuwenhuizen <[email protected]>
Diffstat (limited to 'src/amd/common')
| -rw-r--r-- | src/amd/common/ac_llvm_util.c | 10 | ||||
| -rw-r--r-- | src/amd/common/ac_llvm_util.h | 1 |
2 files changed, 11 insertions, 0 deletions
diff --git a/src/amd/common/ac_llvm_util.c b/src/amd/common/ac_llvm_util.c index 5b701603ebb..c8a8bf146fe 100644 --- 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 ca00540da80..18102be5207 100644 --- 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) |