https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/184800
None >From 4013343b2c623d74af9062face97422860a7cc56 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Thu, 5 Mar 2026 13:26:35 +0100 Subject: [PATCH] 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 bb5e8ab08a711..85eb5ba40213f 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -200,4 +200,6 @@ shared/min.cl shared/vload.cl shared/vstore.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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
