Skip to content

Commit dd69db2

Browse files
arthwarthw
authored
sycl : support MUL_MAT and OUT_PROD with Q1_0 (ggml-org#24721)
1 parent 6ec59dd commit dd69db2

7 files changed

Lines changed: 247 additions & 18 deletions

File tree

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -642,6 +642,8 @@ static void convert_unary_sycl(const void * vx, dst_t * y, const int64_t k, dpct
642642

643643
to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
644644
switch (type) {
645+
case GGML_TYPE_Q1_0:
646+
return dequantize_block_sycl<QK1_0, QR1_0, dequantize_q1_0>;
645647
case GGML_TYPE_Q4_0:
646648
if (dst->src[0]->extra &&
647649
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
@@ -724,6 +726,8 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
724726

725727
to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
726728
switch (type) {
729+
case GGML_TYPE_Q1_0:
730+
return dequantize_block_sycl<QK1_0, QR1_0, dequantize_q1_0>;
727731
case GGML_TYPE_Q4_0:
728732
if (dst->src[0]->extra &&
729733
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
@@ -830,6 +834,8 @@ to_fp16_nc_sycl_t ggml_get_to_fp16_nc_sycl(ggml_type type) {
830834
case GGML_TYPE_BF16:
831835
return convert_unary_nc_sycl<sycl::ext::oneapi::bfloat16>;
832836
#endif
837+
case GGML_TYPE_Q1_0:
838+
return dequantize_block_nc_sycl<QK1_0, QR1_0, dequantize_q1_0>;
833839
case GGML_TYPE_Q4_0:
834840
return dequantize_block_nc_sycl<QK4_0, QR4_0, dequantize_q4_0>;
835841
case GGML_TYPE_Q4_1:

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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -70,6 +70,21 @@ static __dpct_inline__ void dequantize_q4_0_reorder(const void *d_ptr, const int
7070
#endif // GGML_SYCL_F16
7171
}
7272

73+
static __dpct_inline__ void dequantize_q1_0_reorder(const void *d_ptr, const int64_t ib, const void *qs,
74+
const int iqs, dfloat2 &v) {
75+
// Q1_0 reorder layout: scale values followed by quantized bits
76+
const dfloat d = (const dfloat)*((const sycl::half*)d_ptr+ib);
77+
78+
const int bit_index_0 = iqs + 0;
79+
const int bit_index_1 = iqs + 1;
80+
81+
const int bit_0 = (*((const uint8_t *)qs + bit_index_0 / 8) >> (bit_index_0 % 8)) & 1;
82+
const int bit_1 = (*((const uint8_t *)qs + bit_index_1 / 8) >> (bit_index_1 % 8)) & 1;
83+
84+
v.x() = (2 * bit_0 - 1) * d;
85+
v.y() = (2 * bit_1 - 1) * d;
86+
}
87+
7388
static __dpct_inline__ void dequantize_q4_1(const void *vx, const int64_t ib,
7489
const int iqs, dfloat2 &v) {
7590
const block_q4_1 * x = (const block_q4_1 *) vx;

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

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1423,6 +1423,50 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
14231423
}
14241424
}
14251425

1426+
static void dequantize_mul_mat_vec_q1_0_sycl_reorder(const void *vx, const dfloat *y,
1427+
float *dst, const int ncols,
1428+
const int nrows,
1429+
dpct::queue_ptr stream) {
1430+
GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
1431+
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
1432+
// the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
1433+
const sycl::range<3> block_nums(1, 1, block_num_y);
1434+
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
1435+
{
1436+
dpct::has_capability_or_fail(stream->get_device(),
1437+
{sycl::aspect::fp16});
1438+
1439+
stream->parallel_for(
1440+
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1441+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
1442+
dequantize_mul_mat_vec_reorder<QK1_0, QR1_0, dequantize_q1_0_reorder>(
1443+
vx, y, dst, ncols, nrows, item_ct1);
1444+
});
1445+
}
1446+
}
1447+
1448+
static void dequantize_mul_mat_vec_q1_0_sycl(const void *vx, const dfloat *y,
1449+
float *dst, const int ncols,
1450+
const int nrows,
1451+
dpct::queue_ptr stream) {
1452+
GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
1453+
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
1454+
// the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
1455+
const sycl::range<3> block_nums(1, 1, block_num_y);
1456+
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
1457+
{
1458+
dpct::has_capability_or_fail(stream->get_device(),
1459+
{sycl::aspect::fp16});
1460+
1461+
stream->parallel_for(
1462+
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1463+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
1464+
dequantize_mul_mat_vec<QK1_0, QR1_0, dequantize_q1_0>(
1465+
vx, y, dst, ncols, nrows, item_ct1);
1466+
});
1467+
}
1468+
}
1469+
14261470
static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
14271471
float *dst, const int ncols,
14281472
const int nrows,
@@ -1759,6 +1803,7 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
17591803
sycl::half *src1_dfloat = nullptr; // dfloat == half
17601804

17611805
bool src1_convert_f16 =
1806+
src0->type == GGML_TYPE_Q1_0 ||
17621807
src0->type == GGML_TYPE_Q4_0 || src0->type == GGML_TYPE_Q4_1 ||
17631808
src0->type == GGML_TYPE_Q5_0 || src0->type == GGML_TYPE_Q5_1 ||
17641809
src0->type == GGML_TYPE_Q8_0 || src0->type == GGML_TYPE_F16 ||
@@ -1777,6 +1822,14 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
17771822
#endif // GGML_SYCL_F16
17781823

17791824
switch (src0->type) {
1825+
case GGML_TYPE_Q1_0:
1826+
if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
1827+
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
1828+
dequantize_mul_mat_vec_q1_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
1829+
} else {
1830+
dequantize_mul_mat_vec_q1_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
1831+
}
1832+
break;
17801833
case GGML_TYPE_Q4_0:
17811834
if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
17821835
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {

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

Lines changed: 19 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -976,6 +976,7 @@ static int64_t get_row_rounding(ggml_type type, const std::array<float, GGML_SYC
976976
}
977977

978978
switch(type) {
979+
case GGML_TYPE_Q1_0:
979980
case GGML_TYPE_Q4_0:
980981
case GGML_TYPE_Q4_1:
981982
return max_compute_capability >= VER_GEN9 ? 128 : 64;
@@ -3507,6 +3508,7 @@ inline bool ggml_sycl_supports_mmq(enum ggml_type type) {
35073508

35083509
inline bool ggml_sycl_supports_reorder_mul_mat_sycl(enum ggml_type type) {
35093510
switch (type) {
3511+
case GGML_TYPE_Q1_0:
35103512
case GGML_TYPE_Q4_0:
35113513
case GGML_TYPE_Q8_0:
35123514
return true;
@@ -3522,6 +3524,7 @@ inline bool ggml_sycl_supports_reorder_mul_mat_sycl(enum ggml_type type) {
35223524

35233525
inline bool ggml_sycl_supports_reorder_dmmv(enum ggml_type type) {
35243526
switch (type) {
3527+
case GGML_TYPE_Q1_0:
35253528
case GGML_TYPE_Q4_0:
35263529
case GGML_TYPE_Q8_0:
35273530
return true;
@@ -3532,6 +3535,7 @@ inline bool ggml_sycl_supports_reorder_dmmv(enum ggml_type type) {
35323535

35333536
inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) {
35343537
switch (type) {
3538+
case GGML_TYPE_Q1_0:
35353539
case GGML_TYPE_Q4_0:
35363540
case GGML_TYPE_Q8_0:
35373541
case GGML_TYPE_Q3_K:
@@ -3546,6 +3550,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) {
35463550

35473551
static bool ggml_sycl_supports_dmmv(enum ggml_type type) {
35483552
switch (type) {
3553+
case GGML_TYPE_Q1_0:
35493554
case GGML_TYPE_Q4_0:
35503555
case GGML_TYPE_Q4_1:
35513556
case GGML_TYPE_Q5_0:
@@ -5385,7 +5390,7 @@ static ggml_backend_buffer_t ggml_backend_sycl_device_buffer_from_host_ptr(ggml_
53855390
return nullptr;
53865391
}
53875392

5388-
static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const ggml_tensor * op) {
5393+
static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const ggml_tensor * op) {
53895394
ggml_backend_sycl_device_context *sycl_ctx =
53905395
(ggml_backend_sycl_device_context *)dev->context;
53915396
int device = sycl_ctx->device;
@@ -5450,19 +5455,12 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
54505455
struct ggml_tensor * a = op->src[0];
54515456
struct ggml_tensor * b = op->src[1];
54525457

5453-
// disable Q1_0 until implementation
5454-
if (a->type == GGML_TYPE_Q1_0 || b->type == GGML_TYPE_Q1_0) {
5455-
return false;
5456-
}
5457-
54585458
if (a->ne[3] != b->ne[3]) {
54595459
return false;
54605460
}
54615461

54625462
ggml_type src0_type = op->src[0]->type;
54635463

5464-
5465-
54665464
// TODO: The configuration below needs more work to be supported with oneDNN
54675465
if (ggml_is_permuted(a) && !ggml_is_contiguous(a) &&
54685466
a->ne[2] > 1 && a->ne[3] > 1 && src0_type == GGML_TYPE_F16) {
@@ -5472,12 +5470,17 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
54725470
// TODO: This specific configuration can fail with oneDNN and needs more debugging
54735471
if (!ggml_is_permuted(a) && ggml_is_permuted(b) && b->ne[2] > 1 && b->ne[3] > 1 &&
54745472
a->ne[0] > 128 && a->ne[2] == 1 && src0_type == GGML_TYPE_F16) {
5473+
printf("zjy 2\n");
54755474
return false;
54765475
}
54775476
return true;
54785477
}
54795478
case GGML_OP_OUT_PROD:
5480-
return op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->ne[2] == 1 && op->ne[3] == 1;
5479+
return op->type == GGML_TYPE_F32 &&
5480+
(op->src[0]->type == GGML_TYPE_F32 ||
5481+
(op->src[0]->type == GGML_TYPE_Q1_0 && op->src[0]->ne[2] == op->src[1]->ne[2] &&
5482+
op->src[0]->ne[3] == op->src[1]->ne[3])) &&
5483+
op->src[1]->type == GGML_TYPE_F32;
54815484
case GGML_OP_GET_ROWS:
54825485
{
54835486
switch (op->src[0]->type) {
@@ -5734,6 +5737,13 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
57345737
GGML_UNUSED(dev);
57355738
}
57365739

5740+
static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const ggml_tensor * op) {
5741+
bool res = do_ggml_backend_sycl_device_supports_op(dev, op);
5742+
GGML_SYCL_DEBUG("[SYCL] call %s op->op=%s op->type=%s -> %s\n", __func__, ggml_op_name(op->op),
5743+
ggml_type_name(op->type), res ? "true" : "false");
5744+
return res;
5745+
}
5746+
57375747
static bool ggml_backend_sycl_device_supports_buft(ggml_backend_dev_t dev, ggml_backend_buffer_type_t buft) {
57385748
if (buft->iface.get_name != ggml_backend_sycl_buffer_type_get_name) {
57395749
return false;

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

Lines changed: 74 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1194,6 +1194,66 @@ static void mul_mat_vec_q8_0_q8_1_sycl_switch_ncols(
11941194
}
11951195
}
11961196

1197+
static void mul_mat_vec_q1_0_q8_1_sycl(const void * vx, const void * vy,
1198+
float * dst, const int ncols,
1199+
const int nrows,
1200+
dpct::queue_ptr stream) {
1201+
GGML_ASSERT(ncols % QK1_0 == 0);
1202+
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
1203+
const sycl::range<3> block_nums(1, 1, block_num_y);
1204+
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
1205+
1206+
stream->submit([&](sycl::handler & cgh) {
1207+
cgh.parallel_for(
1208+
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1209+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
1210+
mul_mat_vec_q<QK1_0, QI1_0, block_q1_0,
1211+
VDR_Q1_0_Q8_1_MMVQ, vec_dot_q1_0_q8_1>(
1212+
vx, vy, dst, ncols, nrows, item_ct1);
1213+
});
1214+
});
1215+
}
1216+
1217+
template <int ncols_dst>
1218+
static void mul_mat_vec_q1_0_q8_1_sycl_ncols(
1219+
const void * vx, const void * vy, float * dst,
1220+
const int ncols, const int nrows,
1221+
const int stride_col_y, const int stride_col_dst,
1222+
dpct::queue_ptr stream) {
1223+
GGML_ASSERT(ncols % QK1_0 == 0);
1224+
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
1225+
const sycl::range<3> block_nums(1, 1, block_num_y);
1226+
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
1227+
1228+
stream->submit([&](sycl::handler & cgh) {
1229+
cgh.parallel_for(
1230+
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1231+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
1232+
mul_mat_vec_q_ncols<QK1_0, QI1_0, block_q1_0,
1233+
VDR_Q1_0_Q8_1_MMVQ, vec_dot_q1_0_q8_1, ncols_dst>(
1234+
vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, item_ct1);
1235+
});
1236+
});
1237+
}
1238+
1239+
static void mul_mat_vec_q1_0_q8_1_sycl_switch_ncols(
1240+
const void * vx, const void * vy, float * dst,
1241+
const int ncols, const int nrows, const int ncols_dst,
1242+
const int stride_col_y, const int stride_col_dst,
1243+
dpct::queue_ptr stream) {
1244+
switch (ncols_dst) {
1245+
case 1: mul_mat_vec_q1_0_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break;
1246+
case 2: mul_mat_vec_q1_0_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1247+
case 3: mul_mat_vec_q1_0_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1248+
case 4: mul_mat_vec_q1_0_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1249+
case 5: mul_mat_vec_q1_0_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1250+
case 6: mul_mat_vec_q1_0_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1251+
case 7: mul_mat_vec_q1_0_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1252+
case 8: mul_mat_vec_q1_0_q8_1_sycl_ncols<8>(vx, vy, dst, ncols, nrows, stride_col_y, stride_col_dst, stream); break;
1253+
default: GGML_ABORT("unsupported ncols_dst=%d for Q1_0 multi-col MMVQ", ncols_dst);
1254+
}
1255+
}
1256+
11971257
static void mul_mat_vec_q2_K_q8_1_sycl(const void *vx, const void *vy,
11981258
float *dst, const int ncols,
11991259
const int nrows,
@@ -2120,6 +2180,20 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens
21202180
mul_mat_vec_q8_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
21212181
}
21222182
break;
2183+
case GGML_TYPE_Q1_0:
2184+
if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) {
2185+
const int stride_col_y = src1_padded_col_size / QK8_1;
2186+
const int stride_col_dst = dst->ne[0];
2187+
GGML_SYCL_DEBUG("Calling mul_mat_vec_q1_0_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols);
2188+
mul_mat_vec_q1_0_q8_1_sycl_switch_ncols(
2189+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff,
2190+
src1_ncols, stride_col_y, stride_col_dst, stream);
2191+
return;
2192+
} else if (i == 0 || src1_ncols == 1) {
2193+
GGML_SYCL_DEBUG("Calling mul_mat_vec_q1_0_q8_1_sycl\n");
2194+
mul_mat_vec_q1_0_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream);
2195+
}
2196+
break;
21232197
case GGML_TYPE_Q2_K:
21242198
if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) {
21252199
const int stride_col_y = src1_padded_col_size / QK8_1;

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

Lines changed: 45 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,12 @@
11
#include "outprod.hpp"
2+
#include "convert.hpp"
23

34
void ggml_sycl_op_out_prod(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
45
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
56
const ggml_tensor *src0 = dst->src[0];
67
const ggml_tensor *src1 = dst->src[1];
78

8-
GGML_ASSERT(src0->type == GGML_TYPE_F32);
9+
GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_Q1_0);
910
GGML_ASSERT(src1->type == GGML_TYPE_F32);
1011
GGML_ASSERT(dst->type == GGML_TYPE_F32);
1112
GGML_ASSERT(ggml_is_contiguous(src0));
@@ -20,11 +21,31 @@ void ggml_sycl_op_out_prod(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
2021
GGML_ASSERT(ne01 == ne11); // Inner dimensions must match
2122
GGML_ASSERT(ne0 == ne00); // Output rows match src0 rows
2223
GGML_ASSERT(ne1 == ne10); // Output cols match src1 cols
24+
GGML_ASSERT(ne2 == ne12);
25+
GGML_ASSERT(ne3 == ne13);
26+
GGML_ASSERT(ne2 % ne02 == 0);
27+
GGML_ASSERT(ne3 % ne03 == 0);
2328

2429
// Get data pointers
25-
const float* src0_d = (const float*)src0->data;
26-
const float* src1_d = (const float*)src1->data;
27-
float* dst_d = (float*)dst->data;
30+
const float * src0_d = (const float *) src0->data;
31+
const float * src1_d = (const float *) src1->data;
32+
float * dst_d = (float *) dst->data;
33+
34+
ggml_sycl_pool_alloc<float> src0_as_f32(ctx.pool());
35+
int64_t src0_nb02 = nb02;
36+
int64_t src0_nb03 = nb03;
37+
if (src0->type == GGML_TYPE_Q1_0) {
38+
scope_op_debug_print scope_dbg_print(__func__, "/to_fp32_sycl", dst, /*num_src=*/2,
39+
" : converting src0 Q1_0 to fp32");
40+
src0_d = src0_as_f32.alloc(ne00 * ne01 * ne02 * ne03);
41+
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(src0->type, dst);
42+
GGML_ASSERT(to_fp32_sycl != nullptr);
43+
to_fp32_sycl(src0->data, const_cast<float *>(src0_d), ne00 * ne01 * ne02 * ne03, stream);
44+
45+
// Dequantized src0 buffer is contiguous fp32 [ne00, ne01, ne02, ne03].
46+
src0_nb02 = ne00 * ne01 * (int64_t) sizeof(float);
47+
src0_nb03 = ne00 * ne01 * ne02 * (int64_t) sizeof(float);
48+
}
2849

2950
// GEMM parameters
3051
const float alpha = 1.0f;
@@ -35,12 +56,27 @@ void ggml_sycl_op_out_prod(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
3556
const oneapi::mkl::transpose src1_op = src1_T ? oneapi::mkl::transpose::nontrans : oneapi::mkl::transpose::trans;
3657
const int64_t ldb = (src1_T ? nb10 : nb11) / sizeof(float);
3758

59+
const int64_t r2 = ne2 / ne02;
60+
const int64_t r3 = ne3 / ne03;
61+
3862
try {
39-
// Perform matrix multiplication using oneMKL GEMM
40-
oneapi::mkl::blas::column_major::gemm(*stream, oneapi::mkl::transpose::nontrans, src1_op,
41-
ne0, ne1, ne01, alpha, src0_d, ne00, src1_d, ldb, beta, dst_d, ne0);
42-
}
43-
catch (sycl::exception const& exc) {
63+
// OUT_PROD applies independently to each (i2, i3) destination plane.
64+
for (int64_t i3 = 0; i3 < ne3; ++i3) {
65+
for (int64_t i2 = 0; i2 < ne2; ++i2) {
66+
const int64_t i03 = i3 / r3;
67+
const int64_t i02 = i2 / r2;
68+
69+
const float * src0_plane = (const float *) ((const char *) src0_d + i02 * src0_nb02 + i03 * src0_nb03);
70+
const float * src1_plane = (const float *) ((const char *) src1_d + i2 * nb12 + i3 * nb13);
71+
float * dst_plane = (float *) ((char *) dst_d + i2 * nb2 + i3 * nb3);
72+
73+
// Perform matrix multiplication using oneMKL GEMM
74+
oneapi::mkl::blas::column_major::gemm(*stream, oneapi::mkl::transpose::nontrans, src1_op,
75+
ne0, ne1, ne01, alpha, src0_plane, ne00,
76+
src1_plane, ldb, beta, dst_plane, ne0);
77+
}
78+
}
79+
} catch (sycl::exception const& exc) {
4480
std::cerr << exc.what() << std::endl;
4581
GGML_ASSERT(false);
4682
}

0 commit comments

Comments
 (0)