[libclc] afba440 - libclc: Fix amdgpu get_enqueued_local_size (#185171)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-03-07T12:02:41+01:00
New Revision: afba44054184622afefad28905466a9b7129d27b

URL: https://github.com/llvm/llvm-project/commit/afba44054184622afefad28905466a9b7129d27b
DIFF: https://github.com/llvm/llvm-project/commit/afba44054184622afefad28905466a9b7129d27b.diff

LOG: libclc: Fix amdgpu get_enqueued_local_size (#185171)

This should not be the same as get_local_size

Added: 
    

Modified: 
    libclc/clc/lib/amdgcn/workitem/clc_get_enqueued_local_size.cl

Removed: 
    


################################################################################
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..102294e6dbb25 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
@@ -7,8 +7,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;
 }