shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
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.
This commit is contained in:
Matt Arsenault 2025-02-25 18:23:04 +07:00 committed by GitHub
parent f95ad44068
commit b57e63b07a
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 3 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
libclc/amdgcn/lib/workitem/get_global_size.cl
View File

@ -1,17 +1,13 @@
#include <clc/clc.h> #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) { _CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
switch (dim) { switch (dim) {
case 0: case 0:
return __clc_amdgcn_get_global_size_x(); return __builtin_amdgcn_grid_size_x();
case 1: case 1:
return __clc_amdgcn_get_global_size_y(); return __builtin_amdgcn_grid_size_y();
case 2: case 2:
return __clc_amdgcn_get_global_size_z(); return __builtin_amdgcn_grid_size_z();
default: default:
return 1; return 1;
} }