[libclc] libclc: Stop using asm declarations for r600 on amdgcn for get_global_size (PR #128692)

[Matt Arsenault via cfe-commits]

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

Comparing the case where each dimension is used alone, the only codegen difference
is a missed addressing mode fold for the constant offset in the old version due to
an ancient bug.

>From a6b39981322eb0d603852e8330776756f740db5b Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Tue, 25 Feb 2025 17:28:27 +0700
Subject: [PATCH] libclc: Stop using asm declarations for r600 on amdgcn for
 get_global_size

Comparing the case where each dimension is used alone, the only codegen difference
is a missed addressing mode fold for the constant offset in the old version due to
an ancient bug.
---
 libclc/amdgcn/lib/workitem/get_global_size.cl | 10 +++-------
 1 file changed, 3 insertions(+), 7 deletions(-)

diff --git a/libclc/amdgcn/lib/workitem/get_global_size.cl b/libclc/amdgcn/lib/workitem/get_global_size.cl
index 2f28ca6066654..0fec7e24966fd 100644
--- a/libclc/amdgcn/lib/workitem/get_global_size.cl
+++ b/libclc/amdgcn/lib/workitem/get_global_size.cl
@@ -1,17 +1,13 @@
 #include <clc/clc.h>
 
-uint __clc_amdgcn_get_global_size_x(void) __asm("llvm.r600.read.global.size.x");
-uint __clc_amdgcn_get_global_size_y(void) __asm("llvm.r600.read.global.size.y");
-uint __clc_amdgcn_get_global_size_z(void) __asm("llvm.r600.read.global.size.z");
-
 _CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
   switch (dim) {
   case 0:
-    return __clc_amdgcn_get_global_size_x();
+    return __builtin_amdgcn_grid_size_x();
   case 1:
-    return __clc_amdgcn_get_global_size_y();
+    return __builtin_amdgcn_grid_size_y();
   case 2:
-    return __clc_amdgcn_get_global_size_z();
+    return __builtin_amdgcn_grid_size_z();
   default:
     return 1;
   }