[libomptarget][nfc] Provide target_impl malloc/free

[JonChesterfield]

Summary

[libomptarget][nfc] Provide target_impl malloc/free

Sufficient to build support.cu for amdgcn

[JonChesterfield]

openmp/libomptarget/deviceRTLs/common/omptarget.h

19–20

stdlib is only used for malloc/free in the deviceRTL

openmp/libomptarget/deviceRTLs/common/src/support.cu

251

All calls to malloc/free are routed via these wrappers

[jdoerfert]

LGTM.

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

set(cuda_sources
   ${devicertl_base_directory}/common/src/critical.cu
   ${devicertl_base_directory}/common/src/data_sharing.cu
   ${devicertl_base_directory}/common/src/libcall.cu
   ${devicertl_base_directory}/common/src/loop.cu
   ${devicertl_base_directory}/common/src/omp_data.cu
   ${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)
 
 set(h_files
   ${CMAKE_CURRENT_SOURCE_DIR}/src/amdgcn_interface.h
   ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_atomics.h
   ${CMAKE_CURRENT_SOURCE_DIR}/src/target_impl.h
   ${devicertl_base_directory}/common/debug.h

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

 
 // Locks
 EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock);
 EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock);
 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
 __assertfail(const char *, const char *, unsigned, const char *, size_t);
 INLINE static void __assert_fail(const char *__message, const char *__file,
                                  unsigned int __line, const char *__function) {
   __assertfail(__message, __file, __line, __function, sizeof(char));
 }
 DEVICE int printf(const char *, ...);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/common/omptarget.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_H
 #define OMPTARGET_H
 
 // std includes
 #include <inttypes.h>
 #include <math.h>
-#include <stdlib.h>
 
 // local includes
 #include "target_impl.h"
 #include "common/debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "common/state-queue.h"
 #include "common/support.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1

openmp/libomptarget/deviceRTLs/common/src/support.cu

DEVICE unsigned long PadBytes(unsigned long size,
                               unsigned long alignment) // must be a power of 2
 {
   // compute the necessary padding to satisfy alignment constraint
   ASSERT(LT_FUSSY, (alignment & (alignment - 1)) == 0,
          "alignment %lu is not a power of 2\n", alignment);
   return (~(unsigned long)size + 1) & (alignment - 1);
 }
 
 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;
 }
 
 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;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
 
 DEVICE unsigned int *GetTeamsReductionTimestamp() {

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

 //===------------ target_impl.h - NVPTX OpenMP GPU options ------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 #ifndef _TARGET_IMPL_H_
 #define _TARGET_IMPL_H_
 
 #include <cuda.h>
+#include <stdlib.h>
+
 #include "nvptx_interface.h"
 
 #define DEVICE __device__
 #define INLINE __forceinline__ DEVICE
 #define NOINLINE __noinline__ DEVICE
 #define SHARED __shared__
 #define ALIGN(N) __align__(N)
 

 
 // Locks
 EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock);
 EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock);
 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