diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES index 76c3266e3af7..53bbe388f7df 100644 --- a/libclc/clc/lib/amdgcn/SOURCES +++ b/libclc/clc/lib/amdgcn/SOURCES @@ -5,4 +5,6 @@ workitem/clc_get_global_offset.cl workitem/clc_get_global_size.cl workitem/clc_get_group_id.cl workitem/clc_get_local_id.cl +workitem/clc_get_local_size.cl +workitem/clc_get_max_sub_group_size.cl workitem/clc_get_work_dim.cl diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl new file mode 100644 index 000000000000..1e749404168d --- /dev/null +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// +// 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_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) { + switch (dim) { + case 0: + return __builtin_amdgcn_workgroup_size_x(); + case 1: + return __builtin_amdgcn_workgroup_size_y(); + case 2: + return __builtin_amdgcn_workgroup_size_z(); + default: + return 1; + } +} diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl new file mode 100644 index 000000000000..cc56f8d9c325 --- /dev/null +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.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_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() { + return __builtin_amdgcn_wavefrontsize(); +}