================ @@ -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()) { ---------------- wenju-he wrote:
could you move implementation into __clc_get_global_linear_id and call __clc_get_global_linear_id here? https://github.com/llvm/llvm-project/pull/184800 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
