Neko 1.99.3
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 <math.h>
44#include <device/device_config.h>
45#include <device/opencl/jit.h>
46#include <device/opencl/prgm_lib.h>
47#include <device/opencl/check.h>
48
49#include "coef_kernel.cl.h"
50
54void opencl_coef_generate_geo(void *G11, void *G12, void *G13,
55 void *G22, void *G23, void *G33,
56 void *drdx, void *drdy, void *drdz,
57 void *dsdx, void *dsdy, void *dsdz,
58 void *dtdx, void *dtdy, void *dtdz,
59 void *jacinv, void *w3, int *nel,
60 int *lx, int *gdim) {
61
62 cl_int err;
63 int i;
64 if (coef_program == NULL)
65 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
66
67 const size_t global_item_size = 256 * (*nel);
68 const size_t local_item_size = 256;
69
70#define STR(X) #X
71#define GEO_CASE(LX) \
72 case LX: \
73 { \
74 cl_kernel kernel = clCreateKernel(coef_program, \
75 STR(coef_generate_geo_kernel_lx##LX), &err); \
76 CL_CHECK(err); \
77 \
78 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &G11)); \
79 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &G12)); \
80 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &G13)); \
81 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &G22)); \
82 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &G23)); \
83 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &G33)); \
84 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &drdx)); \
85 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &drdy)); \
86 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &drdz)); \
87 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dsdx)); \
88 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dsdy)); \
89 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dsdz)); \
90 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dtdx)); \
91 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &dtdy)); \
92 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dtdz)); \
93 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &jacinv)); \
94 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &w3)); \
95 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), gdim)); \
96 \
97 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
98 kernel, 1, NULL, &global_item_size, \
99 &local_item_size, 0, NULL, NULL)); \
100 CL_CHECK(clReleaseKernel(kernel)); \
101 } \
102 break
103
104 switch(*lx) {
105 GEO_CASE(2);
106 GEO_CASE(3);
107 GEO_CASE(4);
108 GEO_CASE(5);
109 GEO_CASE(6);
110 GEO_CASE(7);
111 GEO_CASE(8);
112 GEO_CASE(9);
113 GEO_CASE(10);
114 GEO_CASE(11);
115 GEO_CASE(12);
116 GEO_CASE(13);
117 GEO_CASE(14);
118 GEO_CASE(15);
119 GEO_CASE(16);
120 }
121}
122
126void opencl_coef_generate_dxyzdrst(void *drdx, void *drdy, void *drdz,
127 void *dsdx, void *dsdy, void *dsdz,
128 void *dtdx, void *dtdy, void *dtdz,
129 void *dxdr, void *dydr, void *dzdr,
130 void *dxds, void *dyds, void *dzds,
131 void *dxdt, void *dydt, void *dzdt,
132 void *dx, void *dy, void *dz,
133 void *x, void *y, void *z,
134 void *jacinv, void *jac,
135 int *lx, int *nel) {
136
137 cl_int err;
138 int i;
139 if (coef_program == NULL)
140 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
141
142 const int n = (*nel) * (*lx) * (*lx) * (*lx);
143 const size_t global_item_size_dxyz = 256 * (*nel);
144 const size_t global_item_size_drst = 256 * n;
145 const size_t local_item_size = 256;
146
147#define STR(X) #X
148#define DXYZDRST_CASE(LX) \
149 case LX: \
150 { \
151 cl_kernel kernel = clCreateKernel(coef_program, \
152 STR(coef_generate_dxyz_kernel_lx##LX), &err); \
153 CL_CHECK(err); \
154 \
155 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &dxdr)); \
156 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &dydr)); \
157 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &dzdr)); \
158 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dxds)); \
159 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dyds)); \
160 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dzds)); \
161 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dxdt)); \
162 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dydt)); \
163 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dzdt)); \
164 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dx)); \
165 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dy)); \
166 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dz)); \
167 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &x)); \
168 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &y)); \
169 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &z)); \
170 \
171 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
172 kernel, 1, NULL, &global_item_size_dxyz, \
173 &local_item_size, 0, NULL, NULL)); \
174 CL_CHECK(clReleaseKernel(kernel)); \
175 } \
176 break
177
178 switch(*lx) {
179 DXYZDRST_CASE(2);
180 DXYZDRST_CASE(3);
181 DXYZDRST_CASE(4);
182 DXYZDRST_CASE(5);
183 DXYZDRST_CASE(6);
184 DXYZDRST_CASE(7);
185 DXYZDRST_CASE(8);
186 DXYZDRST_CASE(9);
187 DXYZDRST_CASE(10);
188 DXYZDRST_CASE(11);
189 DXYZDRST_CASE(12);
190 DXYZDRST_CASE(13);
191 DXYZDRST_CASE(14);
192 DXYZDRST_CASE(15);
193 DXYZDRST_CASE(16);
194 }
195
196 cl_kernel kernel = clCreateKernel(coef_program,
197 "coef_generate_drst_kernel", &err);
198 CL_CHECK(err);
199
200 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &jac));
201 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &jacinv));
202 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &drdx));
203 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &drdy));
204 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &drdz));
205 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dsdx));
206 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dsdy));
207 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dsdz));
208 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dtdx));
209 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dtdy));
210 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dtdz));
211 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dxdr));
212 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dydr));
213 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &dzdr));
214 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dxds));
215 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &dyds));
216 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_mem), (void *) &dzds));
217 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_mem), (void *) &dxdt));
218 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_mem), (void *) &dydt));
219 CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_mem), (void *) &dzdt));
220 CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &n));
221
222 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue,
223 kernel, 1, NULL, &global_item_size_drst,
224 &local_item_size, 0, NULL, NULL));
225 CL_CHECK(clReleaseKernel(kernel));
226}
227
231void opencl_coef_generate_mass(void *B, void *Binv, void *jac,
232 void *w3, int *lxyz, int *nel) {
233 cl_int err;
234 int i;
235 if (coef_program == NULL)
236 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
237
238 const size_t global_item_size = 256 * (*nel);
239 const size_t local_item_size = 256;
240
241 cl_kernel kernel = clCreateKernel(coef_program,
242 "coef_generate_mass_kernel", &err);
243 CL_CHECK(err);
244
245 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &B));
246 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &Binv));
247 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &jac));
248 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &w3));
249 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), lxyz));
250 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), nel));
251 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue,
252 kernel, 1, NULL, &global_item_size,
253 &local_item_size, 0, NULL, NULL));
254 CL_CHECK(clReleaseKernel(kernel));
255}
256
257
261void opencl_coef_generate_area_and_normal(void *area,
262 void *nx, void *ny, void *nz,
263 void *dxdr, void *dydr, void *dzdr,
264 void *dxds, void *dyds, void *dzds,
265 void *dxdt, void *dydt, void *dzdt,
266 void *wx, void *wy, void *wz,
267 int *lx, int *nel, real eps) {
268
269 cl_int err;
270 if (coef_program == NULL)
271 opencl_kernel_jit(coef_kernel, (cl_program *) &coef_program);
272
273 const size_t global_item_size = 256 * (*nel);
274 const size_t local_item_size = 256;
275
276#define STR(X) #X
277#define AREA_CASE(LX) \
278 case LX: \
279 { \
280 printf("%s\n", STR(coef_generate_area_and_normal_kernel_lx##LX)); \
281 cl_kernel kernel = \
282 clCreateKernel(coef_program, \
283 STR(coef_generate_area_and_normal_kernel_lx##LX), &err); \
284 CL_CHECK(err); \
285 printf("Done\n"); \
286 \
287 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &area)); \
288 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &nx)); \
289 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &ny)); \
290 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &nz)); \
291 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dxdr)); \
292 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dydr)); \
293 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dzdr)); \
294 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dxds)); \
295 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dyds)); \
296 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dzds)); \
297 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dxdt)); \
298 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dydt)); \
299 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &dzdt)); \
300 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &wx)); \
301 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &wy)); \
302 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_mem), (void *) &wz)); \
303 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(real), &eps)); \
304 \
305 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, \
306 kernel, 1, NULL, &global_item_size, \
307 &local_item_size, 0, NULL, NULL)); \
308 CL_CHECK(clReleaseKernel(kernel)); \
309 } \
310 break
311
312 switch(*lx) {
313 AREA_CASE(2);
314 AREA_CASE(3);
315 AREA_CASE(4);
316 AREA_CASE(5);
317 AREA_CASE(6);
318 AREA_CASE(7);
319 AREA_CASE(8);
320 AREA_CASE(9);
321 AREA_CASE(10);
322 AREA_CASE(11);
323 AREA_CASE(12);
324 AREA_CASE(13);
325 AREA_CASE(14);
326 AREA_CASE(15);
327 AREA_CASE(16);
328 }
329}
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:231
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:54
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:126
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:261
DXYZDRST_CASE
#define DXYZDRST_CASE(LX)
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:106
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:113
dirichlet_apply_scalar_kernel
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
Definition dirichlet_kernel.h:42
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