26 auto edges = get_edges();
32 Tscal Rker2 = kernel_radius * kernel_radius;
38 count, shamsys::instance::get_compute_scheduler_ptr());
40 count, shamsys::instance::get_compute_scheduler_ptr());
45 pcache, edges.xyz.get_spans().get(
id), edges.hpart.get_spans().get(
id)},
50 const auto ploop_ptrs,
53 u32 *neigh_count_true,
54 u32 *neigh_count_all) {
57 Tscal h_a = hpart[id_a];
58 Tvec xyz_a = xyz[id_a];
62 particle_looper.for_each_object(id_a, [&](
u32 id_b) {
65 Tvec r_ab = xyz_a - xyz[id_b];
66 Tscal rab2 = sycl::dot(r_ab, r_ab);
67 Tscal h_b = hpart[id_b];
69 if (rab2 > h_a * h_a * Rker2 && rab2 > h_b * h_b * Rker2) {
76 neigh_count_all[id_a] = cnt_all;
77 neigh_count_true[id_a] = cnt_true;
81 std::vector<u32> neigh_cnt_vec = neigh_count_true.
copy_to_stdvec();
83 double max = *std::max_element(neigh_cnt_vec.begin(), neigh_cnt_vec.end());
84 double min = *std::min_element(neigh_cnt_vec.begin(), neigh_cnt_vec.end());
85 double mean = std::accumulate(neigh_cnt_vec.begin(), neigh_cnt_vec.end(), 0.0)
86 / neigh_cnt_vec.size();
87 double stddev = std::sqrt(
89 neigh_cnt_vec.begin(),
92 [mean](
double acc,
double x) {
93 return acc + (x - mean) * (x - mean);
95 / neigh_cnt_vec.size());
96 logger::raw_ln(
"(true) max, min, mean, stddev =", max, min, mean, stddev);
100 std::vector<u32> neigh_cnt_vec = neigh_count_all.
copy_to_stdvec();
102 double max = *std::max_element(neigh_cnt_vec.begin(), neigh_cnt_vec.end());
103 double min = *std::min_element(neigh_cnt_vec.begin(), neigh_cnt_vec.end());
104 double mean = std::accumulate(neigh_cnt_vec.begin(), neigh_cnt_vec.end(), 0.0)
105 / neigh_cnt_vec.size();
106 double stddev = std::sqrt(
108 neigh_cnt_vec.begin(),
111 [mean](
double acc,
double x) {
112 return acc + (x - mean) * (x - mean);
114 / neigh_cnt_vec.size());
115 logger::raw_ln(
"(all) max, min, mean, stddev =", max, min, mean, stddev);
void kernel_call(sham::DeviceQueue &q, RefIn in, RefOut in_out, u32 n, Functor &&func, SourceLocation &&callsite=SourceLocation{})
Submit a kernel to a SYCL queue.