https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185171
This should not be the same as get_local_size >From 3adbe80894df1841b8c6e9724773730f7e5f2f9c Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Sat, 7 Mar 2026 10:39:39 +0100 Subject: [PATCH] libclc: Fix amdgpu get_enqueued_local_size This should not be the same as get_local_size --- .../clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl index c7226241694b3..9426b2624caf2 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl @@ -8,7 +8,11 @@ #include "clc/workitem/clc_get_enqueued_local_size.h" #include "clc/workitem/clc_get_local_size.h" +#include <amdhsa_abi.h> _CLC_OVERLOAD _CLC_DEF size_t __clc_get_enqueued_local_size(uint dim) { - return __clc_get_local_size(dim); + __constant amdhsa_implicit_kernarg_v5 *args = + (__constant amdhsa_implicit_kernarg_v5 *) + __builtin_amdgcn_implicitarg_ptr(); + return dim < 3 ? args->group_size[dim] : 1; } _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
