Skip to content

Commit 986e89a

Browse files
RbiessyRbiessy
committed
sycl: Add more debug prints
1 parent 725f23f commit 986e89a

File tree

16 files changed

+124
-105
lines changed

16 files changed

+124
-105
lines changed

‎ggml/src/ggml-sycl/binbcast.cpp‎

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -208,32 +208,27 @@ inline void ggml_sycl_op_repeat(ggml_backend_sycl_context & ctx, ggml_tensor *ds
208208

209209

210210
void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
211-
GGML_SYCL_DEBUG("call %s\n", __func__);
211+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
212212
ggml_sycl_op_add(ctx, dst);
213-
GGML_SYCL_DEBUG("call %s done\n", __func__);
214213
}
215214

216215
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
217-
GGML_SYCL_DEBUG("call %s\n", __func__);
216+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
218217
ggml_sycl_op_sub(ctx, dst);
219-
GGML_SYCL_DEBUG("call %s done\n", __func__);
220218
}
221219

222220
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
223-
GGML_SYCL_DEBUG("call %s\n", __func__);
221+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
224222
ggml_sycl_op_mul(ctx, dst);
225-
GGML_SYCL_DEBUG("call %s done\n", __func__);
226223
}
227224

228225
void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
229-
GGML_SYCL_DEBUG("call %s\n", __func__);
226+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
230227
ggml_sycl_op_div(ctx, dst);
231-
GGML_SYCL_DEBUG("call %s done\n", __func__);
232228
}
233229

234230
void ggml_sycl_repeat(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
235-
GGML_SYCL_DEBUG("call %s\n", __func__);
231+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
236232
ggml_sycl_op_repeat(ctx, dst);
237-
GGML_SYCL_DEBUG("call %s done\n", __func__);
238233
}
239234

‎ggml/src/ggml-sycl/common.hpp‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515

1616
#include <fstream>
1717
#include <iostream>
18+
#include <string>
1819

1920
#include "dpct/helper.hpp"
2021
#include "ggml-sycl.h"
@@ -490,4 +491,59 @@ constexpr size_t ceil_div(const size_t m, const size_t n) {
490491
}
491492

492493
bool gpu_has_xmx(sycl::device &dev);
494+
495+
template <int N, class T>
496+
void debug_print_array(const std::string& prefix, const T array[N]) {
497+
std::stringstream ss;
498+
ss << prefix << "=[";
499+
for (std::size_t i = 0; i < N - 1; ++i) {
500+
ss << array[i] << ", ";
501+
}
502+
if constexpr (N > 0) {
503+
ss << array[N - 1];
504+
}
505+
ss << "]";
506+
GGML_SYCL_DEBUG("%s", ss.str().c_str());
507+
}
508+
509+
inline void debug_print_tensor(const std::string& prefix, const ggml_tensor* tensor) {
510+
GGML_SYCL_DEBUG("%s=", prefix.c_str());
511+
if (tensor) {
512+
GGML_SYCL_DEBUG("'%s':type=%s", tensor->name, ggml_type_name(tensor->type));
513+
debug_print_array<GGML_MAX_DIMS>(";ne", tensor->ne);
514+
debug_print_array<GGML_MAX_DIMS>(";nb", tensor->nb);
515+
if (!ggml_is_contiguous(tensor)) {
516+
GGML_SYCL_DEBUG(";strided");
517+
}
518+
if (ggml_is_permuted(tensor)) {
519+
GGML_SYCL_DEBUG(";permuted");
520+
}
521+
} else {
522+
GGML_SYCL_DEBUG("nullptr");
523+
}
524+
}
525+
526+
struct scope_op_debug_print {
527+
scope_op_debug_print(const std::string& func, const ggml_tensor* dst, std::size_t num_src, const std::string& suffix = "") : func(func) {
528+
if (!g_ggml_sycl_debug) {
529+
return;
530+
}
531+
GGML_SYCL_DEBUG("call %s:", func.c_str());
532+
debug_print_tensor(" dst", dst);
533+
if (dst) {
534+
for (std::size_t i = 0; i < num_src; ++i) {
535+
debug_print_tensor("\tsrc" + std::to_string(i), dst->src[i]);
536+
}
537+
}
538+
GGML_SYCL_DEBUG("%s\n", suffix.c_str());
539+
}
540+
541+
~scope_op_debug_print() {
542+
GGML_SYCL_DEBUG("call %s done\n", func.c_str());
543+
}
544+
545+
private:
546+
std::string func;
547+
};
548+
493549
#endif // GGML_SYCL_COMMON_HPP

‎ggml/src/ggml-sycl/concat.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -159,6 +159,7 @@ static void concat_f32_sycl_non_cont(
159159
}
160160

161161
void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
162+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
162163
const ggml_tensor *src0 = dst->src[0];
163164
const ggml_tensor *src1 = dst->src[1];
164165
queue_ptr stream = ctx.stream();

‎ggml/src/ggml-sycl/conv.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,7 @@ static void conv_transpose_1d_f32_f32_sycl(
7272
}
7373

7474
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
75+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
7576
const ggml_tensor *src0 = dst->src[0];
7677
const ggml_tensor *src1 = dst->src[1];
7778
const float * src0_d = (const float *)src0->data;

‎ggml/src/ggml-sycl/cpy.cpp‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -616,6 +616,8 @@ static void ggml_cpy_i32_i32_sycl(const char * cx, char * cdst, const int ne, co
616616
}
617617

618618
void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1) try {
619+
// Unlike other operators ggml_sycl_cpy takes 2 distinct tensors instead of a dst ggml_tensor and rely on its src field
620+
scope_op_debug_print scope_dbg_print(__func__, src1, /*num_src=*/0, std::string(" src0 type=") + ggml_type_name(src0->type));
619621
const int64_t ne = ggml_nelements(src0);
620622
GGML_ASSERT(ne == ggml_nelements(src1));
621623

@@ -629,8 +631,6 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
629631

630632
char * src0_ddc = (char *) src0->data;
631633
char * src1_ddc = (char *) src1->data;
632-
GGML_SYCL_DEBUG("[SYCL] %s: Tensor supplied: %s to %s\n", __func__, ggml_type_name(src0->type),
633-
ggml_type_name(src1->type));
634634

635635
if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
636636
ggml_cpy_f32_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
@@ -694,8 +694,6 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
694694
}
695695

696696
void ggml_sycl_dup(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
697-
// TODO: why do we pass dst as src1 here?
698-
GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
697+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
699698
ggml_sycl_cpy(ctx, dst->src[0], dst);
700-
GGML_SYCL_DEBUG("[SYCL] call %s done\n", __func__);
701699
}

‎ggml/src/ggml-sycl/dmmv.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1092,6 +1092,7 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
10921092
src0->type == GGML_TYPE_Q8_0 || src0->type == GGML_TYPE_F16;
10931093

10941094
if (src1_convert_f16) {
1095+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2, " : converting src1 to fp16");
10951096
src1_dfloat = src1_dfloat_a.alloc(ne00);
10961097
const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type, dst);
10971098
GGML_ASSERT(to_fp16_sycl != nullptr);

‎ggml/src/ggml-sycl/element_wise.cpp‎

Lines changed: 24 additions & 49 deletions
Original file line numberDiff line numberDiff line change
@@ -1391,146 +1391,121 @@ inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor *dst)
13911391

13921392

13931393
void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1394-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1394+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
13951395
ggml_sycl_op_sqrt(ctx, dst);
1396-
GGML_SYCL_DEBUG("call %s done\n", __func__);
13971396
}
13981397

13991398
void ggml_sycl_sin(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1400-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1399+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14011400
ggml_sycl_op_sin(ctx, dst);
1402-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14031401
}
14041402

14051403
void ggml_sycl_cos(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1406-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1404+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14071405
ggml_sycl_op_cos(ctx, dst);
1408-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14091406
}
14101407

14111408
void ggml_sycl_acc(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1412-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1409+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
14131410
ggml_sycl_op_acc(ctx, dst);
1414-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14151411
}
14161412

14171413
void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1418-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1414+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14191415
ggml_sycl_op_gelu(ctx, dst);
1420-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14211416
}
14221417

14231418
void ggml_sycl_silu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1424-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1419+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14251420
ggml_sycl_op_silu(ctx, dst);
1426-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14271421
}
14281422

14291423
void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1430-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1424+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14311425
ggml_sycl_op_gelu_quick(ctx, dst);
1432-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14331426
}
14341427

14351428
void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1436-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1429+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14371430
ggml_sycl_op_tanh(ctx, dst);
1438-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14391431
}
14401432

14411433
void ggml_sycl_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1442-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1434+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14431435
ggml_sycl_op_relu(ctx, dst);
1444-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14451436
}
14461437

14471438
void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1448-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1439+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14491440
ggml_sycl_op_sigmoid(ctx, dst);
1450-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14511441
}
14521442

14531443
void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1454-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1444+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14551445
ggml_sycl_op_hardsigmoid(ctx, dst);
1456-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14571446
}
14581447

14591448
void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1460-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1449+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14611450
ggml_sycl_op_hardswish(ctx, dst);
1462-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14631451
}
14641452

1465-
14661453
void ggml_sycl_exp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1467-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1454+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14681455
ggml_sycl_op_exp(ctx, dst);
1469-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14701456
}
14711457

14721458
void ggml_sycl_log(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1473-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1459+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14741460
ggml_sycl_op_log(ctx, dst);
1475-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14761461
}
14771462

14781463
void ggml_sycl_neg(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1479-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1464+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14801465
ggml_sycl_op_neg(ctx, dst);
1481-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14821466
}
14831467

14841468
void ggml_sycl_step(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1485-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1469+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14861470
ggml_sycl_op_step(ctx, dst);
1487-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14881471
}
14891472

14901473
void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1491-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1474+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14921475
ggml_sycl_op_leaky_relu(ctx, dst);
1493-
GGML_SYCL_DEBUG("call %s done\n", __func__);
14941476
}
14951477

14961478
void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1497-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1479+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
14981480
ggml_sycl_op_sqr(ctx, dst);
1499-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15001481
}
15011482

15021483
void ggml_sycl_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1503-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1484+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15041485
ggml_sycl_op_upscale(ctx, dst);
1505-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15061486
}
15071487

15081488
void ggml_sycl_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1509-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1489+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15101490
ggml_sycl_op_pad(ctx, dst);
1511-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15121491
}
15131492

15141493
void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1515-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1494+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15161495
ggml_sycl_op_clamp(ctx, dst);
1517-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15181496
}
15191497

15201498
void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1521-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1499+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15221500
ggml_sycl_op_sgn(ctx, dst);
1523-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15241501
}
15251502

15261503
void ggml_sycl_abs(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1527-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1504+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15281505
ggml_sycl_op_abs(ctx, dst);
1529-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15301506
}
15311507

15321508
void ggml_sycl_elu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1533-
GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
1509+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
15341510
ggml_sycl_op_elu(ctx, dst);
1535-
GGML_SYCL_DEBUG("call %s done\n", __func__);
15361511
}

‎ggml/src/ggml-sycl/getrows.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -258,7 +258,6 @@ static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tens
258258
}
259259

260260
void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
261-
262261
GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I32);
263262
GGML_ASSERT(dst->type == GGML_TYPE_F32);
264263

0 commit comments

Comments
 (0)