Author: Matt Arsenault Date: 2026-03-05T12:30:09+01:00 New Revision: e67360ec319a642e56abb0c3f0c3e627a98dd225
URL: https://github.com/llvm/llvm-project/commit/e67360ec319a642e56abb0c3f0c3e627a98dd225 DIFF: https://github.com/llvm/llvm-project/commit/e67360ec319a642e56abb0c3f0c3e627a98dd225.diff LOG: libclc: Implement address space qualifier functions for amdgpu (#184766) Added: libclc/clc/include/clc/address_space/qualifier.h libclc/clc/lib/amdgcn/address_space/qualifier.cl libclc/clc/lib/generic/shared/clc_qualifier.cl libclc/opencl/lib/generic/address_space/qualifier.cl Modified: libclc/clc/lib/amdgcn/SOURCES libclc/clc/lib/generic/SOURCES libclc/opencl/lib/generic/SOURCES Removed: ################################################################################ diff --git a/libclc/clc/include/clc/address_space/qualifier.h b/libclc/clc/include/clc/address_space/qualifier.h new file mode 100644 index 0000000000000..28b8ab33d5df8 --- /dev/null +++ b/libclc/clc/include/clc/address_space/qualifier.h @@ -0,0 +1,33 @@ +//===----------------------------------------------------------------------===// +// +// 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_ADDRESS_SPACE_CLC_QUALIFIER_H__ +#define __CLC_ADDRESS_SPACE_CLC_QUALIFIER_H__ + +#include <clc/clcfunc.h> + +#if _CLC_GENERIC_AS_SUPPORTED + +_CLC_OVERLOAD _CLC_DECL _CLC_CONST cl_mem_fence_flags __clc_get_fence(void *p); +_CLC_OVERLOAD _CLC_DECL _CLC_CONST cl_mem_fence_flags +__clc_get_fence(const void *p); + +_CLC_OVERLOAD _CLC_DECL _CLC_CONST __global void *__clc_to_global(void *p); +_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __global void * +__clc_to_global(const void *p); + +_CLC_OVERLOAD _CLC_DECL _CLC_CONST __local void *__clc_to_local(void *p); +_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __local void * +__clc_to_local(const void *p); + +_CLC_OVERLOAD _CLC_DECL _CLC_CONST __private void *__clc_to_private(void *p); +_CLC_OVERLOAD _CLC_DECL _CLC_CONST const __private void * +__clc_to_private(const void *p); +#endif // _CLC_GENERIC_AS_SUPPORTED + +#endif // __CLC_ADDRESS_SPACE_CLC_QUALIFIER_H__ diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES index 5040bf1158674..b4557b0a26f70 100644 --- a/libclc/clc/lib/amdgcn/SOURCES +++ b/libclc/clc/lib/amdgcn/SOURCES @@ -1,3 +1,4 @@ +address_space/qualifier.cl math/clc_ldexp.cl mem_fence/clc_mem_fence.cl synchronization/clc_work_group_barrier.cl diff --git a/libclc/clc/lib/amdgcn/address_space/qualifier.cl b/libclc/clc/lib/amdgcn/address_space/qualifier.cl new file mode 100644 index 0000000000000..82d9f5351c446 --- /dev/null +++ b/libclc/clc/lib/amdgcn/address_space/qualifier.cl @@ -0,0 +1,36 @@ +//===----------------------------------------------------------------------===// +// +// 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/address_space/qualifier.h> + +#if _CLC_GENERIC_AS_SUPPORTED + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags +__clc_get_fence(const __generic void *p) { + return __builtin_amdgcn_is_shared(p) ? CLK_LOCAL_MEM_FENCE + : CLK_GLOBAL_MEM_FENCE; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __global void * +__clc_to_global(const __generic void *p) { + return __builtin_amdgcn_is_private(p) || __builtin_amdgcn_is_shared(p) + ? NULL + : (const __global void *)p; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __local void * +__clc_to_local(const __generic void *p) { + return __builtin_amdgcn_is_shared(p) ? (__local void *)p : NULL; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST const __private void * +__clc_to_private(const __generic void *p) { + return __builtin_amdgcn_is_private(p) ? (__private void *)p : NULL; +} + +#endif diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES index ac820239baa17..421eb638720e3 100644 --- a/libclc/clc/lib/generic/SOURCES +++ b/libclc/clc/lib/generic/SOURCES @@ -173,6 +173,7 @@ relational/clc_signbit.cl shared/clc_clamp.cl shared/clc_max.cl shared/clc_min.cl +shared/clc_qualifier.cl shared/clc_vload.cl shared/clc_vstore.cl workitem/clc_get_local_linear_id.cl diff --git a/libclc/clc/lib/generic/shared/clc_qualifier.cl b/libclc/clc/lib/generic/shared/clc_qualifier.cl new file mode 100644 index 0000000000000..978f6677798cd --- /dev/null +++ b/libclc/clc/lib/generic/shared/clc_qualifier.cl @@ -0,0 +1,33 @@ +//===----------------------------------------------------------------------===// +// +// 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/address_space/qualifier.h> + +#if _CLC_GENERIC_AS_SUPPORTED + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags +__clc_get_fence(__generic void *p) { + return __clc_get_fence((const __generic void *)p); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __global void * +__clc_to_global(__generic void *p) { + return (__global void *)__clc_to_global((const __generic void *)p); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __local void * +__clc_to_local(__generic void *p) { + return (__local void *)__clc_to_local((const __generic void *)p); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __private void * +__clc_to_private(__generic void *p) { + return (__private void *)__clc_to_private((const __generic void *)p); +} + +#endif diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index c820c6c3c0c06..bb5e8ab08a711 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -1,3 +1,4 @@ +address_space/qualifier.cl subnormal_config.cl async/async_work_group_copy.cl async/async_work_group_strided_copy.cl diff --git a/libclc/opencl/lib/generic/address_space/qualifier.cl b/libclc/opencl/lib/generic/address_space/qualifier.cl new file mode 100644 index 0000000000000..8ad4a373465bd --- /dev/null +++ b/libclc/opencl/lib/generic/address_space/qualifier.cl @@ -0,0 +1,35 @@ +//===----------------------------------------------------------------------===// +// +// 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/address_space/qualifier.h" + +#if _CLC_GENERIC_AS_SUPPORTED + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST cl_mem_fence_flags +get_fence(__generic void *p) { + return __clc_get_fence(p); +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags +get_fence(const __generic void *p) { + return __clc_get_fence(p); +} + +_CLC_DEF _CLC_CONST __global void *__to_global(__generic void *p) { + return __clc_to_global(p); +} + +_CLC_DEF _CLC_CONST __local void *__to_local(__generic void *p) { + return __clc_to_local(p); +} + +_CLC_DEF _CLC_CONST __private void *__to_private(__generic void *p) { + return __clc_to_private(p); +} + +#endif // _CLC_GENERIC_AS_SUPPORTED _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
