Author: Wenju He Date: 2026-03-12T14:08:34+08:00 New Revision: aa168591623260eda72c27c9b6eb90beda689d62
URL: https://github.com/llvm/llvm-project/commit/aa168591623260eda72c27c9b6eb90beda689d62 DIFF: https://github.com/llvm/llvm-project/commit/aa168591623260eda72c27c9b6eb90beda689d62.diff LOG: [libclc][amdgpu] Add clc qualifier functions with non-const pointer (#186015) Fix unresolved external functions: _Z15__clc_get_fencePv, _Z15__clc_to_globalPv, _Z14__clc_to_localPv, _Z16__clc_to_privatePv Added: libclc/clc/lib/amdgpu/address_space/clc_qualifier.inc Modified: libclc/clc/lib/amdgpu/address_space/clc_qualifier.cl Removed: ################################################################################ diff --git a/libclc/clc/lib/amdgpu/address_space/clc_qualifier.cl b/libclc/clc/lib/amdgpu/address_space/clc_qualifier.cl index fb6de7b5794b9..de18c1f0fa766 100644 --- a/libclc/clc/lib/amdgpu/address_space/clc_qualifier.cl +++ b/libclc/clc/lib/amdgpu/address_space/clc_qualifier.cl @@ -10,27 +10,12 @@ #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; -} +#define __CLC_CONST_TYPE const +#include "clc_qualifier.inc" +#undef __CLC_CONST_TYPE -_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; -} +#define __CLC_CONST_TYPE +#include "clc_qualifier.inc" +#undef __CLC_CONST_TYPE #endif diff --git a/libclc/clc/lib/amdgpu/address_space/clc_qualifier.inc b/libclc/clc/lib/amdgpu/address_space/clc_qualifier.inc new file mode 100644 index 0000000000000..946c2df5ac2ba --- /dev/null +++ b/libclc/clc/lib/amdgpu/address_space/clc_qualifier.inc @@ -0,0 +1,32 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST cl_mem_fence_flags +__clc_get_fence(__CLC_CONST_TYPE __generic void *p) { + return __builtin_amdgcn_is_shared(p) ? CLK_LOCAL_MEM_FENCE + : CLK_GLOBAL_MEM_FENCE; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_CONST_TYPE __global void * +__clc_to_global(__CLC_CONST_TYPE __generic void *p) { + return __builtin_amdgcn_is_private(p) || __builtin_amdgcn_is_shared(p) + ? NULL + : (__CLC_CONST_TYPE __global void *)p; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_CONST_TYPE __local void * +__clc_to_local(__CLC_CONST_TYPE __generic void *p) { + return __builtin_amdgcn_is_shared(p) ? (__CLC_CONST_TYPE __local void *)p + : NULL; +} + +_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_CONST_TYPE __private void * +__clc_to_private(__CLC_CONST_TYPE __generic void *p) { + return __builtin_amdgcn_is_private(p) ? (__CLC_CONST_TYPE __private void *)p + : NULL; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
