Skip to content

Commit d4145ec

Browse files
intwanghaointwanghao
authored
[SYCLomatic] Impltement codepin generate data flow graph feature (#2160)
 Signed-off-by: Wang, Hao3 <[email protected]> Signed-off-by: intwanghao <[email protected]>
1 parent d8ba8ec commit d4145ec

File tree

20 files changed

+385
-27
lines changed

20 files changed

+385
-27
lines changed

‎clang/lib/DPCT/ASTTraversal.cpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8361,11 +8361,14 @@ if (CodePinInstrumentation.find(KCallSpellingRange.first) !=
83618361
}
83628362
}
83638363
}
8364+
std::string KernelName =
8365+
KCall->getCalleeDecl()->getAsFunction()->getNameAsString();
83648366

83658367
auto InstrumentKernel = [&](std::string StreamStr, HeaderType HT,
83668368
dpct::ReplacementType CodePinType) {
83678369
std::string CodePinKernelArgsString = "(\"";
8368-
CodePinKernelArgsString += llvm::sys::path::convert_to_slash(
8370+
CodePinKernelArgsString += KernelName + ":" +
8371+
llvm::sys::path::convert_to_slash(
83698372
KCallSpellingRange.first.printToString(SM)) +
83708373
"\", ";
83718374
CodePinKernelArgsString += StreamStr;

‎clang/runtime/dpct-rt/include/dpct/codepin/codepin.hpp‎

Lines changed: 52 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -46,13 +46,25 @@ class logger {
4646
}
4747

4848
template <class... Args>
49-
void print_CP(const std::string &cp_id, size_t free_byte,
50-
size_t total_byte, float elapse_time,
49+
void print_CP(const std::string &cp_id, std::string device_name,
50+
size_t free_byte, size_t total_byte, float elapse_time,
5151
queue_t queue, Args... args) {
5252
ptr_unique.clear();
5353
auto obj = arr.object();
5454
obj.key("ID");
5555
obj.value(cp_id);
56+
obj.key("Device Name");
57+
obj.value(device_name);
58+
obj.key("Device ID");
59+
#ifdef __NVCC__
60+
int device_id;
61+
cudaGetDevice(&device_id);
62+
obj.value(device_id);
63+
#else
64+
obj.value((int)dpct::get_current_device_id());
65+
#endif
66+
obj.key("Stream Address");
67+
obj.value((void *)queue);
5668
obj.key("Free Device Memory");
5769
obj.value(free_byte);
5870
obj.key("Total Device Memory");
@@ -62,25 +74,36 @@ class logger {
6274
obj.key("CheckPoint");
6375
auto cp_obj =
6476
obj.value<detail::json_stringstream::json_obj>();
65-
print_args(cp_obj, queue, args...);
77+
print_args(cp_obj, queue, 0, args...);
6678
}
6779

68-
void print_args(json_stringstream::json_obj &obj,
69-
queue_t queue) {}
80+
void print_args(json_stringstream::json_obj &obj, queue_t queue,
81+
int index = 0) {}
7082

7183
template <class First, class... RestArgs>
72-
void print_args(json_stringstream::json_obj &obj,
73-
queue_t queue, std::string_view arg_name,
74-
First &arg, RestArgs... args) {
84+
void print_args(json_stringstream::json_obj &obj, queue_t queue, int index,
85+
std::string_view arg_name, First &arg, RestArgs... args) {
7586
obj.key(arg_name);
7687
{
7788
auto type_obj =
7889
obj.value<detail::json_stringstream::json_obj>();
7990
detail::data_ser<First>::print_type_name(type_obj);
91+
type_obj.key("Address");
92+
print_address(type_obj, arg);
93+
type_obj.key("Index");
94+
type_obj.value(index);
8095
type_obj.key("Data");
8196
detail::data_ser<First>::dump(json_ss, arg, queue);
8297
}
83-
print_args(obj, queue, args...);
98+
print_args(obj, queue, index + 1, args...);
99+
}
100+
template <class ArgT>
101+
void print_address(json_stringstream::json_obj &obj, ArgT arg) {
102+
if constexpr (std::is_pointer<ArgT>::value) {
103+
obj.value((void *)arg);
104+
} else {
105+
obj.value((void *)&arg);
106+
}
84107
}
85108

86109
private:
@@ -186,10 +209,11 @@ class data_ser<T, typename std::enable_if<std::is_array<T>::value>::type> {
186209
};
187210

188211
template <class... Args>
189-
void gen_log_API_CP(const std::string &cp_id, size_t free_byte,
190-
size_t total_byte, float elapse_time,
212+
void gen_log_API_CP(const std::string &cp_id, std::string device_name,
213+
size_t free_byte, size_t total_byte, float elapse_time,
191214
queue_t queue, Args... args) {
192-
log.print_CP(cp_id, free_byte, total_byte, elapse_time, queue, args...);
215+
log.print_CP(cp_id, device_name, free_byte, total_byte, elapse_time, queue,
216+
args...);
193217
}
194218
} // namespace detail
195219

@@ -217,21 +241,27 @@ void gen_prolog_API_CP(const std::string &cp_id,
217241
cp_id + ":" + std::to_string(api_index[cp_id]);
218242
size_t free_byte, total_byte;
219243
#ifdef __NVCC__
244+
int device;
245+
cudaGetDevice(&device);
246+
cudaDeviceProp deviceProp;
247+
cudaGetDeviceProperties(&deviceProp, device);
248+
std::string device_name(deviceProp.name);
220249
cudaMemGetInfo(&free_byte, &total_byte);
221250
cudaEvent_t event;
222251
cudaEventCreate(&event);
223252
cudaEventRecord(event, queue);
224253
event_map[event_id] = event;
225254
#else
226255
dpct::get_current_device().get_memory_info(free_byte, total_byte);
256+
std::string device_name = dpct::get_current_device().get_info<sycl::info::device::name>();
227257
#ifdef DPCT_PROFILING_ENABLED
228258
sycl::event event = queue->ext_oneapi_submit_barrier();
229259
event_map[event_id] = event;
230260
#endif //DPCT_PROFILING_ENABLED
231261
#endif
232262

233-
detail::gen_log_API_CP(prolog_tag, free_byte,
234-
total_byte, 0.0f, queue, args...);
263+
detail::gen_log_API_CP(prolog_tag, device_name, free_byte, total_byte, 0.0f,
264+
queue, args...);
235265
}
236266

237267
/// Generate API check point epilog.
@@ -248,6 +278,11 @@ void gen_epilog_API_CP(const std::string &cp_id,
248278
size_t free_byte, total_byte;
249279
float kernel_elapsed_time = 0.0f;
250280
#ifdef __NVCC__
281+
int device;
282+
cudaGetDevice(&device);
283+
cudaDeviceProp deviceProp;
284+
cudaGetDeviceProperties(&deviceProp, device);
285+
std::string device_name(deviceProp.name);
251286
cudaMemGetInfo(&free_byte, &total_byte);
252287
cudaEvent_t event;
253288
cudaEventCreate(&event);
@@ -268,10 +303,11 @@ void gen_epilog_API_CP(const std::string &cp_id,
268303
.get_profiling_info<sycl::info::event_profiling::command_start>()) /
269304
1000000.0f;
270305
#endif //DPCT_PROFILING_ENABLED
306+
std::string device_name = dpct::get_current_device().get_info<sycl::info::device::name>();
271307
dpct::get_current_device().get_memory_info(free_byte, total_byte);
272308
#endif
273-
detail::gen_log_API_CP(
274-
epilog_tag, free_byte, total_byte, kernel_elapsed_time, queue, args...);
309+
detail::gen_log_API_CP(epilog_tag, device_name, free_byte, total_byte,
310+
kernel_elapsed_time, queue, args...);
275311
}
276312

277313
inline std::map<void *, uint32_t> &get_ptr_size_map() {

‎clang/runtime/dpct-rt/include/dpct/codepin/serialization/basic.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -107,6 +107,8 @@ class json_stringstream {
107107
void value(std::string_view value) { js.os << "\"" << value << "\""; };
108108
void value(float value) { js.os << "\"" << value << "\""; };
109109
void value(size_t value) { js.os << "\"" << value << "\""; };
110+
void value(int value) { js.os << "\"" << value << "\""; };
111+
void value(void *value) { js.os << "\"" << value << "\""; };
110112
~json_obj() {
111113
js.indent.resize(js.indent.size() - js.tab_length);
112114
js.os << js.newline;

‎clang/test/dpct/codepin/report/change_floating_point_default_tolerances/cuda.json‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:epilog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "100",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0",

‎clang/test/dpct/codepin/report/change_floating_point_default_tolerances/sycl.json‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:epilog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "100",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0",

‎clang/test/dpct/codepin/report/data_mismatch/cuda.json‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:prolog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "100",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0",
@@ -32,6 +35,9 @@
3235
},
3336
{
3437
"ID": "serial1243:epilog:0",
38+
"Device Name": "GPU",
39+
"Device ID": "0",
40+
"Stream Address": "0xe4bb30",
3541
"Free Device Memory": "200",
3642
"Total Device Memory": "500",
3743
"Elapse Time(ms)": "0.1",

‎clang/test/dpct/codepin/report/data_mismatch/sycl.json‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:prolog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "300",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0",
@@ -32,6 +35,9 @@
3235
},
3336
{
3437
"ID": "serial1243:epilog:0",
38+
"Device Name": "GPU",
39+
"Device ID": "0",
40+
"Stream Address": "0xe4bb30",
3541
"Free Device Memory": "400",
3642
"Total Device Memory": "500",
3743
"Elapse Time(ms)": "0.1",

‎clang/test/dpct/codepin/report/data_missing_in_sycl/cuda.json‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:prolog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "200",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0.0",
@@ -32,6 +35,9 @@
3235
},
3336
{
3437
"ID": "serial1243:epilog:0",
38+
"Device Name": "GPU",
39+
"Device ID": "0",
40+
"Stream Address": "0xe4bb30",
3541
"Free Device Memory": "400",
3642
"Total Device Memory": "500",
3743
"Elapse Time(ms)": "0.5",

‎clang/test/dpct/codepin/report/data_missing_in_sycl/sycl.json‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "serial123:prolog:0",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "200",
58
"Total Device Memory": "500",
69
"Elapse Time(ms)": "0.0",

‎clang/test/dpct/codepin/report/dismatch_format/cuda.json‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
[
22
{
33
"ID": "example.cu:27:3:prolog",
4+
"Device Name": "GPU",
5+
"Device ID": "0",
6+
"Stream Address": "0xe4bb30",
47
"Free Device Memory": "16374431744",
58
"Total Device Memory": "16882663424",
69
"Elapse Time(ms)": "0",
@@ -251,6 +254,9 @@
251254
},
252255
{
253256
"ID": "example.cu:27:3:epilog",
257+
"Device Name": "GPU",
258+
"Device ID": "0",
259+
"Stream Address": "0xe4bb30",
254260
"Free Device Memory": "16374431744",
255261
"Total Device Memory": "16882663424",
256262
"Elapse Time(ms)": "0.21312",

0 commit comments

Comments
 (0)