Author: Matt Arsenault Date: 2025-02-25T18:23:04+07:00 New Revision: b57e63b07a7b70ebfb5f794648e2102b7c1bd3a3
URL: https://github.com/llvm/llvm-project/commit/b57e63b07a7b70ebfb5f794648e2102b7c1bd3a3 DIFF: https://github.com/llvm/llvm-project/commit/b57e63b07a7b70ebfb5f794648e2102b7c1bd3a3.diff LOG: libclc: Stop using asm declarations for r600 on amdgcn for get_global_size (#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. Added: Modified: libclc/amdgcn/lib/workitem/get_global_size.cl Removed: ################################################################################ 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
