Shamrock 2025.10.0
Astrophysical Code
Loading...
Searching...
No Matches
sycl_utils.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
19#include "shambase/exception.hpp"
20#include "shambase/integer.hpp"
21#include "shambase/stacktrace.hpp"
22#include "shambackends/sycl.hpp"
23#include "shambackends/vec.hpp"
24#include "shamcomm/logs.hpp"
25#include <stdexcept>
26
27#ifdef SHAMROCK_USE_NVTX
28 #include <nvtx3/nvtx3.hpp>
29#endif
30
31namespace shambase {
32
42 template<class T>
43 void check_buffer_size(
44 sycl::buffer<T> &buf, u64 max_range, const SourceLocation loc = SourceLocation()) {
45 if (buf.size() < max_range) {
46 throw make_except_with_loc<std::invalid_argument>("buffer is too small", loc);
47 }
48 }
49
56 inline std::string getDevice_type(const sycl::device &Device) {
57 auto DeviceType = Device.get_info<sycl::info::device::device_type>();
58 switch (DeviceType) {
59 case sycl::info::device_type::cpu : return "CPU";
60 case sycl::info::device_type::gpu : return "GPU";
61 case sycl::info::device_type::host : return "HOST";
62 case sycl::info::device_type::accelerator: return "ACCELERATOR";
63 default : return "UNKNOWN";
64 }
65 }
66
74 inline sycl::nd_range<1> make_range(u32 length, const u32 group_size = 32) {
75 u32 group_cnt = shambase::group_count(length, group_size);
76 u32 len = group_cnt * group_size;
77 return sycl::nd_range<1>{len, group_size};
78 }
79
80 enum ParallelForWrapMode { PARALLEL_FOR, PARALLEL_FOR_ROUND, ND_RANGE };
81
82#ifdef SHAMROCK_LOOP_DEFAULT_PARALLEL_FOR
83 constexpr ParallelForWrapMode default_loop_mode = PARALLEL_FOR;
84#endif
85
86#ifdef SHAMROCK_LOOP_DEFAULT_PARALLEL_FOR_ROUND
87 constexpr ParallelForWrapMode default_loop_mode = PARALLEL_FOR_ROUND;
88#endif
89
90#ifdef SHAMROCK_LOOP_DEFAULT_ND_RANGE
91 constexpr ParallelForWrapMode default_loop_mode = ND_RANGE;
92#endif
93
94 constexpr u32 default_gsize = SHAMROCK_LOOP_GSIZE;
95 constexpr u32 default_gsize_2d = 16;
96 constexpr u32 default_gsize_3d = 4;
97
98 template<
99 u32 group_size = default_gsize,
100 ParallelForWrapMode mode = default_loop_mode,
101 class LambdaKernel>
102 inline void parallel_for(sycl::handler &cgh, u32 length, const char *name, LambdaKernel &&ker) {
103
104#ifdef SHAMROCK_USE_NVTX
105 nvtxRangePush(name);
106#endif
107
108 shamlog_debug_sycl_ln("SYCL", sham::format("parallel_for {} N={}", name, length));
109
110 if constexpr (mode == PARALLEL_FOR) {
111
112 cgh.parallel_for(sycl::range<1>{length}, [=](sycl::item<1> id) {
113 ker(id.get_linear_id());
114 });
115
116 } else if constexpr (mode == PARALLEL_FOR_ROUND) {
117
118 u32 len = shambase::group_count(length, group_size) * group_size;
119
120 cgh.parallel_for(sycl::range<1>{len}, [=](sycl::item<1> id) {
121 u64 gid = id.get_linear_id();
122 if (gid >= length)
123 return;
124
125 ker(gid);
126 });
127
128 } else if constexpr (mode == ND_RANGE) {
129
130 cgh.parallel_for(make_range(length, group_size), [=](sycl::nd_item<1> id) {
131 u64 gid = id.get_global_linear_id();
132 if (gid >= length)
133 return;
134
135 ker(gid);
136 });
137
138 } else {
139 throw_unimplemented();
140 }
141
142#ifdef SHAMROCK_USE_NVTX
143 nvtxRangePop();
144#endif
145 }
146
147 template<
148 u32 group_size = default_gsize_2d,
149 ParallelForWrapMode mode = default_loop_mode,
150 class LambdaKernel>
151 inline void parallel_for_2d(
152 sycl::handler &cgh, u32 length_x, u32 length_y, const char *name, LambdaKernel &&ker) {
153
154#ifdef SHAMROCK_USE_NVTX
155 nvtxRangePush(name);
156#endif
157
158 shamlog_debug_sycl_ln(
159 "SYCL", sham::format("parallel_for {} N={} {}", name, length_x, length_y));
160
161 if constexpr (mode == PARALLEL_FOR) {
162
163 cgh.parallel_for(sycl::range<2>{length_x, length_y}, [=](sycl::item<2> id) {
164 ker(id.get_id(0), id.get_id(1));
165 });
166
167 } else if constexpr (mode == PARALLEL_FOR_ROUND) {
168
169 u32 len_x = shambase::group_count(length_x, group_size) * group_size;
170 u32 len_y = shambase::group_count(length_y, group_size) * group_size;
171
172 cgh.parallel_for(sycl::range<2>{len_x, len_y}, [=](sycl::item<2> id) {
173 if (id.get_id(0) >= length_x || id.get_id(1) >= length_y)
174 return;
175
176 ker(id.get_id(0), id.get_id(1));
177 });
178
179 } else if constexpr (mode == ND_RANGE) {
180
181 sycl::nd_range<1> rx = make_range(length_x, group_size);
182 sycl::nd_range<1> ry = make_range(length_y, group_size);
183
184 sycl::range<2> tmp_s{rx.get_global_range().size(), ry.get_global_range().size()};
185 sycl::range<2> tmp_g{rx.get_group_range().size(), ry.get_group_range().size()};
186
187 cgh.parallel_for(sycl::nd_range<2>{tmp_s, tmp_g}, [=](sycl::nd_item<2> id) {
188 if (id.get_global_id(0) >= length_x || id.get_global_id(1) >= length_y)
189 return;
190
191 ker(id.get_global_id(0), id.get_global_id(1));
192 });
193
194 } else {
195 throw_unimplemented();
196 }
197
198#ifdef SHAMROCK_USE_NVTX
199 nvtxRangePop();
200#endif
201 }
202
203 template<
204 u32 group_size = default_gsize_3d,
205 ParallelForWrapMode mode = default_loop_mode,
206 class LambdaKernel>
207 inline void parallel_for_3d(
208 sycl::handler &cgh,
209 u32 length_x,
210 u32 length_y,
211 u32 length_z,
212 const char *name,
213 LambdaKernel &&ker) {
214
215#ifdef SHAMROCK_USE_NVTX
216 nvtxRangePush(name);
217#endif
218
219 shamlog_debug_sycl_ln(
220 "SYCL", sham::format("parallel_for {} N={} {} {}", name, length_x, length_y, length_z));
221
222 if constexpr (mode == PARALLEL_FOR) {
223
224 cgh.parallel_for(sycl::range<3>{length_x, length_y, length_z}, [=](sycl::item<3> id) {
225 ker(id.get_id(0), id.get_id(1), id.get_id(2));
226 });
227
228 } else if constexpr (mode == PARALLEL_FOR_ROUND) {
229
230 u32 len_x = shambase::group_count(length_x, group_size) * group_size;
231 u32 len_y = shambase::group_count(length_y, group_size) * group_size;
232 u32 len_z = shambase::group_count(length_z, group_size) * group_size;
233
234 cgh.parallel_for(sycl::range<3>{len_x, len_y, len_z}, [=](sycl::item<3> id) {
235 if (id.get_id(0) >= length_x || id.get_id(1) >= length_y
236 || id.get_id(2) >= length_z)
237 return;
238
239 ker(id.get_id(0), id.get_id(1), id.get_id(2));
240 });
241
242 } else if constexpr (mode == ND_RANGE) {
243
244 sycl::nd_range<1> rx = make_range(length_x, group_size);
245 sycl::nd_range<1> ry = make_range(length_y, group_size);
246 sycl::nd_range<1> rz = make_range(length_z, group_size);
247
248 sycl::range<3> tmp_s{
249 rx.get_global_range().size(),
250 ry.get_global_range().size(),
251 rz.get_global_range().size()};
252 sycl::range<3> tmp_g{
253 rx.get_group_range().size(),
254 ry.get_group_range().size(),
255 rz.get_group_range().size()};
256
257 cgh.parallel_for(sycl::nd_range<3>{tmp_s, tmp_g}, [=](sycl::nd_item<3> id) {
258 if (id.get_global_id(0) >= length_x || id.get_global_id(1) >= length_y
259 || id.get_global_id(2) >= length_z)
260 return;
261
262 ker(id.get_global_id(0), id.get_global_id(1), id.get_global_id(2));
263 });
264
265 } else {
266 throw_unimplemented();
267 }
268
269#ifdef SHAMROCK_USE_NVTX
270 nvtxRangePop();
271#endif
272 }
273
274 template<ParallelForWrapMode mode = default_loop_mode, class LambdaKernel>
275 inline void parallel_for_gsize(
276 sycl::handler &cgh, u32 length, u32 group_size, const char *name, LambdaKernel &&ker) {
277
278#ifdef SHAMROCK_USE_NVTX
279 nvtxRangePush(name);
280#endif
281
282 if constexpr (mode == PARALLEL_FOR) {
283
284 cgh.parallel_for(sycl::range<1>{length}, [=](sycl::item<1> id) {
285 ker(id.get_linear_id());
286 });
287
288 } else if constexpr (mode == PARALLEL_FOR_ROUND) {
289
290 u32 len = shambase::group_count(length, group_size) * group_size;
291
292 cgh.parallel_for(sycl::range<1>{len}, [=](sycl::item<1> id) {
293 u64 gid = id.get_linear_id();
294 if (gid >= length)
295 return;
296
297 ker(gid);
298 });
299
300 } else if constexpr (mode == ND_RANGE) {
301
302 cgh.parallel_for(make_range(length, group_size), [=](sycl::nd_item<1> id) {
303 u64 gid = id.get_global_linear_id();
304 if (gid >= length)
305 return;
306
307 ker(gid);
308 });
309
310 } else {
311 throw_unimplemented();
312 }
313
314#ifdef SHAMROCK_USE_NVTX
315 nvtxRangePop();
316#endif
317 }
318
319 inline void check_queue_state(sycl::queue &q, SourceLocation loc = SourceLocation()) {
320 shamlog_debug_sycl_ln("SYCL", "checking queue state", loc.format_one_line());
321 q.wait_and_throw();
322 shamlog_debug_sycl_ln("SYCL", "checking queue state : OK !");
323 }
324
325} // namespace shambase
u32
std::uint32_t u32
32 bit unsigned integer
Definition aliases_int.hpp:27
u64
std::uint64_t u64
64 bit unsigned integer
Definition aliases_int.hpp:26
exception.hpp
This header file contains utility functions related to exception handling in the code.
shambase
namespace for basic c++ utilities
Definition sycl_builtins.hpp:25
shambase::group_count
constexpr u32 group_count(u32 len, u32 group_size)
Calculates the number of groups based on the length and group size.
Definition integer.hpp:125
shambase::make_range
sycl::nd_range< 1 > make_range(u32 length, const u32 group_size=32)
Generate a sycl nd range out of a group size and length.
Definition sycl_utils.hpp:74
shambase::check_buffer_size
void check_buffer_size(sycl::buffer< T > &buf, u64 max_range, const SourceLocation loc=SourceLocation())
check that the size of a sycl buffer is below or equal to the value of max range throw if it is not t...
Definition sycl_utils.hpp:43
shambase::getDevice_type
std::string getDevice_type(const sycl::device &Device)
Get the Device Type Name.
Definition sycl_utils.hpp:56
shambase::make_except_with_loc
ExcptTypes make_except_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Create an exception with a message and a location.
Definition exception.hpp:105
shambase::throw_unimplemented
void throw_unimplemented(SourceLocation loc=SourceLocation{})
Throw a std::runtime_error saying that the function is unimplemented.
Definition exception.hpp:153
stacktrace.hpp
This file contains the definition for the stacktrace related functionality.
SourceLocation
provide information about the source location
Definition SourceLocation.hpp:61
Generated on for Shamrock by doxygen 1.15.0