Add CUDA_HOST_DEVICE to Tape::destroy_element() allowing the function to be called from device code too

Fixes #1624

Author: Vedant2005goyal

include/clad/Differentiator/Tape.h

@@ -457,12 +457,12 @@ class tape_impl {
 #endif
   }
 
-  DiskInfo& getDiskInfo() {
+  CUDA_HOST_DEVICE DiskInfo& getDiskInfo() {
     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast)
     return *reinterpret_cast<DiskInfo*>(&m_state);
   }
 
-  void check_and_evict_impl(std::true_type) {
+  CUDA_HOST_DEVICE void check_and_evict_impl(std::true_type) {
     DiskInfo& info = getDiskInfo();
     if (GpuOffload && info.m_ActiveVramSlabs >= info.m_MaxVramSlabs) {
       Slab* candidate = m_head;
@@ -539,9 +539,9 @@ class tape_impl {
       }
     }
   }
-  void check_and_evict_impl(std::false_type) {}
+  CUDA_HOST_DEVICE void check_and_evict_impl(std::false_type) {}
 
-  void ensure_loaded_impl(Slab* slab, std::true_type) {
+  CUDA_HOST_DEVICE void ensure_loaded_impl(Slab* slab, std::true_type) {
     DiskInfo& info = getDiskInfo();
     if (GpuOffload) {
       // Already in host RAM — return immediately
@@ -604,14 +604,14 @@ class tape_impl {
     }
   }
 
-  void ensure_loaded_impl(Slab* slab, std::false_type) {}
+  CUDA_HOST_DEVICE void ensure_loaded_impl(Slab* slab, std::false_type) {}
 
-  void check_and_evict() {
+  CUDA_HOST_DEVICE void check_and_evict() {
     check_and_evict_impl(std::integral_constant < bool,
                          DiskOffload || GpuOffload > {});
   }
 
-  void ensure_loaded(Slab* slab) {
+  CUDA_HOST_DEVICE void ensure_loaded(Slab* slab) {
     ensure_loaded_impl(slab, std::integral_constant < bool,
                        DiskOffload || GpuOffload > {});
   }
@@ -923,8 +923,11 @@ class tape_impl {
     m_capacity = SBO_SIZE;
   }
 
-  template <typename ElTy> void destroy_element(ElTy* elem) { elem->~ElTy(); }
-  template <typename ElTy, size_t N> void destroy_element(ElTy (*arr)[N]) {
+  template <typename ElTy> CUDA_HOST_DEVICE void destroy_element(ElTy* elem) {
+    elem->~ElTy();
+  }
+  template <typename ElTy, size_t N>
+  CUDA_HOST_DEVICE void destroy_element(ElTy (*arr)[N]) {
     for (size_t i = 0; i < N; ++i)
       (*arr)[i].~ElTy();
   }

test/Regressions/issue-1624.cu

@@ -0,0 +1,31 @@
+// RUN: %cladclang_cuda -fsyntax-only -I%S/../../include --cuda-path=%cudapath \
+// RUN:     --cuda-gpu-arch=%cudaarch -Xclang -verify %s
+//
+// REQUIRES: cuda-runtime
+// expected-no-diagnostics
+
+#include "clad/Differentiator/Tape.h"
+
+struct TrackedValue {
+    double val;
+    __host__ __device__ TrackedValue() : val(0.0) {}
+    
+    __host__ __device__ ~TrackedValue() {}
+};
+
+template <class T>
+__device__ void exercise_tape(clad::tape_impl<T>& t) {
+  t.emplace_back();
+  
+  auto& back = t.back();
+  (void)back;
+  
+  // This explicitly triggers the `destroy_element` template 
+  // via `pop_back()`.
+  t.pop_back(); 
+}
+
+__global__ void my_kernel() {
+  clad::tape_impl<TrackedValue> t;
+  exercise_tape(t);
+}