root/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2024-12-25 02:44:36 +00:00
Code Issues Actions 9 Packages Projects Releases Wiki Activity

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

Browse Source
This commit is contained in:
AidanBeltonS 2024-07-10 16:10:49 +01:00 committed by GitHub
parent 6b2a849d1f
commit f4444d992c
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
5 changed files with 103 additions and 97 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
ggml/src/ggml-sycl/common.hpp
View File

@ -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); 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 #endif // GGML_SYCL_COMMON_HPP

2
ggml/src/ggml-sycl/convert.cpp
View File

@ -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::range<3>(1, 1, 32)), sycl::range<3>(1, 1, 32)),
[=](sycl::nd_item<3> item_ct1) { [=](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);
}); });
}); });
} }

184
ggml/src/ggml-sycl/mmq.cpp
View File

@ -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>( mul_mat_q4_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q4_0_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q4_0_acc_ct1),
tile_x_d_q4_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q4_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q4_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q4_0_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q4_0_acc_ct1),
tile_x_d_q4_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q4_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q4_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q4_1_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q4_1_acc_ct1),
tile_x_dm_q4_1_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q4_1_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q4_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q4_1_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q4_1_acc_ct1),
tile_x_dm_q4_1_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q4_1_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_0_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_0_acc_ct1),
tile_x_d_q5_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q5_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_0_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_0_acc_ct1),
tile_x_d_q5_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q5_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_1_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_1_acc_ct1),
tile_x_dm_q5_1_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q5_1_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_1<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_1_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_1_acc_ct1),
tile_x_dm_q5_1_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q5_1_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q8_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q8_0_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q8_0_acc_ct1),
tile_x_d_q8_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q8_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q8_0<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_qs_q8_0_acc_ct1.get_pointer(), get_pointer(tile_x_qs_q8_0_acc_ct1),
tile_x_d_q8_0_acc_ct1.get_pointer(), get_pointer(tile_x_d_q8_0_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q2_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q2_K_acc_ct1),
tile_x_dm_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q2_K_acc_ct1),
tile_x_sc_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q2_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q2_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q2_K_acc_ct1),
tile_x_dm_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q2_K_acc_ct1),
tile_x_sc_q2_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q2_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q3_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q3_K_acc_ct1),
tile_x_dm_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q3_K_acc_ct1),
tile_x_qh_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_qh_q3_K_acc_ct1),
tile_x_sc_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q3_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q3_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q3_K_acc_ct1),
tile_x_dm_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q3_K_acc_ct1),
tile_x_qh_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_qh_q3_K_acc_ct1),
tile_x_sc_q3_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q3_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q4_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q4_K_acc_ct1),
tile_x_dm_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q4_K_acc_ct1),
tile_x_sc_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q4_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q4_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q4_K_acc_ct1),
tile_x_dm_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q4_K_acc_ct1),
tile_x_sc_q4_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q4_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_K_acc_ct1),
tile_x_dm_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q5_K_acc_ct1),
tile_x_sc_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q5_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q5_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_ql_q5_K_acc_ct1),
tile_x_dm_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_dm_q5_K_acc_ct1),
tile_x_sc_q5_K_acc_ct1.get_pointer(), get_pointer(tile_x_sc_q5_K_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q6_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_acc_ct1.get_pointer(), get_pointer(tile_x_ql_acc_ct1),
tile_x_dm_acc_ct1.get_pointer(), get_pointer(tile_x_dm_acc_ct1),
tile_x_sc_acc_ct1.get_pointer(), get_pointer(tile_x_sc_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }
@ -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>( mul_mat_q6_K<need_check>(
vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y,
nrows_dst, item_ct1, nrows_dst, item_ct1,
tile_x_ql_acc_ct1.get_pointer(), get_pointer(tile_x_ql_acc_ct1),
tile_x_dm_acc_ct1.get_pointer(), get_pointer(tile_x_dm_acc_ct1),
tile_x_sc_acc_ct1.get_pointer(), get_pointer(tile_x_sc_acc_ct1),
tile_y_qs_acc_ct1.get_pointer(), get_pointer(tile_y_qs_acc_ct1),
tile_y_ds_acc_ct1.get_pointer()); get_pointer(tile_y_ds_acc_ct1));
}); });
}); });
} }

6
ggml/src/ggml-sycl/norm.cpp
View File

@ -218,7 +218,7 @@ static void norm_f32_sycl(const float* x, float* dst, const int ncols,
[=](sycl::nd_item<3> item_ct1) [=](sycl::nd_item<3> item_ct1)
[[intel::reqd_sub_group_size(WARP_SIZE)]] { [[intel::reqd_sub_group_size(WARP_SIZE)]] {
norm_f32(x, dst, ncols, eps, item_ct1, norm_f32(x, dst, ncols, eps, item_ct1,
s_sum_acc_ct1.get_pointer(), work_group_size); get_pointer(s_sum_acc_ct1), work_group_size);
}); });
}); });
} }
@ -265,7 +265,7 @@ static void group_norm_f32_sycl(const float* x, float* dst,
[[intel::reqd_sub_group_size(WARP_SIZE)]] { [[intel::reqd_sub_group_size(WARP_SIZE)]] {
group_norm_f32(x, dst, group_size, ne_elements, group_norm_f32(x, dst, group_size, ne_elements,
eps_ct4, item_ct1, eps_ct4, item_ct1,
s_sum_acc_ct1.get_pointer(), work_group_size); get_pointer(s_sum_acc_ct1), work_group_size);
}); });
}); });
} }
@ -306,7 +306,7 @@ static void rms_norm_f32_sycl(const float* x, float* dst, const int ncols,
[=](sycl::nd_item<3> item_ct1) [=](sycl::nd_item<3> item_ct1)
[[intel::reqd_sub_group_size(WARP_SIZE)]] { [[intel::reqd_sub_group_size(WARP_SIZE)]] {
rms_norm_f32(x, dst, ncols, eps, item_ct1, rms_norm_f32(x, dst, ncols, eps, item_ct1,
s_sum_acc_ct1.get_pointer(), work_group_size); get_pointer(s_sum_acc_ct1), work_group_size);
}); });
}); });
} }

2
ggml/src/ggml-sycl/softmax.cpp
View File

@ -136,7 +136,7 @@ static void soft_max_f32_submitter(const float * x, const float * mask, float *
soft_max_f32<vals_smem, ncols_template, block_size_template>(x, mask, dst, ncols_par, soft_max_f32<vals_smem, ncols_template, block_size_template>(x, mask, dst, ncols_par,
nrows_y, scale, max_bias, m0, nrows_y, scale, max_bias, m0,
m1, n_head_log2, item_ct1, m1, n_head_log2, item_ct1,
local_buf_acc.get_pointer()); get_pointer(local_buf_acc));
}); });
}); });
} }