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

[JonChesterfield]

Summary

[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.

[tianshilei1992]

The change itself looks good to me, but I'm just curious where you would expect to have a more definitive definition later, especially the two functions are so target specific.

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

55

unrelated change

[JonChesterfield]

I'm expecting to have a few in flight locally while trying different designs. I don't expect this to be a long lasting patch.

There's some chance trunk will end up using a different allocator to rocm, at least temporarily, in which case this might make @ronlieb's merging slightly easier.

Longer term I'm hopeful that we can drop the heap allocation from the runtime entirely, but we'll still need malloc() for applications so local experimentation would not be wasted.

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

55

git-clang-format appears to ignore *.hip . Really should get around to renaming these

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

EXTERN double __kmpc_impl_get_wtime() {
   return 0;
 }
 
 // Warp vote function
 EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   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;
 
   // Partial barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
   // reached the barrier, and one to count how many times the barrier has been
   // passed. These are packed in a single atomically accessed 32 bit integer.
unsigned long long __kmpc_atomic_exchange(unsigned long long *Address,
   return R;
 }
 unsigned long long __kmpc_atomic_add(unsigned long long *Address,
                                      unsigned long long Val) {
   return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 
 // 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));
   hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
 }
 
 EXTERN uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   return (((uint64_t)hi) << 32) | (uint64_t)lo;