DFT-EFE
 
Loading...
Searching...
No Matches
DeviceKernelLauncherHelpers.h
Go to the documentation of this file.
1// ---------------------------------------------------------------------
2//
3// Copyright (c) 2017-2025 The Regents of the University of Michigan and DFT-FE
4// authors.
5//
6// This file is part of the DFT-FE code.
7//
8// The DFT-FE code is free software; you can use it, redistribute
9// it, and/or modify it under the terms of the GNU Lesser General
10// Public License as published by the Free Software Foundation; either
11// version 2.1 of the License, or (at your option) any later version.
12// The full text of the license can be found in the file LICENSE at
13// the top level of the DFT-FE distribution.
14//
15// ---------------------------------------------------------------------
16
17/*
18 * @author Ian C. Lin., Sambit Das
19 */
20#ifndef dftefeDeviceKernelLauncherHelpers_h
21#define dftefeDeviceKernelLauncherHelpers_h
22
23#include <tuple>
24
25#ifdef DFTEFE_WITH_DEVICE
26# ifdef DFTEFE_WITH_DEVICE_NVIDIA
27namespace dftefe
28{
29 namespace utils
30 {
31 static const int DEVICE_WARP_SIZE = 32;
32 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
33 static const int DEVICE_BLOCK_SIZE = 256;
34
35 } // namespace utils
36} // namespace dftefe
37
38# elif DFTEFE_WITH_DEVICE_AMD
39
40namespace dftefe
41{
42 namespace utils
43 {
44 static const int DEVICE_WARP_SIZE = 64;
45 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
46 static const int DEVICE_BLOCK_SIZE = 512;
47
48 } // namespace utils
49} // namespace dftefe
50
51# elif DFTEFE_WITH_DEVICE_INTEL
52
53namespace dftefe
54{
55 namespace utils
56 {
57 static const int DEVICE_WARP_SIZE = 32;
58 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
59 static const int DEVICE_BLOCK_SIZE = 256;
60
61 } // namespace utils
62} // namespace dftefe
63
64# endif
65# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
66# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
67 do \
68 { \
69 kernel<<<grid, block, 0, stream>>>(__VA_ARGS__); \
70 } \
71 while (0)
72# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
73# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
74 do \
75 { \
76 hipLaunchKernelGGL( \
77 HIP_KERNEL_NAME(kernel), grid, block, 0, stream, __VA_ARGS__); \
78 } \
79 while (0)
80# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
81 // __VA_ARGS__ is evaluated eagerly via std::make_tuple, then unpacked back
82 // into a plain parameter pack via std::apply *on the host*, before the
83 // kernel lambda is created. The kernel lambda itself ends up an ordinary
84 // capture-by-value lambda over N plain values -- never `this`, a member,
85 // or a MemoryStorage object that a call site's argument expression (e.g.
86 // `d_x`, `storage.data()`) would otherwise require capturing in order to
87 // re-evaluate on the device. Crucially, no std::apply/std::tuple survives
88 // *inside* the kernel body or its capture list: unpacking inside the
89 // kernel (rather than before it's built) compiles and links fine but can
90 // silently break AOT (-fsycl-targets=spir64_gen) kernel enumeration, so
91 // the SYCL runtime can't find the device binary for it at launch time.
92# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
93 std::apply( \
94 [&](auto &&...dftefe_launch_kernel_args_) { \
95 dftefe::utils::queueRegistry.find(stream)->second.parallel_for( \
96 sycl::nd_range<1>((grid) * (block), block), \
97 [=](sycl::nd_item<1> ind) { \
98 kernel(ind, dftefe_launch_kernel_args_...); \
99 }); \
100 }, \
101 std::make_tuple(__VA_ARGS__))
102# else
103# error \
104 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
105# endif
106
107# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
108# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
109 kernel, grid, block, smemtype, smemcount, stream, ...) \
110 do \
111 { \
112 kernel<<<grid, block, smemcount * sizeof(smemtype), stream>>>( \
113 __VA_ARGS__); \
114 } \
115 while (0)
116# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
117# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
118 kernel, grid, block, smemtype, smemcount, stream, ...) \
119 do \
120 { \
121 hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel), \
122 grid, \
123 block, \
124 smemcount * sizeof(smemtype), \
125 stream, \
126 __VA_ARGS__); \
127 } \
128 while (0)
129# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
130# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
131 kernel, grid, block, smemtype, smemcount, stream, ...) \
132 std::apply( \
133 [&](auto &&...dftefe_launch_kernel_args_) { \
134 dftefe::utils::queueRegistry.find(stream)->second.submit( \
135 [=](sycl::handler &cgh) { \
136 sycl::local_accessor<smemtype, 1> SMem_acc(smemcount, cgh); \
137 cgh.parallel_for(sycl::nd_range<1>((grid) * (block), block), \
138 [=](sycl::nd_item<1> ind) { \
139 kernel(ind, \
140 SMem_acc.get_pointer(), \
141 dftefe_launch_kernel_args_...); \
142 }); \
143 }); \
144 }, \
145 std::make_tuple(__VA_ARGS__))
146# else
147# error \
148 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
149# endif
150
151# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
152# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
153 kernel, grid, block, smemtype, smemcount, stream, ...) \
154 do \
155 { \
156 kernel<<<grid, block, 0, stream>>>(__VA_ARGS__); \
157 } \
158 while (0)
159# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
160# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
161 kernel, grid, block, smemtype, smemcount, stream, ...) \
162 do \
163 { \
164 hipLaunchKernelGGL( \
165 HIP_KERNEL_NAME(kernel), grid, block, 0, stream, __VA_ARGS__); \
166 } \
167 while (0)
168# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
169# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
170 kernel, grid, block, smemtype, smemcount, stream, ...) \
171 std::apply( \
172 [&](auto &&...dftefe_launch_kernel_args_) { \
173 dftefe::utils::queueRegistry.find(stream)->second.submit( \
174 [=](sycl::handler &cgh) { \
175 sycl::local_accessor<smemtype, 1> SMem_acc(smemcount, cgh); \
176 cgh.parallel_for(sycl::nd_range<1>((grid) * (block), block), \
177 [=](sycl::nd_item<1> ind) { \
178 kernel(ind, \
179 SMem_acc.get_pointer(), \
180 dftefe_launch_kernel_args_...); \
181 }); \
182 }); \
183 }, \
184 std::make_tuple(__VA_ARGS__))
185# else
186# error \
187 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
188# endif
189
190
191# define DFTEFE_KERNEL_ARGUMENT(...) __VA_ARGS__
192
193
194# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
195 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
196# define DFTEFE_CREATE_KERNEL(RET, NAME, BODY, ...) \
197 __global__ RET NAME(__VA_ARGS__) \
198 { \
199 const size_type globalThreadId = \
200 blockIdx.x * blockDim.x + threadIdx.x; \
201 const size_type nThreadsPerBlock = blockDim.x; \
202 const size_type nThreadBlock = gridDim.x; \
203 BODY \
204 }
205# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
206# define DFTEFE_CREATE_KERNEL(RET, NAME, BODY, ...) \
207 RET NAME(sycl::nd_item<1> ind, __VA_ARGS__) \
208 { \
209 const size_type globalThreadId = ind.get_global_id(0); \
210 const size_type nThreadsPerBlock = ind.get_local_range(0); \
211 const size_type nThreadBlock = ind.get_group_range(0); \
212 BODY \
213 }
214# else
215# error \
216 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
217# endif
218
219# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
220 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
221# define DFTEFE_CREATE_KERNEL_SMEM_D(SMEMTYPE, RET, NAME, BODY, ...) \
222 __global__ RET NAME(__VA_ARGS__) \
223 { \
224 extern __shared__ SMEMTYPE smem[]; \
225 const size_type globalThreadId = \
226 blockIdx.x * blockDim.x + threadIdx.x; \
227 const size_type threadId = threadIdx.x; \
228 const size_type blockId = blockIdx.x; \
229 const size_type nThreadsPerBlock = blockDim.x; \
230 const size_type nThreadBlock = gridDim.x; \
231 BODY \
232 }
233# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
234# define DFTEFE_CREATE_KERNEL_SMEM_D(SMEMTYPE, RET, NAME, BODY, ...) \
235 RET NAME(sycl::nd_item<1> ind, SMEMTYPE *smem, __VA_ARGS__) \
236 { \
237 const size_type globalThreadId = ind.get_global_id(0); \
238 const size_type threadId = ind.get_local_id(0); \
239 const size_type blockId = ind.get_group(0); \
240 const size_type nThreadsPerBlock = ind.get_local_range(0); \
241 const size_type nThreadBlock = ind.get_group_range(0); \
242 BODY \
243 }
244# else
245# error \
246 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
247# endif
248
249# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
250 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
251# define DFTEFE_CREATE_KERNEL_SMEM_S( \
252 SMEMTYPE, SMEMCOUNT, RET, NAME, BODY, ...) \
253 __global__ RET NAME(__VA_ARGS__) \
254 { \
255 __shared__ SMEMTYPE smem[SMEMCOUNT]; \
256 const size_type globalThreadId = \
257 blockIdx.x * blockDim.x + threadIdx.x; \
258 const size_type threadId = threadIdx.x; \
259 const size_type blockId = blockIdx.x; \
260 const size_type nThreadsPerBlock = blockDim.x; \
261 const size_type nThreadBlock = gridDim.x; \
262 BODY \
263 }
264# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
265# define DFTEFE_CREATE_KERNEL_SMEM_S( \
266 SMEMTYPE, SMEMCOUNT, RET, NAME, BODY, ...) \
267 RET NAME(sycl::nd_item<1> ind, SMEMTYPE *smem, __VA_ARGS__) \
268 { \
269 const size_type globalThreadId = ind.get_global_id(0); \
270 const size_type threadId = ind.get_local_id(0); \
271 const size_type blockId = ind.get_group(0); \
272 const size_type nThreadsPerBlock = ind.get_local_range(0); \
273 const size_type nThreadBlock = ind.get_group_range(0); \
274 BODY \
275 }
276# else
277# error \
278 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
279# endif
280
281
282# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
283 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
284# define SYNCTHREADS __syncthreads()
285# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
286# define SYNCTHREADS sycl::group_barrier(ind.get_group());
287# else
288# error \
289 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
290# endif
291
292# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
293 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
294
295# define DFTEFE_DEVICE __device__
296# define DFTEFE_HOST_DEVICE __host__ __device__
297# define DFTEFE_FORCEINLINE __forceinline__
298
299# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
300
301// SYCL: no host/device qualifiers
302# define DFTEFE_DEVICE
303# define DFTEFE_HOST_DEVICE
304# if defined(__clang__)
305# define DFTEFE_FORCEINLINE inline __attribute__((always_inline))
306# else
307# define DFTEFE_FORCEINLINE inline
308# endif
309
310# else
311# error "DFTEFE_WITH_DEVICE is set but no backend (CUDA/HIP/SYCL) defined"
312# endif
313
314#else
315
316# define DFTEFE_HOST_DEVICE
317# define DFTEFE_FORCEINLINE inline
318
319#endif // DFTEFE_WITH_DEVICE
320
321#define DFTEFE_HOST_DEVICE_FUNC DFTEFE_FORCEINLINE DFTEFE_HOST_DEVICE
322
323#endif // dftefeDeviceKernelLauncherHelpers_h
dftefe
dealii includes
Definition: AtomFieldDataSpherical.cpp:31