[libclc] libclc: Implement get_sub_group_id with get_local_linear_id (PR #188713)

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/188713

None

>From 066af6166f4a8ab712ca7dd0b9a07dfe39be3318 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 26 Mar 2026 10:23:06 +0100
Subject: [PATCH] libclc: Implement get_sub_group_id with get_local_linear_id

---
 .../lib/generic/workitem/clc_get_sub_group_id.cl | 16 +++-------------
 1 file changed, 3 insertions(+), 13 deletions(-)

diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
index 681151b5e92c5..02391c52ca813 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
@@ -6,19 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "clc/workitem/clc_get_local_id.h"
-#include "clc/workitem/clc_get_local_size.h"
+#include "clc/workitem/clc_get_local_linear_id.h"
 #include "clc/workitem/clc_get_max_sub_group_size.h"
-#include "clc/workitem/clc_get_sub_group_id.h"
 
-_CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_id() {
-  // sreg.warpid is volatile and doesn't represent virtual warp index
-  // see https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
-  size_t id_x = __clc_get_local_id(0);
-  size_t id_y = __clc_get_local_id(1);
-  size_t id_z = __clc_get_local_id(2);
-  size_t size_x = __clc_get_local_size(0);
-  size_t size_y = __clc_get_local_size(1);
-  uint sg_size = __clc_get_max_sub_group_size();
-  return (id_z * size_y * size_x + id_y * size_x + id_x) / sg_size;
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_id(void) {
+  return __clc_get_local_linear_id() / __clc_get_max_sub_group_size();
 }