Author: Wenju He Date: 2025-08-18T06:51:01+08:00 New Revision: bce14c69db0ba2b8c54b250a0bd309879fb5bc78
URL: https://github.com/llvm/llvm-project/commit/bce14c69db0ba2b8c54b250a0bd309879fb5bc78 DIFF: https://github.com/llvm/llvm-project/commit/bce14c69db0ba2b8c54b250a0bd309879fb5bc78.diff LOG: [libclc] Fix out-of-bound value for workitem functions according to OpenCL spec (#153784) Added: Modified: libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl Removed: ################################################################################ diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl index aea927c3460b4..4dab7905ba301 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl @@ -17,6 +17,6 @@ _CLC_DEF _CLC_OVERLOAD size_t __clc_get_group_id(uint dim) { case 2: return __builtin_amdgcn_workgroup_id_z(); default: - return 1; + return 0; } } diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl index b7b7a43e735d3..8d1d16d4762fd 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl @@ -17,6 +17,6 @@ _CLC_DEF _CLC_OVERLOAD size_t __clc_get_local_id(uint dim) { case 2: return __builtin_amdgcn_workitem_id_z(); default: - return 1; + return 0; } } diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl index 4525c85f1e382..2d547d3af249e 100644 --- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl +++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl @@ -17,6 +17,6 @@ _CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) { case 2: return __nvvm_read_ptx_sreg_ntid_z(); default: - return 0; + return 1; } } diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl index 495864751ef68..87151ea4be625 100644 --- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl +++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl @@ -17,6 +17,6 @@ _CLC_OVERLOAD _CLC_DEF size_t __clc_get_num_groups(uint dim) { case 2: return __nvvm_read_ptx_sreg_nctaid_z(); default: - return 0; + return 1; } } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
