Neko 1.99.2
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
rhs_maker.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2022, 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#ifdef __APPLE__
36#include <OpenCL/cl.h>
37#else
38#include <CL/cl.h>
39#endif
40
41#include <stdio.h>
42#include <device/device_config.h>
43#include <device/opencl/jit.h>
44#include <device/opencl/prgm_lib.h>
45#include <device/opencl/check.h>
46
47#include "rhs_maker_kernel.cl.h"
48
49void rhs_maker_sumab_opencl(void *u, void *v, void *w,
50 void *uu, void *vv, void *ww,
51 void *ulag1, void *ulag2, void *vlag1,
52 void *vlag2, void *wlag1, void *wlag2,
53 real *ext1, real *ext2, real *ext3, int *nab, int *n) {
54 cl_int err;
55
56 if (rhs_maker_program == NULL)
57 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
58
59 cl_kernel kernel = clCreateKernel(rhs_maker_program, "sumab_kernel", &err);
60 CL_CHECK(err);
61
62 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &u));
63 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &v));
64 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &w));
65 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &uu));
66 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &vv));
67 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &ww));
68 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &ulag1));
69 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &ulag2));
70 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &vlag1));
71 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &vlag2));
72 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &wlag1));
73 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &wlag2));
74 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(real), ext1));
75 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(real), ext2));
76 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(real), ext3));
77 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), nab));
78 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), n));
79
80 const int nb = ((*n) + 256 - 1) / 256;
81 const size_t global_item_size = 256 * nb;
82 const size_t local_item_size = 256;
83
84 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
85 NULL, &global_item_size, &local_item_size,
86 0, NULL, NULL));
87 CL_CHECK(clReleaseKernel(kernel));
88}
89
90void rhs_maker_ext_opencl(void *abx1, void *aby1, void *abz1,
91 void *abx2, void *aby2, void *abz2,
92 void *bfx, void *bfy, void *bfz,
93 real *rho, real *ext1, real *ext2, real *ext3, int *n) {
94 cl_int err;
95
96 if (rhs_maker_program == NULL)
97 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
98
99 cl_kernel kernel = clCreateKernel(rhs_maker_program, "makeext_kernel", &err);
100 CL_CHECK(err);
101
102 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &abx1));
103 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &aby1));
104 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &abz1));
105 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &abx2));
106 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &aby2));
107 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &abz2));
108 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &bfx));
109 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &bfy));
110 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &bfz));
111 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(real), rho));
112 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(real), ext1));
113 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(real), ext2));
114 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(real), ext3));
115 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), n));
116
117 const int nb = ((*n) + 256 - 1) / 256;
118 const size_t global_item_size = 256 * nb;
119 const size_t local_item_size = 256;
120
121 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
122 NULL, &global_item_size, &local_item_size,
123 0, NULL, NULL));
124
125 CL_CHECK(clReleaseKernel(kernel));
126}
127
128void scalar_rhs_maker_ext_opencl(void *fs_lag, void *fs_laglag, void *fs,
129 real *rho, real *ext1, real *ext2,
130 real *ext3, int *n) {
131 cl_int err;
132
133 if (rhs_maker_program == NULL)
134 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
135
136 cl_kernel kernel = clCreateKernel(rhs_maker_program, "scalar_makeext_kernel", &err);
137 CL_CHECK(err);
138
139 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &fs_lag));
140 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &fs_laglag));
141 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &fs));
142 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), rho));
143 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(real), ext1));
144 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(real), ext2));
145 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(real), ext3));
146 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), n));
147
148 const int nb = ((*n) + 256 - 1) / 256;
149 const size_t global_item_size = 256 * nb;
150 const size_t local_item_size = 256;
151
152 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
153 NULL, &global_item_size, &local_item_size,
154 0, NULL, NULL));
155
156 CL_CHECK(clReleaseKernel(kernel));
157}
158
159void rhs_maker_bdf_opencl(void *ulag1, void *ulag2, void *vlag1,
160 void *vlag2, void *wlag1, void *wlag2,
161 void *bfx, void *bfy, void *bfz,
162 void *u, void *v, void *w, void *B,
163 real *rho, real *dt, real *bd2,
164 real *bd3, real *bd4, int *nbd, int *n) {
165 cl_int err;
166
167 if (rhs_maker_program == NULL)
168 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
169
170 cl_kernel kernel = clCreateKernel(rhs_maker_program, "makebdf_kernel", &err);
171 CL_CHECK(err);
172
173 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &ulag1));
174 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &ulag2));
175 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &vlag1));
176 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &vlag2));
177 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &wlag1));
178 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &wlag2));
179 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &bfx));
180 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &bfy));
181 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &bfz));
182 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &u));
183 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &v));
184 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &w));
185 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &B));
186 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(real), rho));
187 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(real), dt));
188 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(real), bd2));
189 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(real), bd3));
190 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(real), bd4));
191 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), nbd));
192 CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), n));
193
194 const int nb = ((*n) + 256 - 1) / 256;
195 const size_t global_item_size = 256 * nb;
196 const size_t local_item_size = 256;
197
198 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
199 NULL, &global_item_size, &local_item_size,
200 0, NULL, NULL));
201
202 CL_CHECK(clReleaseKernel(kernel));
203}
204
205void scalar_rhs_maker_bdf_opencl(void *s_lag, void *s_laglag, void *fs,
206 void *s, void *B, real *rho, real *dt,
207 real *bd2, real *bd3, real *bd4,
208 int *nbd, int *n) {
209 cl_int err;
210
211 if (rhs_maker_program == NULL)
212 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
213
214 cl_kernel kernel = clCreateKernel(rhs_maker_program, "scalar_makebdf_kernel", &err);
215 CL_CHECK(err);
216
217 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &s_lag));
218 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &s_laglag));
219 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &fs));
220 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &s));
221 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &B));
222 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(real), rho));
223 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(real), dt));
224 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(real), bd2));
225 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(real), bd3));
226 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(real), bd4));
227 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), nbd));
228 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), n));
229
230 const int nb = ((*n) + 256 - 1) / 256;
231 const size_t global_item_size = 256 * nb;
232 const size_t local_item_size = 256;
233
234 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
235 NULL, &global_item_size, &local_item_size,
236 0, NULL, NULL));
237
238 CL_CHECK(clReleaseKernel(kernel));
239}
240
241void rhs_maker_oifs_opencl(void *phi_x, void *phi_y, void *phi_z,
242 void *bf_x, void *bf_y, void *bf_z,
243 real *rho, real *dt, int *n) {
244 cl_int err;
245
246 if (rhs_maker_program == NULL)
247 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
248
249 cl_kernel kernel = clCreateKernel(rhs_maker_program, "makeoifs_kernel", &err);
250 CL_CHECK(err);
251
252 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &phi_x));
253 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &phi_y));
254 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &phi_z));
255 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &bf_x));
256 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &bf_y));
257 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &bf_z));
258 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(real), rho));
259 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(real), dt));
260 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), n));
261
262 const int nb = ((*n) + 256 - 1) / 256;
263 const size_t global_item_size = 256 * nb;
264 const size_t local_item_size = 256;
265
266 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
267 NULL, &global_item_size, &local_item_size,
268 0, NULL, NULL));
269
270 CL_CHECK(clReleaseKernel(kernel));
271}
272
273void scalar_rhs_maker_oifs_opencl(void *phi_s, void *bf_s,
274 real *rho, real *dt, int *n) {
275 cl_int err;
276
277 if (rhs_maker_program == NULL)
278 opencl_kernel_jit(rhs_maker_kernel, (cl_program *) &rhs_maker_program);
279
280 cl_kernel kernel = clCreateKernel(rhs_maker_program, "scalar_makeoifs_kernel", &err);
281 CL_CHECK(err);
282
283 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &phi_s));
284 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &bf_s));
285 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(real), rho));
286 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), dt));
287 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), n));
288
289 const int nb = ((*n) + 256 - 1) / 256;
290 const size_t global_item_size = 256 * nb;
291 const size_t local_item_size = 256;
292
293 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
294 NULL, &global_item_size, &local_item_size,
295 0, NULL, NULL));
296
297 CL_CHECK(clReleaseKernel(kernel));
298}
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
Definition cdtp_kernel.h:109
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
rhs_maker_program
void * rhs_maker_program
scalar_rhs_maker_bdf_opencl
void scalar_rhs_maker_bdf_opencl(void *s_lag, void *s_laglag, void *fs, void *s, void *B, real *rho, real *dt, real *bd2, real *bd3, real *bd4, int *nbd, int *n)
Definition rhs_maker.c:205
rhs_maker_sumab_opencl
void rhs_maker_sumab_opencl(void *u, void *v, void *w, void *uu, void *vv, void *ww, void *ulag1, void *ulag2, void *vlag1, void *vlag2, void *wlag1, void *wlag2, real *ext1, real *ext2, real *ext3, int *nab, int *n)
Definition rhs_maker.c:49
scalar_rhs_maker_oifs_opencl
void scalar_rhs_maker_oifs_opencl(void *phi_s, void *bf_s, real *rho, real *dt, int *n)
Definition rhs_maker.c:273
rhs_maker_ext_opencl
void rhs_maker_ext_opencl(void *abx1, void *aby1, void *abz1, void *abx2, void *aby2, void *abz2, void *bfx, void *bfy, void *bfz, real *rho, real *ext1, real *ext2, real *ext3, int *n)
Definition rhs_maker.c:90
rhs_maker_oifs_opencl
void rhs_maker_oifs_opencl(void *phi_x, void *phi_y, void *phi_z, void *bf_x, void *bf_y, void *bf_z, real *rho, real *dt, int *n)
Definition rhs_maker.c:241
scalar_rhs_maker_ext_opencl
void scalar_rhs_maker_ext_opencl(void *fs_lag, void *fs_laglag, void *fs, real *rho, real *ext1, real *ext2, real *ext3, int *n)
Definition rhs_maker.c:128
rhs_maker_bdf_opencl
void rhs_maker_bdf_opencl(void *ulag1, void *ulag2, void *vlag1, void *vlag2, void *wlag1, void *wlag2, void *bfx, void *bfy, void *bfz, void *u, void *v, void *w, void *B, real *rho, real *dt, real *bd2, real *bd3, real *bd4, int *nbd, int *n)
Definition rhs_maker.c:159