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