[libclc] 0c91bc6 - libclc: Stop using r600 asm intrinsic declarations for amdgcn (#181975)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-02-18T14:56:17+01:00
New Revision: 0c91bc6ae659e306bffed08652ee2f1091efcdb9

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

LOG: libclc: Stop using r600 asm intrinsic declarations for amdgcn (#181975)

Really the workitem functions should all be moved to generic code
and use gpuintrin.h. These implementations were copied from there.

Added: 
    

Modified: 
    libclc/opencl/lib/amdgcn/workitem/get_local_size.cl
    libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl

Removed: 
    


################################################################################
diff  --git a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl
index 8aa24201de573..34e4f2f1b4c19 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl
@@ -8,18 +8,14 @@
 
 #include <clc/opencl/opencl-base.h>
 
-uint __clc_amdgcn_get_local_size_x(void) __asm("llvm.r600.read.local.size.x");
-uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y");
-uint __clc_amdgcn_get_local_size_z(void) __asm("llvm.r600.read.local.size.z");
-
 _CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
   switch (dim) {
   case 0:
-    return __clc_amdgcn_get_local_size_x();
+    return __builtin_amdgcn_workgroup_size_x();
   case 1:
-    return __clc_amdgcn_get_local_size_y();
+    return __builtin_amdgcn_workgroup_size_y();
   case 2:
-    return __clc_amdgcn_get_local_size_z();
+    return __builtin_amdgcn_workgroup_size_z();
   default:
     return 1;
   }

diff  --git a/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl b/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl
index 11c1ba373aeff..9e8dddb859064 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_num_groups.cl
@@ -8,18 +8,14 @@
 
 #include <clc/opencl/opencl-base.h>
 
-uint __clc_amdgcn_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x");
-uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y");
-uint __clc_amdgcn_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z");
-
 _CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
   switch (dim) {
   case 0:
-    return __clc_amdgcn_get_num_groups_x();
+    return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
   case 1:
-    return __clc_amdgcn_get_num_groups_y();
+    return __builtin_amdgcn_grid_size_y() / __builtin_amdgcn_workgroup_size_y();
   case 2:
-    return __clc_amdgcn_get_num_groups_z();
+    return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
   default:
     return 1;
   }