https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/184800
>From cf0cda96f2b5ec64ee189b2ef48b1d80efd7569c Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Thu, 5 Mar 2026 13:26:35 +0100 Subject: [PATCH 1/2] libclc: Implement get_global_linear_id and get_local_linear_id --- libclc/opencl/lib/generic/SOURCES | 2 + .../generic/workitem/get_global_linear_id.cl | 40 +++++++++++++++++++ .../generic/workitem/get_local_linear_id.cl | 15 +++++++ 3 files changed, 57 insertions(+) create mode 100644 libclc/opencl/lib/generic/workitem/get_global_linear_id.cl create mode 100644 libclc/opencl/lib/generic/workitem/get_local_linear_id.cl diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index 312657f3bf106..e4de6633d6030 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -204,4 +204,6 @@ 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 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..557c5ab3516ff --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl @@ -0,0 +1,40 @@ +//===----------------------------------------------------------------------===// +// +// 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/opencl/opencl-base.h> + +static size_t get_global_id_no_offset(uint dim) { + return get_group_id(dim) * get_local_size(dim) + 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) * get_global_size(0) + + get_global_linear_id_1d(); +} + +static size_t get_global_linear_id_3d() { + return (get_global_id_no_offset(2) * get_global_size(1) + + get_global_id_no_offset(1)) * + get_global_size(0) + + get_global_id_no_offset(0); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_global_linear_id() { + switch (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/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..10daa3d7b6bd2 --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl @@ -0,0 +1,15 @@ +//===----------------------------------------------------------------------===// +// +// 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/opencl/opencl-base.h> + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_local_linear_id() { + return (get_local_id(2) * get_local_size(1) + get_local_id(1)) * + get_local_size(0) + + get_local_id(0); +} >From 21eca3b4d48e97a15ff606251ede90a98a44cd0d Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Fri, 6 Mar 2026 15:44:49 +0100 Subject: [PATCH 2/2] Move to clc --- .../clc/workitem/clc_get_global_linear_id.h | 16 +++++++ libclc/clc/lib/generic/SOURCES | 1 + .../workitem/clc_get_global_linear_id.cl | 46 +++++++++++++++++++ .../workitem/clc_get_local_linear_id.cl | 12 ++--- .../generic/workitem/get_global_linear_id.cl | 31 +------------ .../generic/workitem/get_local_linear_id.cl | 6 +-- 6 files changed, 73 insertions(+), 39 deletions(-) create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_linear_id.h create mode 100644 libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl 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..ffd7e918a4084 --- /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_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_local_size.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_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/workitem/get_global_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl index 557c5ab3516ff..2d0f039d3a626 100644 --- a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl +++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl @@ -6,35 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/opencl-base.h> - -static size_t get_global_id_no_offset(uint dim) { - return get_group_id(dim) * get_local_size(dim) + 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) * get_global_size(0) + - get_global_linear_id_1d(); -} - -static size_t get_global_linear_id_3d() { - return (get_global_id_no_offset(2) * get_global_size(1) + - get_global_id_no_offset(1)) * - get_global_size(0) + - get_global_id_no_offset(0); -} +#include "clc/workitem/clc_get_global_linear_id.h" _CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_global_linear_id() { - switch (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; - } + 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 index 10daa3d7b6bd2..e7599fa8c75b2 100644 --- a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl +++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl @@ -6,10 +6,8 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/opencl-base.h> +#include "clc/workitem/clc_get_local_linear_id.h" _CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_local_linear_id() { - return (get_local_id(2) * get_local_size(1) + get_local_id(1)) * - get_local_size(0) + - get_local_id(0); + return __clc_get_local_linear_id(); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
