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

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, grokos, ABataev, ronlieb, pdhaliwal, tianshilei1992, ye-luo.
Herald added subscribers: t-tye, tpr, dstuttard, yaxunl, jvesely, kzhuravl.
JonChesterfield requested review of this revision.
Herald added subscribers: openmp-commits, sstefan1, wdng.
Herald added a project: OpenMP.

[libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation

There are a lot of different ways we might implement the devicertl local alloc
and free functions. Via host, local buffers (stack or arena), specialising per
kernel etc. It is not yet clear what the right design is. This change makes the
alloc and free functions weak, so one can override them from local tests while
comparing options.

Not strictly necessary, as a comparable patch can be applied locally each time,
but would be convenient for out of tree dev. Plan would be to drop the weak
attribute at the same time as introducing a working allocator to trunk.


Repository:
  rG LLVM Github Monorepo

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.345440.patch
Type: text/x-patch
Size: 1133 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210514/719830b3/attachment-0001.bin>