[PATCH] D16638: [CUDA] Added device-side std::{malloc/free}

[Artem Belevich via cfe-commits]

tra created this revision.
tra added a reviewer: jlebar.
tra added a subscriber: cfe-commits.

In addition to math functions, we also need to support std::malloc and std::free to match NVCC behavior.

http://reviews.llvm.org/D16638

Files:
  lib/Headers/__clang_cuda_cmath.h
  lib/Headers/__clang_cuda_runtime_wrapper.h

Index: lib/Headers/__clang_cuda_runtime_wrapper.h
===================================================================
--- lib/Headers/__clang_cuda_runtime_wrapper.h
+++ lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -169,6 +169,10 @@
 // CUDA headers. Alas, device_functions.hpp included below needs it.
 static inline __device__ void __brkpt(int c) { __brkpt(); }
 
+// We also need extern "C" decls for device-side allocator functions.
+extern "C" __device__ void free(void *__ptr);
+extern "C" __device__ void *malloc(size_t __size);
+
 // Now include *.hpp with definitions of various GPU functions.  Alas,
 // a lot of thins get declared/defined with __host__ attribute which
 // we don't want and we have to define it out. We also have to include
Index: lib/Headers/__clang_cuda_cmath.h
===================================================================
--- lib/Headers/__clang_cuda_cmath.h
+++ lib/Headers/__clang_cuda_cmath.h
@@ -218,6 +218,9 @@
 __DEVICE__ float trunc(float x) { return ::truncf(x); }
 __DEVICE__ double trunc(double x) { return ::trunc(x); }
 
+__DEVICE__ void free(void *__ptr) { return ::free(__ptr); }
+__DEVICE__ void *malloc(size_t __size) { return ::malloc(__size); }
+
 } // namespace std
 
 #endif


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D16638.46151.patch
Type: text/x-patch
Size: 1233 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160127/3c14820e/attachment.bin>