Skip to content

Commit f4444d9

Browse files
author
AidanBeltonS
authored
[SYCL] Use multi_ptr to clean up deprecated warnings (#8256)
1 parent 6b2a849 commit f4444d9

File tree

5 files changed

+103
-97
lines changed

5 files changed

+103
-97
lines changed

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -346,4 +346,10 @@ inline sycl::vec<Tp, n> vec_aligned_load(const Tp* aligned_ptr) {
346346
return *reinterpret_cast<const sycl::vec<Tp, n>*>(aligned_ptr);
347347
}
348348

349+
// Helper for accessing pointers with no warnings
350+
template <typename Tp, int dim>
351+
static __dpct_inline__ Tp* get_pointer(sycl::local_accessor<Tp, dim> acc) {
352+
return acc.template get_multi_ptr<sycl::access::decorated::no>().get();
353+
}
354+
349355
#endif // GGML_SYCL_COMMON_HPP

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -158,7 +158,7 @@ static void dequantize_row_q4_K_sycl(const void *vx, dst_t *y, const int k,
158158
sycl::range<3>(1, 1, 32),
159159
sycl::range<3>(1, 1, 32)),
160160
[=](sycl::nd_item<3> item_ct1) {
161-
dequantize_block_q4_K(vx, y, scale_local_acc.get_pointer(), item_ct1);
161+
dequantize_block_q4_K(vx, y, get_pointer(scale_local_acc), item_ct1);
162162
});
163163
});
164164
}

‎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
}

0 commit comments

Comments
 (0)