Modify MoveData in Tape.h to make it CUDA compatible

* Add addressof function make the call to std::address CUDA compatible
* Change fprintf to printf
* Add trap function to replace exit function

include/clad/Differentiator/CladConfig.h

@@ -3,10 +3,46 @@
 #ifndef CLAD_CONFIG_H
 #define CLAD_CONFIG_H
 
+#include <cstdlib>
+#include <memory>
+
+// Define CUDA_HOST_DEVICE attribute for adding CUDA support to
+// clad functions
 #ifdef __CUDACC__
 #define CUDA_HOST_DEVICE __host__ __device__
 #else
 #define CUDA_HOST_DEVICE
 #endif
 
+// Define trap function that is a CUDA compatible replacement for
+// exit(int code) function
+#ifdef __CUDACC__
+__device__ void trap(int code) {
+  asm("trap;");
+}
+__host__ void trap(int code) {
+  exit(code);
+}
+#else
+void trap(int code) {
+  exit(code);
+}
+#endif
+
+#ifdef  __CUDACC__
+template<typename T>
+__device__ T* addressof(T& r) {
+  return __builtin_addressof(r);
+}
+template<typename T>
+__host__ T* addressof(T& r) {
+  return std::addressof(r);
+}
+#else
+template<typename T>
+T* addressof(T& r) {
+  return std::addressof(r);
+}
+#endif
+
 #endif // CLAD_CONFIG_H

include/clad/Differentiator/Tape.h

@@ -100,13 +100,13 @@ namespace clad {
       // allocation properly.
       for (; first != last; ++first, (void)++current) {
         auto new_data = ::new (const_cast<void*>(
-            static_cast<const volatile void*>(std::addressof(*current))))
+            static_cast<const volatile void*>(addressof(current))))
             T(std::move(*first));
         if (!new_data) {
           // clean up the memory mess just in case!
           destroy(d_first, current);
-          fprintf(stderr, "Allocation failure during tape resize! Aborting.");
-          exit(EXIT_FAILURE);
+          printf("Allocation failure during tape resize! Aborting.");
+          trap(EXIT_FAILURE);
         }
       }
     }