https://github.com/wenju-he created https://github.com/llvm/llvm-project/pull/186015
Fix unresolved external functions: _Z15__clc_get_fencePv, _Z15__clc_to_globalPv, _Z14__clc_to_localPv, _Z16__clc_to_privatePv >From 37162fb42254a5fd14db11cefe5cb85b0157c586 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Thu, 12 Mar 2026 01:55:33 +0100 Subject: [PATCH] [libclc][amdgpu] Add missing qualifier clc functions with non-const pointer Fix unresolved external functions: _Z15__clc_get_fencePv, _Z15__clc_to_globalPv, _Z14__clc_to_localPv, _Z16__clc_to_privatePv --- .../lib/amdgpu/address_space/clc_qualifier.cl | 27 ++++------------ .../amdgpu/address_space/clc_qualifier.inc | 32 +++++++++++++++++++ 2 files changed, 38 insertions(+), 21 deletions(-) create mode 100644 libclc/clc/lib/amdgpu/address_space/clc_qualifier.inc 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
