Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
syclreduction.hpp
Go to the documentation of this file.
1// -------------------------------------------------------//
2//
3// SHAMROCK code for hydrodynamics
4// Copyright (c) 2021-2026 Timothée David--Cléris <tim.shamrock@proton.me>
5// SPDX-License-Identifier: CeCILL Free Software License Agreement v2.1
6// Shamrock is licensed under the CeCILL 2.1 License, see LICENSE for more information
7//
8// -------------------------------------------------------//
9
10#pragma once
11
18#include "shambase/exception.hpp"
19#include "shambackends/DeviceBuffer.hpp"
20#include "shambackends/sycl.hpp"
21#include <memory>
22#include <stdexcept>
23
24namespace syclalg {
25
26 // TODO to optimize
27 template<class T>
28 [[deprecated("please use the shamalgs library instead")]]
29 inline T get_max(sycl::queue &queue, const std::unique_ptr<sycl::buffer<T>> &buf, u32 len) {
30
31 T accum;
32
33 if (buf) {
34 accum = buf->get_host_access()[0];
35
36 {
37 sycl::host_accessor acc{*buf, sycl::read_only};
38
39 // queue.submit([&](sycl::handler &cgh) {
40 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
41
42 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
43 // u32 i = (u32)item.get_id(0);
44
45 // });
46 // });
47
48 for (u32 i = 0; i < len; i++) {
49 accum = sycl::max(accum, acc[i]);
50 }
51 }
52 } else {
53 throw shambase::make_except_with_loc<std::runtime_error>(
54 "syclalg::get_max : input buffer not allocated");
55 }
56
57 return accum;
58 }
59
60 template<class T>
61 [[deprecated("please use the shamalgs library instead")]]
62 inline T get_max(sham::DeviceBuffer<T> &buf, u32 len) {
63
64 T accum;
65
66 if (!buf.is_empty()) {
67 auto vec = buf.copy_to_stdvec();
68 auto accum = vec[0];
69
70 {
71
72 // queue.submit([&](sycl::handler &cgh) {
73 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
74
75 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
76 // u32 i = (u32)item.get_id(0);
77
78 // });
79 // });
80
81 for (u32 i = 0; i < len; i++) {
82 accum = sycl::max(accum, vec[i]);
83 }
84 }
85 } else {
86 throw shambase::make_except_with_loc<std::runtime_error>(
87 "syclalg::get_max : input buffer not allocated");
88 }
89
90 return accum;
91 }
92
93 // TODO to optimize
94 template<class T>
95 [[deprecated("please use the shamalgs library instead")]]
96 inline T get_min(sycl::queue &queue, const std::unique_ptr<sycl::buffer<T>> &buf, u32 len) {
97
98 T accum;
99
100 if (buf) {
101
102 accum = buf->get_host_access()[0];
103
104 {
105 sycl::host_accessor acc{*buf, sycl::read_only};
106
107 // queue.submit([&](sycl::handler &cgh) {
108 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
109
110 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
111 // u32 i = (u32)item.get_id(0);
112
113 // });
114 // });
115
116 for (u32 i = 0; i < len; i++) {
117 accum = sycl::min(accum, acc[i]);
118 }
119 }
120
121 } else {
122 throw shambase::make_except_with_loc<std::runtime_error>(
123 "syclalg::get_min : input buffer not allocated");
124 }
125
126 return accum;
127 }
128
129 template<class T>
130 [[deprecated("please use the shamalgs library instead")]]
131 inline T get_min(sham::DeviceBuffer<T> &buf, u32 len) {
132
133 T accum;
134
135 if (!buf.is_empty()) {
136 auto vec = buf.copy_to_stdvec();
137 auto accum = vec[0];
138
139 {
140
141 // queue.submit([&](sycl::handler &cgh) {
142 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
143
144 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
145 // u32 i = (u32)item.get_id(0);
146
147 // });
148 // });
149
150 for (u32 i = 0; i < len; i++) {
151 accum = sycl::min(accum, vec[i]);
152 }
153 }
154 } else {
155 throw shambase::make_except_with_loc<std::runtime_error>(
156 "syclalg::get_min : input buffer not allocated");
157 }
158
159 return accum;
160 }
161
162 template<class T>
163 [[deprecated("please use the shamalgs library instead")]]
164 inline std::tuple<T, T> get_min_max(
165 sycl::queue &queue, const std::unique_ptr<sycl::buffer<T>> &buf, u32 len) {
166
167 T accum_min, accum_max;
168
169 if (buf) {
170
171 accum_min = buf->get_host_access()[0];
172 accum_max = buf->get_host_access()[0];
173
174 {
175 sycl::host_accessor acc{*buf, sycl::read_only};
176
177 // queue.submit([&](sycl::handler &cgh) {
178 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
179
180 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
181 // u32 i = (u32)item.get_id(0);
182
183 // });
184 // });
185
186 for (u32 i = 0; i < len; i++) {
187 accum_min = sycl::min(accum_min, acc[i]);
188 accum_max = sycl::max(accum_max, acc[i]);
189 }
190 }
191
192 } else {
193 throw shambase::make_except_with_loc<std::runtime_error>(
194 "syclalg::get_max : input buffer not allocated");
195 }
196
197 return {accum_min, accum_max};
198 }
199
200 template<class T>
201 [[deprecated("please use the shamalgs library instead")]]
202 inline std::tuple<T, T> get_min_max(sham::DeviceBuffer<T> &buf, u32 len) {
203
204 T accum_min, accum_max;
205
206 if (!buf.is_empty()) {
207 auto vec = buf.copy_to_stdvec();
208
209 auto accum_min = vec[0];
210 auto accum_max = vec[0];
211
212 {
213
214 // queue.submit([&](sycl::handler &cgh) {
215 // auto acc = buf->get_access<sycl::access::mode::read>(cgh);
216
217 // cgh.parallel_for(sycl::range(buf->size()), [=](sycl::item<1> item) {
218 // u32 i = (u32)item.get_id(0);
219
220 // });
221 // });
222
223 for (u32 i = 0; i < len; i++) {
224 accum_min = sycl::min(accum_min, vec[i]);
225 accum_max = sycl::max(accum_max, vec[i]);
226 }
227 }
228
229 } else {
230 throw shambase::make_except_with_loc<std::runtime_error>(
231 "syclalg::get_max : input buffer not allocated");
232 }
233
234 return {accum_min, accum_max};
235 }
236
237} // namespace syclalg
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
sham::DeviceBuffer
A buffer allocated in USM (Unified Shared Memory)
Definition DeviceBuffer.hpp:106
sham::DeviceBuffer::copy_to_stdvec
std::vector< T > copy_to_stdvec() const
Copy the content of the buffer to a std::vector.
Definition DeviceBuffer.hpp:510
sham::DeviceBuffer::is_empty
bool is_empty() const
Check if the buffer is empty.
Definition DeviceBuffer.hpp:490
exception.hpp
This header file contains utility functions related to exception handling in the code.
shambase::throw_with_loc
void throw_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Throw an exception and append the source location to it.
Definition exception.hpp:132
Generated on Wed May 13 2026 13:21:49 for Shamrock by doxygen 1.9.8