Skip to content

Commit 1ae14e0

Browse files
RbiessyRbiessy
committed
clang-format
1 parent 45a027a commit 1ae14e0

File tree

6 files changed

+82
-77
lines changed

6 files changed

+82
-77
lines changed

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

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -47,11 +47,11 @@ extern int g_ggml_sycl_prioritize_dmmv;
4747

4848
#if defined(__clang__) && __has_builtin(__builtin_expect)
4949
// Hint the optimizer to pipeline the more likely following instruction in branches
50-
#define LIKELY(expr) __builtin_expect(expr, true)
51-
#define UNLIKELY(expr) __builtin_expect(expr, false)
50+
# define LIKELY(expr) __builtin_expect(expr, true)
51+
# define UNLIKELY(expr) __builtin_expect(expr, false)
5252
#else
53-
#define LIKELY(expr) (expr)
54-
#define UNLIKELY(expr) (expr)
53+
# define LIKELY(expr) (expr)
54+
# define UNLIKELY(expr) (expr)
5555
#endif
5656

5757
#define GGML_SYCL_DEBUG(...) \
@@ -540,23 +540,23 @@ inline void debug_print_tensor(const std::string & prefix, const ggml_tensor * t
540540
}
541541

542542
struct scope_op_debug_print {
543-
scope_op_debug_print(const std::string& func, const ggml_tensor* dst, std::size_t num_src, const std::string& suffix = "") : func(func) {
544-
if (LIKELY(!g_ggml_sycl_debug)) {
545-
return;
546-
}
547-
GGML_SYCL_DEBUG("[SYCL][OP] call %s:", func.c_str());
548-
debug_print_tensor(" dst", dst);
549-
if (dst) {
550-
for (std::size_t i = 0; i < num_src; ++i) {
551-
debug_print_tensor("\tsrc" + std::to_string(i), dst->src[i]);
543+
scope_op_debug_print(const std::string & func, const ggml_tensor * dst, std::size_t num_src,
544+
const std::string & suffix = "") :
545+
func(func) {
546+
if (LIKELY(!g_ggml_sycl_debug)) {
547+
return;
552548
}
549+
GGML_SYCL_DEBUG("[SYCL][OP] call %s:", func.c_str());
550+
debug_print_tensor(" dst", dst);
551+
if (dst) {
552+
for (std::size_t i = 0; i < num_src; ++i) {
553+
debug_print_tensor("\tsrc" + std::to_string(i), dst->src[i]);
554+
}
555+
}
556+
GGML_SYCL_DEBUG("%s\n", suffix.c_str());
553557
}
554-
GGML_SYCL_DEBUG("%s\n", suffix.c_str());
555-
}
556558

557-
~scope_op_debug_print() {
558-
GGML_SYCL_DEBUG("[SYCL][OP] call %s done\n", func.c_str());
559-
}
559+
~scope_op_debug_print() { GGML_SYCL_DEBUG("[SYCL][OP] call %s done\n", func.c_str()); }
560560

561561
private:
562562
std::string func;

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

Lines changed: 31 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -159,40 +159,37 @@ 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);
163-
const ggml_tensor *src0 = dst->src[0];
164-
const ggml_tensor *src1 = dst->src[1];
165-
queue_ptr stream = ctx.stream();
166-
167-
const int32_t dim = ((int32_t *)dst->op_params)[0];
168-
169-
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
170-
const float *src0_d = (const float *)src0->data;
171-
const float *src1_d = (const float *)src1->data;
172-
173-
float *dst_d = (float *)dst->data;
174-
175-
if (dim != 3) {
176-
for (int i3 = 0; i3 < dst->ne[3]; i3++) {
177-
concat_f32_sycl(
178-
src0_d + i3 * (src0->nb[3] / 4), src1_d + i3 * (src1->nb[3] / 4),
179-
dst_d + i3 * (dst->nb[3] / 4), src0->ne[0], src0->ne[1],
180-
src0->ne[2], dst->ne[0], dst->ne[1], dst->ne[2], dim, stream);
181-
}
162+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
163+
const ggml_tensor * src0 = dst->src[0];
164+
const ggml_tensor * src1 = dst->src[1];
165+
queue_ptr stream = ctx.stream();
166+
167+
const int32_t dim = ((int32_t *) dst->op_params)[0];
168+
169+
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
170+
const float * src0_d = (const float *) src0->data;
171+
const float * src1_d = (const float *) src1->data;
172+
173+
float * dst_d = (float *) dst->data;
174+
175+
if (dim != 3) {
176+
for (int i3 = 0; i3 < dst->ne[3]; i3++) {
177+
concat_f32_sycl(src0_d + i3 * (src0->nb[3] / 4), src1_d + i3 * (src1->nb[3] / 4),
178+
dst_d + i3 * (dst->nb[3] / 4), src0->ne[0], src0->ne[1], src0->ne[2], dst->ne[0],
179+
dst->ne[1], dst->ne[2], dim, stream);
180+
}
181+
} else {
182+
const size_t size0 = ggml_nbytes(src0);
183+
const size_t size1 = ggml_nbytes(src1);
184+
185+
SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0).wait()));
186+
SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d + size0 / 4, src1_d, size1).wait()));
187+
}
182188
} else {
183-
const size_t size0 = ggml_nbytes(src0);
184-
const size_t size1 = ggml_nbytes(src1);
185-
186-
SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0).wait()));
187-
SYCL_CHECK(CHECK_TRY_ERROR(
188-
stream->memcpy(dst_d + size0 / 4, src1_d, size1).wait()));
189+
concat_f32_sycl_non_cont(stream, (const char *) src0->data, (const char *) src1->data, (char *) dst->data,
190+
src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], src0->nb[0], src0->nb[1],
191+
src0->nb[2], src0->nb[3], src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
192+
src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3], dst->ne[0], dst->ne[1], dst->ne[2],
193+
dst->ne[3], dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
189194
}
190-
} else
191-
concat_f32_sycl_non_cont(
192-
stream, (const char *)src0->data, (const char *)src1->data,
193-
(char *)dst->data, src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
194-
src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], src1->ne[0],
195-
src1->ne[1], src1->ne[2], src1->ne[3], src1->nb[0], src1->nb[1],
196-
src1->nb[2], src1->nb[3], dst->ne[0], dst->ne[1], dst->ne[2],
197-
dst->ne[3], dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
198195
}

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -617,7 +617,8 @@ static void ggml_cpy_i32_i32_sycl(const char * cx, char * cdst, const int ne, co
617617

618618
void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1) try {
619619
// 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));
620+
scope_op_debug_print scope_dbg_print(__func__, src1, /*num_src=*/0,
621+
std::string(" src0 type=") + ggml_type_name(src0->type));
621622
const int64_t ne = ggml_nelements(src0);
622623
GGML_ASSERT(ne == ggml_nelements(src1));
623624

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -257,7 +257,7 @@ static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tens
257257
GGML_UNUSED(ctx);
258258
}
259259

260-
void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
260+
void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
261261
GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I32);
262262
GGML_ASSERT(dst->type == GGML_TYPE_F32);
263263

@@ -307,4 +307,3 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
307307
GGML_ABORT("fatal error");
308308
}
309309
}
310-

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

Lines changed: 29 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -503,7 +503,7 @@ ggml_backend_sycl_buffer_cpy_tensor(ggml_backend_buffer_t buffer,
503503
static void ggml_backend_sycl_buffer_clear(ggml_backend_buffer_t buffer,
504504
uint8_t value) try {
505505
GGML_SYCL_DEBUG("[SYCL] call %s: size=%zu\n", __func__, buffer->size);
506-
ggml_backend_sycl_buffer_context * ctx = ( ggml_backend_sycl_buffer_context *)buffer->context;
506+
ggml_backend_sycl_buffer_context * ctx = (ggml_backend_sycl_buffer_context *) buffer->context;
507507

508508
ggml_sycl_set_device(ctx->device);
509509
queue_ptr stream = ctx->stream;
@@ -2036,12 +2036,12 @@ inline void ggml_sycl_op_mul_mat_sycl(
20362036
#else
20372037
bool use_fp16 = false;
20382038
#endif
2039-
if ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) &&
2040-
use_fp16 && ggml_is_contiguous(src0) && row_diff == src0->ne[1] &&
2041-
dst->op_params[0] == GGML_PREC_DEFAULT) {
2039+
if ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && use_fp16 && ggml_is_contiguous(src0) &&
2040+
row_diff == src0->ne[1] && dst->op_params[0] == GGML_PREC_DEFAULT) {
20422041
ggml_sycl_pool_alloc<sycl::half> src0_as_f16(ctx.pool());
20432042
if (src0->type != GGML_TYPE_F16) {
2044-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_sycl", dst, /*num_src=*/2, " : converting src0 to fp16");
2043+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_sycl", dst, /*num_src=*/2,
2044+
" : converting src0 to fp16");
20452045
const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src0->type, dst);
20462046
GGML_ASSERT(to_fp16_sycl != nullptr);
20472047
size_t ne = row_diff*ne00;
@@ -2054,7 +2054,8 @@ inline void ggml_sycl_op_mul_mat_sycl(
20542054

20552055
ggml_sycl_pool_alloc<sycl::half> src1_as_f16(ctx.pool());
20562056
if (src1->type != GGML_TYPE_F16) {
2057-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_sycl", dst, /*num_src=*/2, " : converting src1 to fp16");
2057+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_sycl", dst, /*num_src=*/2,
2058+
" : converting src1 to fp16");
20582059
const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type, dst);
20592060
GGML_ASSERT(to_fp16_sycl != nullptr);
20602061
size_t ne = src1_ncols*ne10;
@@ -2071,7 +2072,8 @@ inline void ggml_sycl_op_mul_mat_sycl(
20712072
DnnlGemmWrapper::row_gemm(ctx, src1_ncols, row_diff, ne10, src1_ptr,
20722073
DnnlGemmWrapper::to_dt<sycl::half>(), src0_ptr, DnnlGemmWrapper::to_dt<sycl::half>(),
20732074
dst_f16.get(), DnnlGemmWrapper::to_dt<sycl::half>(), stream);
2074-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2, " : converting dst to fp32");
2075+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2,
2076+
" : converting dst to fp32");
20752077
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
20762078
to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff* src1_ncols, stream);
20772079
}
@@ -2087,23 +2089,25 @@ inline void ggml_sycl_op_mul_mat_sycl(
20872089
src1_ptr, dpct::library_data_t::real_half, ne10, &beta_f16,
20882090
dst_f16.get(), dpct::library_data_t::real_half, ldc,
20892091
dpct::library_data_t::real_half)));
2090-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2, " : converting dst to fp32");
2092+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2,
2093+
" : converting dst to fp32");
20912094
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
20922095
to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff*src1_ncols, stream);
20932096
}
2094-
}
2095-
else {
2097+
} else {
20962098
ggml_sycl_pool_alloc<float> src0_ddq_as_f32(ctx.pool());
20972099
ggml_sycl_pool_alloc<float> src1_ddq_as_f32(ctx.pool());
20982100
if (src0->type != GGML_TYPE_F32) {
2099-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2, " : converting src0 to fp32");
2101+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2,
2102+
" : converting src0 to fp32");
21002103
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(src0->type, dst);
21012104
GGML_ASSERT(to_fp32_sycl != nullptr);
21022105
src0_ddq_as_f32.alloc(row_diff*ne00);
21032106
to_fp32_sycl(src0_dd_i, src0_ddq_as_f32.get(), row_diff*ne00, stream);
21042107
}
21052108
if (src1->type != GGML_TYPE_F32) {
2106-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2, " : converting src1 to fp32");
2109+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp32_sycl", dst, /*num_src=*/2,
2110+
" : converting src1 to fp32");
21072111
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(src1->type, dst);
21082112
GGML_ASSERT(to_fp32_sycl != nullptr);
21092113
src1_ddq_as_f32.alloc(src1_ncols*ne10);
@@ -2139,7 +2143,7 @@ catch (sycl::exception const &exc) {
21392143
std::exit(1);
21402144
}
21412145

2142-
static void ggml_sycl_op_pool2d(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
2146+
static void ggml_sycl_op_pool2d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
21432147
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
21442148
GGML_ASSERT( dst->type == GGML_TYPE_F32);
21452149
dpct::queue_ptr main_stream = ctx.stream();
@@ -2191,7 +2195,7 @@ inline void ggml_sycl_op_sum(ggml_backend_sycl_context & ctx, ggml_tensor *dst)
21912195
sum_rows_f32_sycl(src0_dd, dst_dd, ne, 1, main_stream);
21922196
}
21932197

2194-
inline void ggml_sycl_op_sum_rows(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
2198+
inline void ggml_sycl_op_sum_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
21952199
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
21962200
GGML_ASSERT( dst->type == GGML_TYPE_F32);
21972201
dpct::queue_ptr main_stream = ctx.stream();
@@ -2222,7 +2226,7 @@ inline void ggml_sycl_op_argsort(ggml_backend_sycl_context & ctx, ggml_tensor *
22222226
argsort_f32_i32_sycl(src0_dd, (int *) dst_dd, ncols, nrows, order, main_stream);
22232227
}
22242228

2225-
inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
2229+
inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
22262230
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
22272231
GGML_ASSERT( dst->type == GGML_TYPE_I32);
22282232

@@ -2237,7 +2241,7 @@ inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor *ds
22372241
argmax_f32_i32_sycl(src0_dd, dst_dd, ncols, nrows, main_stream);
22382242
}
22392243

2240-
inline void ggml_sycl_op_diag_mask_inf(ggml_backend_sycl_context & ctx,ggml_tensor *dst) {
2244+
inline void ggml_sycl_op_diag_mask_inf(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
22412245
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
22422246
GGML_ASSERT( dst->type == GGML_TYPE_F32);
22432247
dpct::queue_ptr main_stream = ctx.stream();
@@ -2254,7 +2258,7 @@ inline void ggml_sycl_op_diag_mask_inf(ggml_backend_sycl_context & ctx,ggml_tens
22542258
diag_mask_inf_f32_sycl(src0_dd, dst_dd, ne00, nrows0, ne01, n_past, main_stream);
22552259
}
22562260

2257-
inline void ggml_sycl_op_scale(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
2261+
inline void ggml_sycl_op_scale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
22582262
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
22592263
GGML_ASSERT( dst->type == GGML_TYPE_F32);
22602264
dpct::queue_ptr main_stream = ctx.stream();
@@ -2441,7 +2445,8 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten
24412445
dev[i].src1_ddq = dev[i].src1_ddq_alloc.alloc(ctx.pool(i), nrows1*src1_padded_col_size*q8_1_ts/q8_1_bs);
24422446

24432447
if (src1_on_device && src1_is_contiguous) {
2444-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/quantize_row_q8_1_sycl", dst, /*num_src=*/2, " : converting src1 to Q8_1");
2448+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/quantize_row_q8_1_sycl", dst,
2449+
/*num_src=*/2, " : converting src1 to Q8_1");
24452450
quantize_row_q8_1_sycl(dev[i].src1_ddf, dev[i].src1_ddq, ne10, nrows1, src1_padded_col_size, stream);
24462451
/*
24472452
DPCT1010:90: SYCL uses exceptions to report errors and does not
@@ -2546,7 +2551,8 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten
25462551
}
25472552

25482553
if (convert_src1_to_q8_1 && !src1_is_contiguous) {
2549-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/quantize_row_q8_1_sycl", dst, /*num_src=*/2, " : converting src1 to Q8_1");
2554+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/quantize_row_q8_1_sycl", dst,
2555+
/*num_src=*/2, " : converting src1 to Q8_1");
25502556
quantize_row_q8_1_sycl(src1_ddf_i, src1_ddq_i, ne10, src1_ncols, src1_padded_col_size, stream);
25512557
/*
25522558
DPCT1010:92: SYCL uses exceptions to report errors and does
@@ -2790,7 +2796,8 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
27902796

27912797
// convert src1 to fp16
27922798
if (src1->type != GGML_TYPE_F16) {
2793-
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_nc_sycl", dst, /*num_src=*/2, " : converting src1 to fp16");
2799+
scope_op_debug_print scope_dbg_print(std::string(__func__) + "/to_fp16_nc_sycl", dst, /*num_src=*/2,
2800+
" : converting src1 to fp16");
27942801
const to_fp16_nc_sycl_t to_fp16_nc_sycl = get_to_fp16_nc_sycl(src1->type);
27952802
GGML_ASSERT(to_fp16_nc_sycl != nullptr);
27962803
const int64_t ne_src1 = ggml_nelements(src1);
@@ -3787,7 +3794,8 @@ static bool ggml_backend_sycl_cpy_tensor_async(ggml_backend_t backend,
37873794
const ggml_tensor *src,
37883795
ggml_tensor *dst) try {
37893796
ggml_backend_sycl_context * sycl_ctx = (ggml_backend_sycl_context *)backend->context;
3790-
bool is_cpy_supported = dst->buffer->buft == ggml_backend_sycl_buffer_type(sycl_ctx->device) && ggml_backend_buffer_is_sycl(src->buffer);
3797+
bool is_cpy_supported = dst->buffer->buft == ggml_backend_sycl_buffer_type(sycl_ctx->device) &&
3798+
ggml_backend_buffer_is_sycl(src->buffer);
37913799
GGML_SYCL_DEBUG("[SYCL] call %s", __func__);
37923800
debug_print_tensor(": dst=", dst);
37933801
debug_print_tensor(" src=", src);

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,7 @@ static void timestep_embedding_f32_sycl(
5757

5858
void ggml_sycl_op_timestep_embedding(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
5959
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
60-
const ggml_tensor *src0 = dst->src[0];
60+
const ggml_tensor * src0 = dst->src[0];
6161
const float * src0_d = (const float *)src0->data;
6262
float * dst_d = (float *)dst->data;
6363
dpct::queue_ptr stream = ctx.stream();

0 commit comments

Comments
 (0)