Author: Matt Arsenault Date: 2026-03-07T09:42:51Z New Revision: 704c87bb948aff1bec718d56ad52b9b5d9c49cfb
URL: https://github.com/llvm/llvm-project/commit/704c87bb948aff1bec718d56ad52b9b5d9c49cfb DIFF: https://github.com/llvm/llvm-project/commit/704c87bb948aff1bec718d56ad52b9b5d9c49cfb.diff LOG: libclc: Implement get_global_linear_id and get_local_linear_id (#184800) Added: libclc/clc/include/clc/workitem/clc_get_global_linear_id.h libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl libclc/opencl/lib/generic/workitem/get_global_linear_id.cl libclc/opencl/lib/generic/workitem/get_local_linear_id.cl Modified: libclc/clc/lib/generic/SOURCES libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl libclc/opencl/lib/generic/SOURCES Removed: ################################################################################ diff --git a/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h b/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h new file mode 100644 index 0000000000000..fab5040e497ad --- /dev/null +++ b/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__ +#define __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__ + +#include "clc/internal/clc.h" + +_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_global_linear_id(); + +#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__ diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES index 421eb638720e3..b352382fbe6d9 100644 --- a/libclc/clc/lib/generic/SOURCES +++ b/libclc/clc/lib/generic/SOURCES @@ -176,6 +176,7 @@ shared/clc_min.cl shared/clc_qualifier.cl shared/clc_vload.cl shared/clc_vstore.cl +workitem/clc_get_global_linear_id.cl workitem/clc_get_local_linear_id.cl workitem/clc_get_num_sub_groups.cl workitem/clc_get_sub_group_id.cl diff --git a/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl new file mode 100644 index 0000000000000..0a7d5ed57f1e8 --- /dev/null +++ b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "clc/workitem/clc_get_enqueued_local_size.h" +#include "clc/workitem/clc_get_global_linear_id.h" +#include "clc/workitem/clc_get_global_size.h" +#include "clc/workitem/clc_get_group_id.h" +#include "clc/workitem/clc_get_local_id.h" +#include "clc/workitem/clc_get_work_dim.h" + +static size_t get_global_id_no_offset(uint dim) { + return __clc_get_group_id(dim) * __clc_get_enqueued_local_size(dim) + + __clc_get_local_id(dim); +} + +static size_t get_global_linear_id_1d() { return get_global_id_no_offset(0); } + +static size_t get_global_linear_id_2d() { + return get_global_id_no_offset(1) * __clc_get_global_size(0) + + get_global_linear_id_1d(); +} + +static size_t get_global_linear_id_3d() { + return (get_global_id_no_offset(2) * __clc_get_global_size(1) + + get_global_id_no_offset(1)) * + __clc_get_global_size(0) + + get_global_id_no_offset(0); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t __clc_get_global_linear_id() { + switch (__clc_get_work_dim()) { + case 1: + return get_global_linear_id_1d(); + case 2: + return get_global_linear_id_2d(); + case 3: + return get_global_linear_id_3d(); + default: + return 0; + } +} diff --git a/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl b/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl index fd7905568d595..a8fcb0059f768 100644 --- a/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl +++ b/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl @@ -6,13 +6,13 @@ // //===----------------------------------------------------------------------===// -#include <clc/workitem/clc_get_local_id.h> -#include <clc/workitem/clc_get_local_linear_id.h> -#include <clc/workitem/clc_get_local_size.h> +#include "clc/workitem/clc_get_local_id.h" +#include "clc/workitem/clc_get_local_linear_id.h" +#include "clc/workitem/clc_get_local_size.h" -_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_linear_id() { - return __clc_get_local_id(2) * __clc_get_local_size(1) * +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t __clc_get_local_linear_id() { + return (__clc_get_local_id(2) * __clc_get_local_size(1) + + __clc_get_local_id(1)) * __clc_get_local_size(0) + - __clc_get_local_id(1) * __clc_get_local_size(0) + __clc_get_local_id(0); } diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index dd51d5c084d51..cdc1d8321dfeb 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -204,5 +204,7 @@ subgroup/sub_group_broadcast.cl synchronization/work_group_barrier.cl workitem/get_enqueued_local_size.cl workitem/get_global_id.cl +workitem/get_global_linear_id.cl workitem/get_global_size.cl +workitem/get_local_linear_id.cl workitem/get_num_groups.cl diff --git a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl new file mode 100644 index 0000000000000..2d0f039d3a626 --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl @@ -0,0 +1,13 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "clc/workitem/clc_get_global_linear_id.h" + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_global_linear_id() { + return __clc_get_global_linear_id(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl new file mode 100644 index 0000000000000..e7599fa8c75b2 --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl @@ -0,0 +1,13 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "clc/workitem/clc_get_local_linear_id.h" + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_local_linear_id() { + return __clc_get_local_linear_id(); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
