[Openmp-commits] [PATCH] D102499: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation
Jon Chesterfield via Phabricator via Openmp-commits
openmp-commits at lists.llvm.org
Fri May 14 07:46:34 PDT 2021
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>
More information about the Openmp-commits
mailing list