[Openmp-commits] [PATCH] D102499: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation

[Jon Chesterfield via Phabricator via Openmp-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGd54712ab4deb: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation (authored by JonChesterfield).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D102499/new/

https://reviews.llvm.org/D102499

Files:
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip


Index: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
===================================================================
--- openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
+++ 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));


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D102499.347039.patch
Type: text/x-patch
Size: 1133 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210521/eb31b54b/attachment.bin>