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
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
