diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip --- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip +++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip @@ -52,8 +52,7 @@ return __builtin_amdgcn_read_exec(); } -static void pteam_mem_barrier(uint32_t num_threads, uint32_t * barrier_state) -{ +static void pteam_mem_barrier(uint32_t num_threads, uint32_t *barrier_state) { __atomic_thread_fence(__ATOMIC_ACQUIRE); uint32_t num_waves = (num_threads + WARPSIZE - 1) / WARPSIZE; @@ -178,8 +177,12 @@ } // Stub implementations -EXTERN void *__kmpc_impl_malloc(size_t) { return nullptr; } -EXTERN void __kmpc_impl_free(void *) {} +// Weak to allow overriding by local versions while comparing different +// potential implementations +__attribute__((weak)) EXTERN void *__kmpc_impl_malloc(size_t) { + return nullptr; +} +__attribute__((weak)) EXTERN void __kmpc_impl_free(void *) {} EXTERN void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) { lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));