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.arsena...@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;
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
