Skip to content

Commit 4be214f

Browse files
author
Aidan
committed
Use get_multi_ptr
1 parent 023b880 commit 4be214f

File tree

3 files changed

+100
-95
lines changed

3 files changed

+100
-95
lines changed

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -351,4 +351,9 @@ static __dpct_inline__ float warp_reduce_max(float x,
351351
return x;
352352
}
353353

354+
template <typename Tp, int dim = 1>
355+
static __dpct_inline__ Tp* get_pointer(sycl::local_accessor<Tp, dim> acc) {
356+
return acc.template get_multi_ptr<sycl::access::decorated::no>().get();
357+
}
358+
354359
#endif // GGML_SYCL_COMMON_HPP

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

Lines changed: 92 additions & 92 deletions
Original file line numberDiff line numberDiff line change
@@ -1835,10 +1835,10 @@ static void ggml_mul_mat_q4_0_q8_1_sycl(const void *vx, const void *vy,
18351835
mul_mat_q4_0<need_check>(
18361836
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
18371837
nrows_dst, item_ct1,
1838-
tile_x_qs_q4_0_acc_ct1.get_pointer(),
1839-
tile_x_d_q4_0_acc_ct1.get_pointer(),
1840-
tile_y_qs_acc_ct1.get_pointer(),
1841-
tile_y_ds_acc_ct1.get_pointer());
1838+
get_pointer(tile_x_qs_q4_0_acc_ct1),
1839+
get_pointer(tile_x_d_q4_0_acc_ct1),
1840+
get_pointer(tile_y_qs_acc_ct1),
1841+
get_pointer(tile_y_ds_acc_ct1));
18421842
});
18431843
});
18441844
}
@@ -1870,10 +1870,10 @@ static void ggml_mul_mat_q4_0_q8_1_sycl(const void *vx, const void *vy,
18701870
mul_mat_q4_0<need_check>(
18711871
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
18721872
nrows_dst, item_ct1,
1873-
tile_x_qs_q4_0_acc_ct1.get_pointer(),
1874-
tile_x_d_q4_0_acc_ct1.get_pointer(),
1875-
tile_y_qs_acc_ct1.get_pointer(),
1876-
tile_y_ds_acc_ct1.get_pointer());
1873+
get_pointer(tile_x_qs_q4_0_acc_ct1),
1874+
get_pointer(tile_x_d_q4_0_acc_ct1),
1875+
get_pointer(tile_y_qs_acc_ct1),
1876+
get_pointer(tile_y_ds_acc_ct1));
18771877
});
18781878
});
18791879
}
@@ -1950,10 +1950,10 @@ static void ggml_mul_mat_q4_1_q8_1_sycl(const void *vx, const void *vy,
19501950
mul_mat_q4_1<need_check>(
19511951
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
19521952
nrows_dst, item_ct1,
1953-
tile_x_qs_q4_1_acc_ct1.get_pointer(),
1954-
tile_x_dm_q4_1_acc_ct1.get_pointer(),
1955-
tile_y_qs_acc_ct1.get_pointer(),
1956-
tile_y_ds_acc_ct1.get_pointer());
1953+
get_pointer(tile_x_qs_q4_1_acc_ct1),
1954+
get_pointer(tile_x_dm_q4_1_acc_ct1),
1955+
get_pointer(tile_y_qs_acc_ct1),
1956+
get_pointer(tile_y_ds_acc_ct1));
19571957
});
19581958
});
19591959
}
@@ -1985,10 +1985,10 @@ static void ggml_mul_mat_q4_1_q8_1_sycl(const void *vx, const void *vy,
19851985
mul_mat_q4_1<need_check>(
19861986
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
19871987
nrows_dst, item_ct1,
1988-
tile_x_qs_q4_1_acc_ct1.get_pointer(),
1989-
tile_x_dm_q4_1_acc_ct1.get_pointer(),
1990-
tile_y_qs_acc_ct1.get_pointer(),
1991-
tile_y_ds_acc_ct1.get_pointer());
1988+
get_pointer(tile_x_qs_q4_1_acc_ct1),
1989+
get_pointer(tile_x_dm_q4_1_acc_ct1),
1990+
get_pointer(tile_y_qs_acc_ct1),
1991+
get_pointer(tile_y_ds_acc_ct1));
19921992
});
19931993
});
19941994
}
@@ -2065,10 +2065,10 @@ static void ggml_mul_mat_q5_0_q8_1_sycl(const void *vx, const void *vy,
20652065
mul_mat_q5_0<need_check>(
20662066
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
20672067
nrows_dst, item_ct1,
2068-
tile_x_ql_q5_0_acc_ct1.get_pointer(),
2069-
tile_x_d_q5_0_acc_ct1.get_pointer(),
2070-
tile_y_qs_acc_ct1.get_pointer(),
2071-
tile_y_ds_acc_ct1.get_pointer());
2068+
get_pointer(tile_x_ql_q5_0_acc_ct1),
2069+
get_pointer(tile_x_d_q5_0_acc_ct1),
2070+
get_pointer(tile_y_qs_acc_ct1),
2071+
get_pointer(tile_y_ds_acc_ct1));
20722072
});
20732073
});
20742074
}
@@ -2100,10 +2100,10 @@ static void ggml_mul_mat_q5_0_q8_1_sycl(const void *vx, const void *vy,
21002100
mul_mat_q5_0<need_check>(
21012101
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
21022102
nrows_dst, item_ct1,
2103-
tile_x_ql_q5_0_acc_ct1.get_pointer(),
2104-
tile_x_d_q5_0_acc_ct1.get_pointer(),
2105-
tile_y_qs_acc_ct1.get_pointer(),
2106-
tile_y_ds_acc_ct1.get_pointer());
2103+
get_pointer(tile_x_ql_q5_0_acc_ct1),
2104+
get_pointer(tile_x_d_q5_0_acc_ct1),
2105+
get_pointer(tile_y_qs_acc_ct1),
2106+
get_pointer(tile_y_ds_acc_ct1));
21072107
});
21082108
});
21092109
}
@@ -2180,10 +2180,10 @@ static void ggml_mul_mat_q5_1_q8_1_sycl(const void *vx, const void *vy,
21802180
mul_mat_q5_1<need_check>(
21812181
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
21822182
nrows_dst, item_ct1,
2183-
tile_x_ql_q5_1_acc_ct1.get_pointer(),
2184-
tile_x_dm_q5_1_acc_ct1.get_pointer(),
2185-
tile_y_qs_acc_ct1.get_pointer(),
2186-
tile_y_ds_acc_ct1.get_pointer());
2183+
get_pointer(tile_x_ql_q5_1_acc_ct1),
2184+
get_pointer(tile_x_dm_q5_1_acc_ct1),
2185+
get_pointer(tile_y_qs_acc_ct1),
2186+
get_pointer(tile_y_ds_acc_ct1));
21872187
});
21882188
});
21892189
}
@@ -2215,10 +2215,10 @@ static void ggml_mul_mat_q5_1_q8_1_sycl(const void *vx, const void *vy,
22152215
mul_mat_q5_1<need_check>(
22162216
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
22172217
nrows_dst, item_ct1,
2218-
tile_x_ql_q5_1_acc_ct1.get_pointer(),
2219-
tile_x_dm_q5_1_acc_ct1.get_pointer(),
2220-
tile_y_qs_acc_ct1.get_pointer(),
2221-
tile_y_ds_acc_ct1.get_pointer());
2218+
get_pointer(tile_x_ql_q5_1_acc_ct1),
2219+
get_pointer(tile_x_dm_q5_1_acc_ct1),
2220+
get_pointer(tile_y_qs_acc_ct1),
2221+
get_pointer(tile_y_ds_acc_ct1));
22222222
});
22232223
});
22242224
}
@@ -2295,10 +2295,10 @@ static void ggml_mul_mat_q8_0_q8_1_sycl(const void *vx, const void *vy,
22952295
mul_mat_q8_0<need_check>(
22962296
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
22972297
nrows_dst, item_ct1,
2298-
tile_x_qs_q8_0_acc_ct1.get_pointer(),
2299-
tile_x_d_q8_0_acc_ct1.get_pointer(),
2300-
tile_y_qs_acc_ct1.get_pointer(),
2301-
tile_y_ds_acc_ct1.get_pointer());
2298+
get_pointer(tile_x_qs_q8_0_acc_ct1),
2299+
get_pointer(tile_x_d_q8_0_acc_ct1),
2300+
get_pointer(tile_y_qs_acc_ct1),
2301+
get_pointer(tile_y_ds_acc_ct1));
23022302
});
23032303
});
23042304
}
@@ -2330,10 +2330,10 @@ static void ggml_mul_mat_q8_0_q8_1_sycl(const void *vx, const void *vy,
23302330
mul_mat_q8_0<need_check>(
23312331
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
23322332
nrows_dst, item_ct1,
2333-
tile_x_qs_q8_0_acc_ct1.get_pointer(),
2334-
tile_x_d_q8_0_acc_ct1.get_pointer(),
2335-
tile_y_qs_acc_ct1.get_pointer(),
2336-
tile_y_ds_acc_ct1.get_pointer());
2333+
get_pointer(tile_x_qs_q8_0_acc_ct1),
2334+
get_pointer(tile_x_d_q8_0_acc_ct1),
2335+
get_pointer(tile_y_qs_acc_ct1),
2336+
get_pointer(tile_y_ds_acc_ct1));
23372337
});
23382338
});
23392339
}
@@ -2412,11 +2412,11 @@ static void ggml_mul_mat_q2_K_q8_1_sycl(const void *vx, const void *vy,
24122412
mul_mat_q2_K<need_check>(
24132413
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
24142414
nrows_dst, item_ct1,
2415-
tile_x_ql_q2_K_acc_ct1.get_pointer(),
2416-
tile_x_dm_q2_K_acc_ct1.get_pointer(),
2417-
tile_x_sc_q2_K_acc_ct1.get_pointer(),
2418-
tile_y_qs_acc_ct1.get_pointer(),
2419-
tile_y_ds_acc_ct1.get_pointer());
2415+
get_pointer(tile_x_ql_q2_K_acc_ct1),
2416+
get_pointer(tile_x_dm_q2_K_acc_ct1),
2417+
get_pointer(tile_x_sc_q2_K_acc_ct1),
2418+
get_pointer(tile_y_qs_acc_ct1),
2419+
get_pointer(tile_y_ds_acc_ct1));
24202420
});
24212421
});
24222422
}
@@ -2450,11 +2450,11 @@ static void ggml_mul_mat_q2_K_q8_1_sycl(const void *vx, const void *vy,
24502450
mul_mat_q2_K<need_check>(
24512451
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
24522452
nrows_dst, item_ct1,
2453-
tile_x_ql_q2_K_acc_ct1.get_pointer(),
2454-
tile_x_dm_q2_K_acc_ct1.get_pointer(),
2455-
tile_x_sc_q2_K_acc_ct1.get_pointer(),
2456-
tile_y_qs_acc_ct1.get_pointer(),
2457-
tile_y_ds_acc_ct1.get_pointer());
2453+
get_pointer(tile_x_ql_q2_K_acc_ct1),
2454+
get_pointer(tile_x_dm_q2_K_acc_ct1),
2455+
get_pointer(tile_x_sc_q2_K_acc_ct1),
2456+
get_pointer(tile_y_qs_acc_ct1),
2457+
get_pointer(tile_y_ds_acc_ct1));
24582458
});
24592459
});
24602460
}
@@ -2537,12 +2537,12 @@ static void ggml_mul_mat_q3_K_q8_1_sycl(const void *vx, const void *vy,
25372537
mul_mat_q3_K<need_check>(
25382538
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
25392539
nrows_dst, item_ct1,
2540-
tile_x_ql_q3_K_acc_ct1.get_pointer(),
2541-
tile_x_dm_q3_K_acc_ct1.get_pointer(),
2542-
tile_x_qh_q3_K_acc_ct1.get_pointer(),
2543-
tile_x_sc_q3_K_acc_ct1.get_pointer(),
2544-
tile_y_qs_acc_ct1.get_pointer(),
2545-
tile_y_ds_acc_ct1.get_pointer());
2540+
get_pointer(tile_x_ql_q3_K_acc_ct1),
2541+
get_pointer(tile_x_dm_q3_K_acc_ct1),
2542+
get_pointer(tile_x_qh_q3_K_acc_ct1),
2543+
get_pointer(tile_x_sc_q3_K_acc_ct1),
2544+
get_pointer(tile_y_qs_acc_ct1),
2545+
get_pointer(tile_y_ds_acc_ct1));
25462546
});
25472547
});
25482548
}
@@ -2578,12 +2578,12 @@ static void ggml_mul_mat_q3_K_q8_1_sycl(const void *vx, const void *vy,
25782578
mul_mat_q3_K<need_check>(
25792579
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
25802580
nrows_dst, item_ct1,
2581-
tile_x_ql_q3_K_acc_ct1.get_pointer(),
2582-
tile_x_dm_q3_K_acc_ct1.get_pointer(),
2583-
tile_x_qh_q3_K_acc_ct1.get_pointer(),
2584-
tile_x_sc_q3_K_acc_ct1.get_pointer(),
2585-
tile_y_qs_acc_ct1.get_pointer(),
2586-
tile_y_ds_acc_ct1.get_pointer());
2581+
get_pointer(tile_x_ql_q3_K_acc_ct1),
2582+
get_pointer(tile_x_dm_q3_K_acc_ct1),
2583+
get_pointer(tile_x_qh_q3_K_acc_ct1),
2584+
get_pointer(tile_x_sc_q3_K_acc_ct1),
2585+
get_pointer(tile_y_qs_acc_ct1),
2586+
get_pointer(tile_y_ds_acc_ct1));
25872587
});
25882588
});
25892589
}
@@ -2663,11 +2663,11 @@ static void ggml_mul_mat_q4_K_q8_1_sycl(const void *vx, const void *vy,
26632663
mul_mat_q4_K<need_check>(
26642664
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
26652665
nrows_dst, item_ct1,
2666-
tile_x_ql_q4_K_acc_ct1.get_pointer(),
2667-
tile_x_dm_q4_K_acc_ct1.get_pointer(),
2668-
tile_x_sc_q4_K_acc_ct1.get_pointer(),
2669-
tile_y_qs_acc_ct1.get_pointer(),
2670-
tile_y_ds_acc_ct1.get_pointer());
2666+
get_pointer(tile_x_ql_q4_K_acc_ct1),
2667+
get_pointer(tile_x_dm_q4_K_acc_ct1),
2668+
get_pointer(tile_x_sc_q4_K_acc_ct1),
2669+
get_pointer(tile_y_qs_acc_ct1),
2670+
get_pointer(tile_y_ds_acc_ct1));
26712671
});
26722672
});
26732673
}
@@ -2701,11 +2701,11 @@ static void ggml_mul_mat_q4_K_q8_1_sycl(const void *vx, const void *vy,
27012701
mul_mat_q4_K<need_check>(
27022702
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
27032703
nrows_dst, item_ct1,
2704-
tile_x_ql_q4_K_acc_ct1.get_pointer(),
2705-
tile_x_dm_q4_K_acc_ct1.get_pointer(),
2706-
tile_x_sc_q4_K_acc_ct1.get_pointer(),
2707-
tile_y_qs_acc_ct1.get_pointer(),
2708-
tile_y_ds_acc_ct1.get_pointer());
2704+
get_pointer(tile_x_ql_q4_K_acc_ct1),
2705+
get_pointer(tile_x_dm_q4_K_acc_ct1),
2706+
get_pointer(tile_x_sc_q4_K_acc_ct1),
2707+
get_pointer(tile_y_qs_acc_ct1),
2708+
get_pointer(tile_y_ds_acc_ct1));
27092709
});
27102710
});
27112711
}
@@ -2784,11 +2784,11 @@ static void ggml_mul_mat_q5_K_q8_1_sycl(const void *vx, const void *vy,
27842784
mul_mat_q5_K<need_check>(
27852785
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
27862786
nrows_dst, item_ct1,
2787-
tile_x_ql_q5_K_acc_ct1.get_pointer(),
2788-
tile_x_dm_q5_K_acc_ct1.get_pointer(),
2789-
tile_x_sc_q5_K_acc_ct1.get_pointer(),
2790-
tile_y_qs_acc_ct1.get_pointer(),
2791-
tile_y_ds_acc_ct1.get_pointer());
2787+
get_pointer(tile_x_ql_q5_K_acc_ct1),
2788+
get_pointer(tile_x_dm_q5_K_acc_ct1),
2789+
get_pointer(tile_x_sc_q5_K_acc_ct1),
2790+
get_pointer(tile_y_qs_acc_ct1),
2791+
get_pointer(tile_y_ds_acc_ct1));
27922792
});
27932793
});
27942794
}
@@ -2822,11 +2822,11 @@ static void ggml_mul_mat_q5_K_q8_1_sycl(const void *vx, const void *vy,
28222822
mul_mat_q5_K<need_check>(
28232823
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
28242824
nrows_dst, item_ct1,
2825-
tile_x_ql_q5_K_acc_ct1.get_pointer(),
2826-
tile_x_dm_q5_K_acc_ct1.get_pointer(),
2827-
tile_x_sc_q5_K_acc_ct1.get_pointer(),
2828-
tile_y_qs_acc_ct1.get_pointer(),
2829-
tile_y_ds_acc_ct1.get_pointer());
2825+
get_pointer(tile_x_ql_q5_K_acc_ct1),
2826+
get_pointer(tile_x_dm_q5_K_acc_ct1),
2827+
get_pointer(tile_x_sc_q5_K_acc_ct1),
2828+
get_pointer(tile_y_qs_acc_ct1),
2829+
get_pointer(tile_y_ds_acc_ct1));
28302830
});
28312831
});
28322832
}
@@ -2905,11 +2905,11 @@ static void ggml_mul_mat_q6_K_q8_1_sycl(const void *vx, const void *vy,
29052905
mul_mat_q6_K<need_check>(
29062906
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
29072907
nrows_dst, item_ct1,
2908-
tile_x_ql_acc_ct1.get_pointer(),
2909-
tile_x_dm_acc_ct1.get_pointer(),
2910-
tile_x_sc_acc_ct1.get_pointer(),
2911-
tile_y_qs_acc_ct1.get_pointer(),
2912-
tile_y_ds_acc_ct1.get_pointer());
2908+
get_pointer(tile_x_ql_acc_ct1),
2909+
get_pointer(tile_x_dm_acc_ct1),
2910+
get_pointer(tile_x_sc_acc_ct1),
2911+
get_pointer(tile_y_qs_acc_ct1),
2912+
get_pointer(tile_y_ds_acc_ct1));
29132913
});
29142914
});
29152915
}
@@ -2943,11 +2943,11 @@ static void ggml_mul_mat_q6_K_q8_1_sycl(const void *vx, const void *vy,
29432943
mul_mat_q6_K<need_check>(
29442944
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
29452945
nrows_dst, item_ct1,
2946-
tile_x_ql_acc_ct1.get_pointer(),
2947-
tile_x_dm_acc_ct1.get_pointer(),
2948-
tile_x_sc_acc_ct1.get_pointer(),
2949-
tile_y_qs_acc_ct1.get_pointer(),
2950-
tile_y_ds_acc_ct1.get_pointer());
2946+
get_pointer(tile_x_ql_acc_ct1),
2947+
get_pointer(tile_x_dm_acc_ct1),
2948+
get_pointer(tile_x_sc_acc_ct1),
2949+
get_pointer(tile_y_qs_acc_ct1),
2950+
get_pointer(tile_y_ds_acc_ct1));
29512951
});
29522952
});
29532953
}

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -214,7 +214,7 @@ static void norm_f32_sycl(const float* x, float* dst, const int ncols,
214214
[=](sycl::nd_item<3> item_ct1)
215215
[[intel::reqd_sub_group_size(WARP_SIZE)]] {
216216
norm_f32(x, dst, ncols, eps, item_ct1,
217-
s_sum_acc_ct1.get_pointer(), work_group_size);
217+
get_pointer(s_sum_acc_ct1), work_group_size);
218218
});
219219
});
220220
}
@@ -261,7 +261,7 @@ static void group_norm_f32_sycl(const float* x, float* dst,
261261
[[intel::reqd_sub_group_size(WARP_SIZE)]] {
262262
group_norm_f32(x, dst, group_size, ne_elements,
263263
eps_ct4, item_ct1,
264-
s_sum_acc_ct1.get_pointer(), work_group_size);
264+
get_pointer(s_sum_acc_ct1), work_group_size);
265265
});
266266
});
267267
}
@@ -302,7 +302,7 @@ static void rms_norm_f32_sycl(const float* x, float* dst, const int ncols,
302302
[=](sycl::nd_item<3> item_ct1)
303303
[[intel::reqd_sub_group_size(WARP_SIZE)]] {
304304
rms_norm_f32(x, dst, ncols, eps, item_ct1,
305-
s_sum_acc_ct1.get_pointer(), work_group_size);
305+
get_pointer(s_sum_acc_ct1), work_group_size);
306306
});
307307
});
308308
}

0 commit comments

Comments
 (0)