[CUDA] Added device-side system call decls and related wrappers.

[tra]

Summary

Added declarations for system calls provided by PTX.
In addition to math functions, we also need to support std::malloc and std::free to match NVCC behavior.

[jlebar]

lib/Headers/__clang_cuda_cmath.h

222 ↗

(On Diff #46151)

Not really math stuff; maybe they should live in the main header?

lib/Headers/__clang_cuda_runtime_wrapper.h

170

Perhaps add a comment explaining where these are implemented?

[tra]

Added few more device-side system calls and related wrapper functions.
Added nothrow attributes on malloc/free.

[jlebar]

lib/Headers/__clang_cuda_runtime_wrapper.h

215

It seems that only vprintf, free, malloc, and __assertfail are syscalls? If so please rejigger this comment so that it's clear what it is and isn't covering.

232

I'd prefer to keep the information from the previous comment: This declaration is there for type-safety, not because things will fail to compile if it is removed.

Otherwise someone may look at this, delete this definition, see that everything still works, and conclude that they can safely remove this line.

234

Please put "// extern "C"

236

Nit: Period

242

Please add "//namespace std"

[tra]

Addressed Justin's comments.

[jlebar]

lg with one question about printf.

lib/Headers/__clang_cuda_runtime_wrapper.h

237

I think we want an attribute on this so that we know it's printf-like, so you get appropriate format-string warnings. Unless the compiler is going to add said attribute automatically based on the just the function name.

[tra]

lib/Headers/__clang_cuda_runtime_wrapper.h

232

I don't quite see what makes vprintf() special compared to other decls here or, let's say, in /usr/include/*.h. I think 'declaration is there for type safety' applies to all of them and does not need to be spelled out for each individual decl.

237

I've verified that lang is already aware that it's printf and checks arguments against format string appropriately.

[tra]

Updated comment.

[jlebar]

lgtm. Thank you.

[tra]

Committed in r259690

lib/Headers/__clang_cuda_runtime_wrapper.h

 // has taken care of builtin variables declared in the file.
 #define __DEVICE_LAUNCH_PARAMETERS_H__
 
 // {math,device}_functions.h only have declarations of the
 // functions. We don't need them as we're going to pull in their
 // definitions from .hpp files.
 #define __DEVICE_FUNCTIONS_H__
 #define __MATH_FUNCTIONS_H__
+#define __COMMON_FUNCTIONS_H__
 
 #undef __CUDACC__
 #define __CUDABE__
 // Disables definitions of device-side runtime support stubs in
 // cuda_device_runtime_api.h
-#define __CUDADEVRT_INTERNAL__
 #include "host_config.h"
 #include "host_defines.h"
 #include "driver_types.h"
-#include "common_functions.h"
-#undef __CUDADEVRT_INTERNAL__
 
 #undef __CUDABE__
 #define __CUDACC__
 #include "cuda_runtime.h"
 
 #undef __CUDACC__
 #define __CUDABE__
 

 static inline float normcdfinv(float a) { return normcdfinvf(a); }
 static inline float normcdf(float a) { return normcdff(a); }
 static inline float erfcx(float a) { return erfcxf(a); }
 
 // For some reason single-argument variant is not always declared by
 // CUDA headers. Alas, device_functions.hpp included below needs it.
 static inline __device__ void __brkpt(int c) { __brkpt(); }
 
 // 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
 // {device,math}_functions.hpp again in order to extract the other
 // branch of #if/else inside.
 
 #define __host__
 #undef __CUDABE__
 #define __CUDACC__

 #if defined(__CUDA_ARCH__)
 // We need to emit IR declaration for non-existing __nvvm_reflect() to
 // let backend know that it should be treated as const nothrow
 // function which is what NVVMReflect pass expects to see.
 extern "C" __device__ __attribute__((const)) int __nvvm_reflect(const void *);
 static __device__ __attribute__((used)) int __nvvm_reflect_anchor() {
   return __nvvm_reflect("NONE");
 }
-
-// The nvptx vprintf syscall.  This doesn't actually need to be declared, but we
-// declare it so that if someone else declares it with a different signature,
-// we'll throw an error.
-extern "C" __device__ int vprintf(const char*, const char*);
 #endif
 
+extern "C" {
+// Device-side CUDA system calls.
+// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
+// We need these declarations and wrappers for device-side
+// malloc/free/printf calls to work without relying on
+// -fcuda-disable-target-call-checks option.
+__device__ int vprintf(const char*, const char*);
+__device__ void free(void *) __attribute((nothrow));
+__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
+__device__ void __assertfail(const char *message, const char *file,
+                             unsigned line, const char *function,
+                             size_t charSize) __attribute__((noreturn));
+
+// In order for standard assert() macro on linux to work we need to
+// provide device-side __assert_fail()
+__device__ static inline void __assert_fail(const char *message,
+                                            const char *file, unsigned line,
+                                            const char *function) {
+  __assertfail(message, file, line, function, sizeof(char));
+}
+
+// Clang will convert printf into vprintf, but we still need
+// device-side declaration for it.
+__device__ int printf(const char *, ...);
+} // extern "C"
+
+// We also need device-side std::malloc and std::free.
+namespace std {
+__device__ static inline void free(void *__ptr) { ::free(__ptr); }
+__device__ static inline void *malloc(size_t __size) {
+  return ::malloc(__size);
+}
+} // namespace std
+
 #include <__clang_cuda_cmath.h>
 
 #endif // __CUDA__
 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__