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