[llvm-branch-commits] [libclc] libclc: Fix amdgpu get_enqueued_local_size (PR #185171)

[Matt Arsenault via llvm-branch-commits]

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 <Matthew.Arsenault at amd.com>
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;
 }