Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
coef.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2022-2026, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34
35
36#ifdef __APPLE__
37#include <OpenCL/cl.h>
38#else
39#include <CL/cl.h>
40#endif
41
42#include <stdio.h>
43#include <stdlib.h>
44#include <math.h>
45#include <device/device_config.h>
46#include <device/opencl/jit.h>
47#include <device/opencl/prgm_lib.h>
48#include <device/opencl/check.h>
49
50#include "coef_kernel.cl.h"
51
55void opencl_coef_generate_geo(void *G11, void *G12, void *G13,
56 void *G22, void *G23, void *G33,
57 void *drdx, void *drdy, void *drdz,
58 void *dsdx, void *dsdy, void *dsdz,
59 void *dtdx, void *dtdy, void *dtdz,
60 void *jacinv, void *w3, int *nel,
61 int *lx, int *gdim) {
62
63 cl_int err;
64 int i;
65 if (coef_program == NULL)
66 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
67
68 const size_t global_item_size = 256 * (*nel);
69 const size_t local_item_size = 256;
70
71#define STR(X) #X
72#define GEO_CASE(LX) \
73 case LX: \
74 { \
75 cl_kernel kernel = clCreateKernel(coef_program, \
76 STR(coef_generate_geo_kernel_lx##LX), &err); \
77 CL_CHECK(err); \
78 \
79 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &G11)); \
80 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &G12)); \
81 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &G13)); \
82 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &G22)); \
83 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &G23)); \
84 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &G33)); \
85 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &drdx)); \
86 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &drdy)); \
87 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &drdz)); \
88 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dsdx)); \
89 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dsdy)); \
90 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dsdz)); \
91 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dtdx)); \
92 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &dtdy)); \
93 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dtdz)); \
94 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &jacinv)); \
95 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &w3)); \
96 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), gdim)); \
97 \
98 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
99 kernel, 1, NULL, &global_item_size, \
100 &local_item_size, 0, NULL, NULL)); \
101 CL_CHECK(clReleaseKernel(kernel)); \
102 } \
103 break
104
105 switch(*lx) {
106 GEO_CASE(2);
107 GEO_CASE(3);
108 GEO_CASE(4);
109 GEO_CASE(5);
110 GEO_CASE(6);
111 GEO_CASE(7);
112 GEO_CASE(8);
113 GEO_CASE(9);
114 GEO_CASE(10);
115 GEO_CASE(11);
116 GEO_CASE(12);
117 GEO_CASE(13);
118 GEO_CASE(14);
119 GEO_CASE(15);
120 GEO_CASE(16);
121 default:
122 {
123 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
124 exit(1);
125 }
126 }
127}
128
132void opencl_coef_generate_dxyzdrst(void *drdx, void *drdy, void *drdz,
133 void *dsdx, void *dsdy, void *dsdz,
134 void *dtdx, void *dtdy, void *dtdz,
135 void *dxdr, void *dydr, void *dzdr,
136 void *dxds, void *dyds, void *dzds,
137 void *dxdt, void *dydt, void *dzdt,
138 void *dx, void *dy, void *dz,
139 void *x, void *y, void *z,
140 void *jacinv, void *jac,
141 int *lx, int *nel) {
142
143 cl_int err;
144 int i;
145 if (coef_program == NULL)
146 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
147
148 const int n = (*nel) * (*lx) * (*lx) * (*lx);
149 const size_t global_item_size_dxyz = 256 * (*nel);
150 const size_t global_item_size_drst = 256 * n;
151 const size_t local_item_size = 256;
152
153#define STR(X) #X
154#define DXYZDRST_CASE(LX) \
155 case LX: \
156 { \
157 cl_kernel kernel = clCreateKernel(coef_program, \
158 STR(coef_generate_dxyz_kernel_lx##LX), &err); \
159 CL_CHECK(err); \
160 \
161 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &dxdr)); \
162 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &dydr)); \
163 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dzdr)); \
164 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dxds)); \
165 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dyds)); \
166 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dzds)); \
167 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dxdt)); \
168 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dydt)); \
169 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dzdt)); \
170 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dx)); \
171 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dy)); \
172 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dz)); \
173 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &x)); \
174 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &y)); \
175 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &z)); \
176 \
177 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
178 kernel, 1, NULL, &global_item_size_dxyz, \
179 &local_item_size, 0, NULL, NULL)); \
180 CL_CHECK(clReleaseKernel(kernel)); \
181 } \
182 break
183
184 switch(*lx) {
185 DXYZDRST_CASE(2);
186 DXYZDRST_CASE(3);
187 DXYZDRST_CASE(4);
188 DXYZDRST_CASE(5);
189 DXYZDRST_CASE(6);
190 DXYZDRST_CASE(7);
191 DXYZDRST_CASE(8);
192 DXYZDRST_CASE(9);
193 DXYZDRST_CASE(10);
194 DXYZDRST_CASE(11);
195 DXYZDRST_CASE(12);
196 DXYZDRST_CASE(13);
197 DXYZDRST_CASE(14);
198 DXYZDRST_CASE(15);
199 DXYZDRST_CASE(16);
200 default:
201 {
202 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
203 exit(1);
204 }
205 }
206
207 cl_kernel kernel = clCreateKernel(coef_program,
208 "coef_generate_drst_kernel", &err);
209 CL_CHECK(err);
210
211 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &jac));
212 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &jacinv));
213 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &drdx));
214 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &drdy));
215 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &drdz));
216 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dsdx));
217 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dsdy));
218 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dsdz));
219 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dtdx));
220 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dtdy));
221 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dtdz));
222 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dxdr));
223 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dydr));
224 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &dzdr));
225 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dxds));
226 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &dyds));
227 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &dzds));
228 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_mem), (void *) &dxdt));
229 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_mem), (void *) &dydt));
230 CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_mem), (void *) &dzdt));
231 CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &n));
232
233 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue,
234 kernel, 1, NULL, &global_item_size_drst,
235 &local_item_size, 0, NULL, NULL));
236 CL_CHECK(clReleaseKernel(kernel));
237}
238
242void opencl_coef_generate_mass(void *B, void *Binv, void *jac,
243 void *w3, int *lxyz, int *nel) {
244 cl_int err;
245 int i;
246 if (coef_program == NULL)
247 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
248
249 const size_t global_item_size = 256 * (*nel);
250 const size_t local_item_size = 256;
251
252 cl_kernel kernel = clCreateKernel(coef_program,
253 "coef_generate_mass_kernel", &err);
254 CL_CHECK(err);
255
256 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &B));
257 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &Binv));
258 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &jac));
259 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &w3));
260 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), lxyz));
261 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), nel));
262 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue,
263 kernel, 1, NULL, &global_item_size,
264 &local_item_size, 0, NULL, NULL));
265 CL_CHECK(clReleaseKernel(kernel));
266}
267
268
272void opencl_coef_generate_area_and_normal(void *area,
273 void *nx, void *ny, void *nz,
274 void *dxdr, void *dydr, void *dzdr,
275 void *dxds, void *dyds, void *dzds,
276 void *dxdt, void *dydt, void *dzdt,
277 void *wx, void *wy, void *wz,
278 int *lx, int *nel, real eps) {
279
280 cl_int err;
281 if (coef_program == NULL)
282 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
283
284 const size_t global_item_size = 256 * (*nel);
285 const size_t local_item_size = 256;
286
287#define STR(X) #X
288#define AREA_CASE(LX) \
289 case LX: \
290 { \
291 cl_kernel kernel = \
292 clCreateKernel(coef_program, \
293 STR(coef_generate_area_and_normal_kernel_lx##LX), &err); \
294 CL_CHECK(err); \
295 \
296 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &area)); \
297 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &nx)); \
298 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &ny)); \
299 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &nz)); \
300 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dxdr)); \
301 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dydr)); \
302 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dzdr)); \
303 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dxds)); \
304 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dyds)); \
305 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dzds)); \
306 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dxdt)); \
307 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dydt)); \
308 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dzdt)); \
309 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &wx)); \
310 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &wy)); \
311 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &wz)); \
312 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(real), &eps)); \
313 \
314 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
315 kernel, 1, NULL, &global_item_size, \
316 &local_item_size, 0, NULL, NULL)); \
317 CL_CHECK(clReleaseKernel(kernel)); \
318 } \
319 break
320
321 switch(*lx) {
322 AREA_CASE(2);
323 AREA_CASE(3);
324 AREA_CASE(4);
325 AREA_CASE(5);
326 AREA_CASE(6);
327 AREA_CASE(7);
328 AREA_CASE(8);
329 AREA_CASE(9);
330 AREA_CASE(10);
331 AREA_CASE(11);
332 AREA_CASE(12);
333 AREA_CASE(13);
334 AREA_CASE(14);
335 AREA_CASE(15);
336 AREA_CASE(16);
337 default:
338 {
339 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
340 exit(1);
341 }
342 }
343}
344
348void opencl_coef_get_normal(void *normal_x, void *normal_y, void *normal_z,
349 void *nx, void *ny, void *nz,
350 void *i_idx, void *j_idx, void *k_idx,
351 void *e_idx, void *facet, int *lx, int *n) {
352
353 if (*n <= 0) {
354 return;
355 }
356
357 cl_int err;
358 if (coef_program == NULL)
359 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
360
361 const size_t local_item_size = 256;
362 const size_t global_item_size = local_item_size *
363 (((size_t) *n + local_item_size - 1) / local_item_size);
364
365 cl_kernel kernel = clCreateKernel(coef_program, "coef_get_normal_kernel",
366 &err);
367 CL_CHECK(err);
368
369 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &normal_x));
370 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &normal_y));
371 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &normal_z));
372 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &nx));
373 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &ny));
374 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &nz));
375 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &i_idx));
376 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &j_idx));
377 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &k_idx));
378 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &e_idx));
379 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &facet));
380 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), lx));
381 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), n));
382
383 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue,
384 kernel, 1, NULL, &global_item_size,
385 &local_item_size, 0, NULL, NULL));
386 CL_CHECK(clReleaseKernel(kernel));
387}
GEO_CASE
#define GEO_CASE(LX)
AREA_CASE
#define AREA_CASE(LX)
opencl_coef_generate_mass
void opencl_coef_generate_mass(void *B, void *Binv, void *jac, void *w3, int *lxyz, int *nel)
Definition coef.c:242
opencl_coef_get_normal
void opencl_coef_get_normal(void *normal_x, void *normal_y, void *normal_z, void *nx, void *ny, void *nz, void *i_idx, void *j_idx, void *k_idx, void *e_idx, void *facet, int *lx, int *n)
Definition coef.c:348
opencl_coef_generate_geo
void opencl_coef_generate_geo(void *G11, void *G12, void *G13, void *G22, void *G23, void *G33, void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *jacinv, void *w3, int *nel, int *lx, int *gdim)
Definition coef.c:55
opencl_coef_generate_dxyzdrst
void opencl_coef_generate_dxyzdrst(void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *dxdr, void *dydr, void *dzdr, void *dxds, void *dyds, void *dzds, void *dxdt, void *dydt, void *dzdt, void *dx, void *dy, void *dz, void *x, void *y, void *z, void *jacinv, void *jac, int *lx, int *nel)
Definition coef.c:132
opencl_coef_generate_area_and_normal
void opencl_coef_generate_area_and_normal(void *area, void *nx, void *ny, void *nz, void *dxdr, void *dydr, void *dzdr, void *dxds, void *dyds, void *dzds, void *dxdt, void *dydt, void *dzdt, void *wx, void *wy, void *wz, int *lx, int *nel, real eps)
Definition coef.c:272
DXYZDRST_CASE
#define DXYZDRST_CASE(LX)
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
drdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdy
Definition ax_helm_full_kernel.h:54
drdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdz
Definition ax_helm_full_kernel.h:55
dsdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdz
Definition ax_helm_full_kernel.h:58
dsdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdy
Definition ax_helm_full_kernel.h:57
i
const int i
Definition ax_helm_full_kernel.h:95
dtdy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdy
Definition ax_helm_full_kernel.h:60
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
drdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdx
Definition ax_helm_full_kernel.h:53
dtdz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdz
Definition ax_helm_full_kernel.h:61
dsdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdx
Definition ax_helm_full_kernel.h:56
dtdx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdx
Definition ax_helm_full_kernel.h:59
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition ax_helm_full_kernel.h:51
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
jacinv
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ jacinv
Definition ax_helm_full_kernel.h:62
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:108
w3
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
Definition cdtp_kernel.h:115
device_config.h
real
double real
Definition device_config.h:12
opencl_kernel_jit
void opencl_kernel_jit(const char *kernel, cl_program *program)
Definition jit.c:50
CL_CHECK
#define CL_CHECK(err)
Definition check.h:12
prgm_lib.h
coef_program
void * coef_program