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#ifdef DFTEFE_WITH_DEVICE
24# ifdef DFTEFE_WITH_DEVICE_NVIDIA
25namespace dftefe
26{
27 namespace utils
28 {
29 static const int DEVICE_WARP_SIZE = 32;
30 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
31 static const int DEVICE_BLOCK_SIZE = 256;
32
33 } // namespace utils
34} // namespace dftefe
35
36# elif DFTEFE_WITH_DEVICE_AMD
37
38namespace dftefe
39{
40 namespace utils
41 {
42 static const int DEVICE_WARP_SIZE = 64;
43 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
44 static const int DEVICE_BLOCK_SIZE = 512;
45
46 } // namespace utils
47} // namespace dftefe
48
49# elif DFTEFE_WITH_DEVICE_INTEL
50
51namespace dftefe
52{
53 namespace utils
54 {
55 static const int DEVICE_WARP_SIZE = 32;
56 static const int DEVICE_MAX_BLOCK_SIZE = 1024;
57 static const int DEVICE_BLOCK_SIZE = 256;
58
59 } // namespace utils
60} // namespace dftefe
61
62# endif
63# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
64# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
65 do \
66 { \
67 kernel<<<grid, block, 0, stream>>>(__VA_ARGS__); \
68 } \
69 while (0)
70# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
71# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
72 do \
73 { \
74 hipLaunchKernelGGL( \
75 HIP_KERNEL_NAME(kernel), grid, block, 0, stream, __VA_ARGS__); \
76 } \
77 while (0)
78# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
79# define DFTEFE_LAUNCH_KERNEL(kernel, grid, block, stream, ...) \
80 do \
81 { \
82 dftefe::utils::queueRegistry.find(stream)->second.parallel_for( \
83 sycl::nd_range<1>((grid) * (block), block), \
84 [=](sycl::nd_item<1> ind) { kernel(ind, __VA_ARGS__); }); \
85 } \
86 while (0)
87# else
88# error \
89 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
90# endif
91
92# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
93# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
94 kernel, grid, block, smemtype, smemcount, stream, ...) \
95 do \
96 { \
97 kernel<<<grid, block, smemcount * sizeof(smemtype), stream>>>( \
98 __VA_ARGS__); \
99 } \
100 while (0)
101# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
102# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
103 kernel, grid, block, smemtype, smemcount, stream, ...) \
104 do \
105 { \
106 hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel), \
107 grid, \
108 block, \
109 smemcount * sizeof(smemtype), \
110 stream, \
111 __VA_ARGS__); \
112 } \
113 while (0)
114# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
115# define DFTEFE_LAUNCH_KERNEL_SMEM_D( \
116 kernel, grid, block, smemtype, smemcount, stream, ...) \
117 do \
118 { \
119 dftefe::utils::queueRegistry.find(stream)->second.submit( \
120 [=](sycl::handler &cgh) { \
121 sycl::local_accessor<smemtype, 1> SMem_acc(smemcount, cgh); \
122 cgh.parallel_for(sycl::nd_range<1>((grid) * (block), block), \
123 [=](sycl::nd_item<1> ind) { \
124 kernel(ind, \
125 SMem_acc.get_pointer(), \
126 __VA_ARGS__); \
127 }); \
128 }); \
129 } \
130 while (0)
131# else
132# error \
133 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
134# endif
135
136# ifdef DFTEFE_WITH_DEVICE_LANG_CUDA
137# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
138 kernel, grid, block, smemtype, smemcount, stream, ...) \
139 do \
140 { \
141 kernel<<<grid, block, 0, stream>>>(__VA_ARGS__); \
142 } \
143 while (0)
144# elif defined(DFTEFE_WITH_DEVICE_LANG_HIP)
145# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
146 kernel, grid, block, smemtype, smemcount, stream, ...) \
147 do \
148 { \
149 hipLaunchKernelGGL( \
150 HIP_KERNEL_NAME(kernel), grid, block, 0, stream, __VA_ARGS__); \
151 } \
152 while (0)
153# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
154# define DFTEFE_LAUNCH_KERNEL_SMEM_S( \
155 kernel, grid, block, smemtype, smemcount, stream, ...) \
156 do \
157 { \
158 dftefe::utils::queueRegistry.find(stream)->second.submit( \
159 [=](sycl::handler &cgh) { \
160 sycl::local_accessor<smemtype, 1> SMem_acc(smemcount, cgh); \
161 cgh.parallel_for(sycl::nd_range<1>((grid) * (block), block), \
162 [=](sycl::nd_item<1> ind) { \
163 kernel(ind, \
164 SMem_acc.get_pointer(), \
165 __VA_ARGS__); \
166 }); \
167 }); \
168 } \
169 while (0)
170# else
171# error \
172 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
173# endif
174
175
176# define DFTEFE_KERNEL_ARGUMENT(...) __VA_ARGS__
177
178
179# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
180 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
181# define DFTEFE_CREATE_KERNEL(RET, NAME, BODY, ...) \
182 __global__ RET NAME(__VA_ARGS__) \
183 { \
184 const size_type globalThreadId = \
185 blockIdx.x * blockDim.x + threadIdx.x; \
186 const size_type nThreadsPerBlock = blockDim.x; \
187 const size_type nThreadBlock = gridDim.x; \
188 BODY \
189 }
190# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
191# define DFTEFE_CREATE_KERNEL(RET, NAME, BODY, ...) \
192 RET NAME(sycl::nd_item<1> ind, __VA_ARGS__) \
193 { \
194 const size_type globalThreadId = ind.get_global_id(0); \
195 const size_type nThreadsPerBlock = ind.get_local_range(0); \
196 const size_type nThreadBlock = ind.get_group_range(0); \
197 BODY \
198 }
199# else
200# error \
201 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
202# endif
203
204# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
205 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
206# define DFTEFE_CREATE_KERNEL_SMEM_D(SMEMTYPE, RET, NAME, BODY, ...) \
207 __global__ RET NAME(__VA_ARGS__) \
208 { \
209 extern __shared__ SMEMTYPE smem[]; \
210 const size_type globalThreadId = \
211 blockIdx.x * blockDim.x + threadIdx.x; \
212 const size_type threadId = threadIdx.x; \
213 const size_type blockId = blockIdx.x; \
214 const size_type nThreadsPerBlock = blockDim.x; \
215 const size_type nThreadBlock = gridDim.x; \
216 BODY \
217 }
218# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
219# define DFTEFE_CREATE_KERNEL_SMEM_D(SMEMTYPE, RET, NAME, BODY, ...) \
220 RET NAME(sycl::nd_item<1> ind, SMEMTYPE *smem, __VA_ARGS__) \
221 { \
222 const size_type globalThreadId = ind.get_global_id(0); \
223 const size_type threadId = ind.get_local_id(0); \
224 const size_type blockId = ind.get_group(0); \
225 const size_type nThreadsPerBlock = ind.get_local_range(0); \
226 const size_type nThreadBlock = ind.get_group_range(0); \
227 BODY \
228 }
229# else
230# error \
231 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
232# endif
233
234# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
235 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
236# define DFTEFE_CREATE_KERNEL_SMEM_S( \
237 SMEMTYPE, SMEMCOUNT, RET, NAME, BODY, ...) \
238 __global__ RET NAME(__VA_ARGS__) \
239 { \
240 __shared__ SMEMTYPE smem[SMEMCOUNT]; \
241 const size_type globalThreadId = \
242 blockIdx.x * blockDim.x + threadIdx.x; \
243 const size_type threadId = threadIdx.x; \
244 const size_type blockId = blockIdx.x; \
245 const size_type nThreadsPerBlock = blockDim.x; \
246 const size_type nThreadBlock = gridDim.x; \
247 BODY \
248 }
249# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
250# define DFTEFE_CREATE_KERNEL_SMEM_S( \
251 SMEMTYPE, SMEMCOUNT, RET, NAME, BODY, ...) \
252 RET NAME(sycl::nd_item<1> ind, SMEMTYPE *smem, __VA_ARGS__) \
253 { \
254 const size_type globalThreadId = ind.get_global_id(0); \
255 const size_type threadId = ind.get_local_id(0); \
256 const size_type blockId = ind.get_group(0); \
257 const size_type nThreadsPerBlock = ind.get_local_range(0); \
258 const size_type nThreadBlock = ind.get_group_range(0); \
259 BODY \
260 }
261# else
262# error \
263 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
264# endif
265
266
267# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
268 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
269# define SYNCTHREADS __syncthreads()
270# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
271# define SYNCTHREADS sycl::group_barrier(ind.get_group());
272# else
273# error \
274 "No device backend defined (DFTEFE_WITH_DEVICE_LANG_CUDA or DFTEFE_WITH_DEVICE_LANG_HIP or DFTEFE_WITH_DEVICE_LANG_SYCL)"
275# endif
276
277# if defined(DFTEFE_WITH_DEVICE_LANG_CUDA) || \
278 defined(DFTEFE_WITH_DEVICE_LANG_HIP)
279
280# define DFTEFE_DEVICE __device__
281# define DFTEFE_HOST_DEVICE __host__ __device__
282# define DFTEFE_FORCEINLINE __forceinline__
283
284# elif defined(DFTEFE_WITH_DEVICE_LANG_SYCL)
285
286// SYCL: no host/device qualifiers
287# define DFTEFE_DEVICE
288# define DFTEFE_HOST_DEVICE
289# if defined(__clang__)
290# define DFTEFE_FORCEINLINE inline __attribute__((always_inline))
291# else
292# define DFTEFE_FORCEINLINE inline
293# endif
294
295# else
296# error "DFTEFE_WITH_DEVICE is set but no backend (CUDA/HIP/SYCL) defined"
297# endif
298
299#else
300
301# define DFTEFE_HOST_DEVICE
302# define DFTEFE_FORCEINLINE inline
303
304#endif // DFTEFE_WITH_DEVICE
305
306#define DFTEFE_HOST_DEVICE_FUNC DFTEFE_FORCEINLINE DFTEFE_HOST_DEVICE
307
308#endif // dftefeDeviceKernelLauncherHelpers_h
dftefe
dealii includes
Definition: AtomFieldDataSpherical.cpp:31