[Openmp-commits] [openmp] 63e2aa5 - [libomptarget][nfc] Provide target_impl malloc/free

[Jon Chesterfield via Openmp-commits]

Author: Jon Chesterfield
Date: 2019-12-19T16:54:28Z
New Revision: 63e2aa5658bd8a4905ae5a85f9046250a5e16a86

URL: https://github.com/llvm/llvm-project/commit/63e2aa5658bd8a4905ae5a85f9046250a5e16a86
DIFF: https://github.com/llvm/llvm-project/commit/63e2aa5658bd8a4905ae5a85f9046250a5e16a86.diff

LOG: [libomptarget][nfc] Provide target_impl malloc/free

Summary:
[libomptarget][nfc] Provide target_impl malloc/free

Sufficient to build support.cu for amdgcn

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: jdoerfert

Subscribers: jvesely, mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71685

Added: 
    

Modified: 
    openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
    openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
    openmp/libomptarget/deviceRTLs/common/omptarget.h
    openmp/libomptarget/deviceRTLs/common/src/support.cu
    openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
index d3df65b734df..4395ccf80b5b 100644
--- a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
@@ -64,6 +64,7 @@ set(cuda_sources
   ${devicertl_base_directory}/common/src/omptarget.cu
   ${devicertl_base_directory}/common/src/parallel.cu
   ${devicertl_base_directory}/common/src/reduction.cu
+  ${devicertl_base_directory}/common/src/support.cu
   ${devicertl_base_directory}/common/src/sync.cu
   ${devicertl_base_directory}/common/src/task.cu)
 

diff  --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
index 9b7b8e3fd95d..ec72800cd69d 100644
--- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
@@ -164,6 +164,10 @@ EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock);
 EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock);
 EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock);
 
+// Memory
+EXTERN void *__kmpc_impl_malloc(size_t x);
+EXTERN void __kmpc_impl_free(void *x);
+
 // DEVICE versions of part of libc
 extern "C" {
 DEVICE __attribute__((noreturn)) void

diff  --git a/openmp/libomptarget/deviceRTLs/common/omptarget.h b/openmp/libomptarget/deviceRTLs/common/omptarget.h
index 8ffa31091d4c..8f577f7364d9 100644
--- a/openmp/libomptarget/deviceRTLs/common/omptarget.h
+++ b/openmp/libomptarget/deviceRTLs/common/omptarget.h
@@ -17,7 +17,6 @@
 // std includes
 #include <inttypes.h>
 #include <math.h>
-#include <stdlib.h>
 
 // local includes
 #include "target_impl.h"

diff  --git a/openmp/libomptarget/deviceRTLs/common/src/support.cu b/openmp/libomptarget/deviceRTLs/common/src/support.cu
index ea1fc3841ad1..d7a0b23667fa 100644
--- a/openmp/libomptarget/deviceRTLs/common/src/support.cu
+++ b/openmp/libomptarget/deviceRTLs/common/src/support.cu
@@ -250,7 +250,7 @@ DEVICE unsigned long PadBytes(unsigned long size,
 
 DEVICE void *SafeMalloc(size_t size, const char *msg) // check if success
 {
-  void *ptr = malloc(size);
+  void *ptr = __kmpc_impl_malloc(size);
   PRINT(LD_MEM, "malloc data of size %llu for %s: 0x%llx\n",
         (unsigned long long)size, msg, (unsigned long long)ptr);
   return ptr;
@@ -258,7 +258,7 @@ DEVICE void *SafeMalloc(size_t size, const char *msg) // check if success
 
 DEVICE void *SafeFree(void *ptr, const char *msg) {
   PRINT(LD_MEM, "free data ptr 0x%llx for %s\n", (unsigned long long)ptr, msg);
-  free(ptr);
+  __kmpc_impl_free(ptr);
   return NULL;
 }
 

diff  --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
index 6d5852882683..c11fe253cc7c 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -13,6 +13,8 @@
 #define _TARGET_IMPL_H_
 
 #include <cuda.h>
+#include <stdlib.h>
+
 #include "nvptx_interface.h"
 
 #define DEVICE __device__
@@ -204,4 +206,8 @@ EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock);
 EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock);
 EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock);
 
+// Memory
+INLINE void *__kmpc_impl_malloc(size_t x) { return malloc(x); }
+INLINE void __kmpc_impl_free(void *x) { free(x); }
+
 #endif