Skip to content

Commit cf64927

Browse files
Vedant2005goyalvgvassilev
authored andcommitted
Enable restore_tracker usage in device kernels
Replace the STL-based implementation of clad::restore_tracker with a device-compatible implementation. The previous implementation relied on std::vector and std::map, which are unavailable in device code, preventing restore_tracker from being used inside GPU kernels.
1 parent c18a748 commit cf64927

2 files changed

Lines changed: 72 additions & 0 deletions

File tree

include/clad/Differentiator/RestoreTracker.h

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,12 @@
77
#include <map>
88
#include <utility>
99
#include <vector>
10+
#ifndef Max_Records
11+
#define Max_Records 64
12+
#endif
13+
#ifndef Max_Bytes
14+
#define Max_Bytes 1024
15+
#endif
1016

1117
namespace clad {
1218

@@ -21,6 +27,41 @@ namespace clad {
2127
/// clad::tape is not viable.
2228
class restore_tracker {
2329
// m_data consists of pairs of memory addresses and bitwise values
30+
#ifdef __CUDACC__
31+
struct MetaData {
32+
char* addr;
33+
size_t size;
34+
size_t off;
35+
};
36+
MetaData m_meta[Max_Records];
37+
uint8_t m_buf[Max_Bytes];
38+
size_t m_cnt = 0, m_off = 0;
39+
40+
public:
41+
__host__ __device__ restore_tracker() = default;
42+
43+
template <typename T> __host__ __device__ void store(const T& val) {
44+
for (size_t i = 0; i < m_cnt; ++i)
45+
if (m_meta[i].addr == (char*)&val)
46+
return;
47+
48+
if (m_cnt >= Max_Records || m_off + sizeof(T) > Max_Bytes) {
49+
// Clad restore_tracker GPU capacity exceeded. Try again with larger value
50+
return;
51+
}
52+
53+
m_meta[m_cnt] = {(char*)&val, sizeof(T), m_off};
54+
std::memcpy(m_buf + m_off, &val, sizeof(T));
55+
m_off += sizeof(T);
56+
m_cnt++;
57+
}
58+
59+
__host__ __device__ void restore() {
60+
for (size_t i = 0; i < m_cnt; ++i)
61+
std::memcpy(m_meta[i].addr, m_buf + m_meta[i].off, m_meta[i].size);
62+
m_cnt = m_off = 0;
63+
}
64+
#else
2465
using RawMemory = std::vector<uint8_t>;
2566
using Address = char*;
2667
std::map<const Address, RawMemory> m_data;
@@ -47,6 +88,7 @@ class restore_tracker {
4788
}
4889
m_data.clear();
4990
}
91+
#endif
5092
};
5193
} // namespace clad
5294

test/CUDA/RestoreTrackerDevice.cu

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// RUN: %cladclang_cuda -I%S/../../include --cuda-path=%cudapath \
2+
// RUN: --cuda-gpu-arch=%cudaarch %cudaldflags -o RestoreTrackerDevice.out %s
3+
//
4+
// RUN: %cudarun ./RestoreTrackerDevice.out | %filecheck_exec %s
5+
//
6+
// REQUIRES: cuda-runtime
7+
#include "clad/Differentiator/Differentiator.h"
8+
#include <iostream>
9+
10+
__global__ void test_kernel() {
11+
clad::restore_tracker tracker;
12+
13+
double val = 3.14;
14+
tracker.store(val);
15+
val = 0.0;
16+
tracker.restore();
17+
18+
if (val == 3.14) {
19+
printf("Working on device!\n");
20+
}
21+
}
22+
23+
int main() {
24+
25+
test_kernel<<<1, 1>>>();
26+
cudaDeviceSynchronize();
27+
28+
// CHECK-EXEC: Working on device!
29+
return 0;
30+
}

0 commit comments

Comments
 (0)