Skip to content

[SYCL] Use multi_ptr to clean up deprecated warnings #8256

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Jul 10, 2024
Merged

[SYCL] Use multi_ptr to clean up deprecated warnings #8256

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
03da9a6
Use get_multi_ptr
Jul 2, 2024
e0b8a57
Update get_pointer ggml-sycl.cpp
Jul 3, 2024
cf94e5d
Remove dim default val
Jul 3, 2024
4a32e6a
Update submitters
Jul 10, 2024
1e5ecc5
Fix rebase
Jul 10, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions ggml/src/ggml-sycl/common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,4 +346,10 @@ inline sycl::vec<Tp, n> vec_aligned_load(const Tp* aligned_ptr) {
return *reinterpret_cast<const sycl::vec<Tp, n>*>(aligned_ptr);
}

// Helper for accessing pointers with no warnings
template <typename Tp, int dim>
static __dpct_inline__ Tp* get_pointer(sycl::local_accessor<Tp, dim> acc) {
return acc.template get_multi_ptr<sycl::access::decorated::no>().get();
}

#endif // GGML_SYCL_COMMON_HPP
2 changes: 1 addition & 1 deletion ggml/src/ggml-sycl/convert.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -158,7 +158,7 @@ static void dequantize_row_q4_K_sycl(const void *vx, dst_t *y, const int k,
sycl::range<3>(1, 1, 32),
sycl::range<3>(1, 1, 32)),
[=](sycl::nd_item<3> item_ct1) {
dequantize_block_q4_K(vx, y, scale_local_acc.get_pointer(), item_ct1);
dequantize_block_q4_K(vx, y, get_pointer(scale_local_acc), item_ct1);
});
});
}
Expand Down
184 changes: 92 additions & 92 deletions ggml/src/ggml-sycl/mmq.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1835,10 +1835,10 @@ static void ggml_mul_mat_q4_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q4_0_acc_ct1.get_pointer(),
tile_x_d_q4_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q4_0_acc_ct1),
get_pointer(tile_x_d_q4_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -1870,10 +1870,10 @@ static void ggml_mul_mat_q4_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q4_0_acc_ct1.get_pointer(),
tile_x_d_q4_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q4_0_acc_ct1),
get_pointer(tile_x_d_q4_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -1950,10 +1950,10 @@ static void ggml_mul_mat_q4_1_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q4_1_acc_ct1.get_pointer(),
tile_x_dm_q4_1_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q4_1_acc_ct1),
get_pointer(tile_x_dm_q4_1_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -1985,10 +1985,10 @@ static void ggml_mul_mat_q4_1_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q4_1_acc_ct1.get_pointer(),
tile_x_dm_q4_1_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q4_1_acc_ct1),
get_pointer(tile_x_dm_q4_1_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2065,10 +2065,10 @@ static void ggml_mul_mat_q5_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_0_acc_ct1.get_pointer(),
tile_x_d_q5_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_0_acc_ct1),
get_pointer(tile_x_d_q5_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2100,10 +2100,10 @@ static void ggml_mul_mat_q5_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_0_acc_ct1.get_pointer(),
tile_x_d_q5_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_0_acc_ct1),
get_pointer(tile_x_d_q5_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2180,10 +2180,10 @@ static void ggml_mul_mat_q5_1_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_1_acc_ct1.get_pointer(),
tile_x_dm_q5_1_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_1_acc_ct1),
get_pointer(tile_x_dm_q5_1_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2215,10 +2215,10 @@ static void ggml_mul_mat_q5_1_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_1_acc_ct1.get_pointer(),
tile_x_dm_q5_1_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_1_acc_ct1),
get_pointer(tile_x_dm_q5_1_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2295,10 +2295,10 @@ static void ggml_mul_mat_q8_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q8_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q8_0_acc_ct1.get_pointer(),
tile_x_d_q8_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q8_0_acc_ct1),
get_pointer(tile_x_d_q8_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2330,10 +2330,10 @@ static void ggml_mul_mat_q8_0_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q8_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_qs_q8_0_acc_ct1.get_pointer(),
tile_x_d_q8_0_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_qs_q8_0_acc_ct1),
get_pointer(tile_x_d_q8_0_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2412,11 +2412,11 @@ static void ggml_mul_mat_q2_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q2_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q2_K_acc_ct1.get_pointer(),
tile_x_dm_q2_K_acc_ct1.get_pointer(),
tile_x_sc_q2_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q2_K_acc_ct1),
get_pointer(tile_x_dm_q2_K_acc_ct1),
get_pointer(tile_x_sc_q2_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2450,11 +2450,11 @@ static void ggml_mul_mat_q2_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q2_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q2_K_acc_ct1.get_pointer(),
tile_x_dm_q2_K_acc_ct1.get_pointer(),
tile_x_sc_q2_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q2_K_acc_ct1),
get_pointer(tile_x_dm_q2_K_acc_ct1),
get_pointer(tile_x_sc_q2_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2537,12 +2537,12 @@ static void ggml_mul_mat_q3_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q3_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q3_K_acc_ct1.get_pointer(),
tile_x_dm_q3_K_acc_ct1.get_pointer(),
tile_x_qh_q3_K_acc_ct1.get_pointer(),
tile_x_sc_q3_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q3_K_acc_ct1),
get_pointer(tile_x_dm_q3_K_acc_ct1),
get_pointer(tile_x_qh_q3_K_acc_ct1),
get_pointer(tile_x_sc_q3_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2578,12 +2578,12 @@ static void ggml_mul_mat_q3_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q3_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q3_K_acc_ct1.get_pointer(),
tile_x_dm_q3_K_acc_ct1.get_pointer(),
tile_x_qh_q3_K_acc_ct1.get_pointer(),
tile_x_sc_q3_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q3_K_acc_ct1),
get_pointer(tile_x_dm_q3_K_acc_ct1),
get_pointer(tile_x_qh_q3_K_acc_ct1),
get_pointer(tile_x_sc_q3_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2663,11 +2663,11 @@ static void ggml_mul_mat_q4_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q4_K_acc_ct1.get_pointer(),
tile_x_dm_q4_K_acc_ct1.get_pointer(),
tile_x_sc_q4_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q4_K_acc_ct1),
get_pointer(tile_x_dm_q4_K_acc_ct1),
get_pointer(tile_x_sc_q4_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2701,11 +2701,11 @@ static void ggml_mul_mat_q4_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q4_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q4_K_acc_ct1.get_pointer(),
tile_x_dm_q4_K_acc_ct1.get_pointer(),
tile_x_sc_q4_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q4_K_acc_ct1),
get_pointer(tile_x_dm_q4_K_acc_ct1),
get_pointer(tile_x_sc_q4_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2784,11 +2784,11 @@ static void ggml_mul_mat_q5_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_K_acc_ct1.get_pointer(),
tile_x_dm_q5_K_acc_ct1.get_pointer(),
tile_x_sc_q5_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_K_acc_ct1),
get_pointer(tile_x_dm_q5_K_acc_ct1),
get_pointer(tile_x_sc_q5_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2822,11 +2822,11 @@ static void ggml_mul_mat_q5_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q5_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_q5_K_acc_ct1.get_pointer(),
tile_x_dm_q5_K_acc_ct1.get_pointer(),
tile_x_sc_q5_K_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_q5_K_acc_ct1),
get_pointer(tile_x_dm_q5_K_acc_ct1),
get_pointer(tile_x_sc_q5_K_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2905,11 +2905,11 @@ static void ggml_mul_mat_q6_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q6_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_acc_ct1.get_pointer(),
tile_x_dm_acc_ct1.get_pointer(),
tile_x_sc_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_acc_ct1),
get_pointer(tile_x_dm_acc_ct1),
get_pointer(tile_x_sc_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down Expand Up @@ -2943,11 +2943,11 @@ static void ggml_mul_mat_q6_K_q8_1_sycl(const void *vx, const void *vy,
mul_mat_q6_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1,
tile_x_ql_acc_ct1.get_pointer(),
tile_x_dm_acc_ct1.get_pointer(),
tile_x_sc_acc_ct1.get_pointer(),
tile_y_qs_acc_ct1.get_pointer(),
tile_y_ds_acc_ct1.get_pointer());
get_pointer(tile_x_ql_acc_ct1),
get_pointer(tile_x_dm_acc_ct1),
get_pointer(tile_x_sc_acc_ct1),
get_pointer(tile_y_qs_acc_ct1),
get_pointer(tile_y_ds_acc_ct1));
});
});
}
Expand Down
Loading
Loading