Neko 1.1.0
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
compressible_ops_update.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2025, 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#include "compressible_ops_update_kernel.cl.h"
47
48void opencl_update_uvw(void *u, void *v, void *w,
49 void *m_x, void *m_y, void *m_z,
50 void *rho, int n) {
51
52 cl_int err;
53
54 if (compressible_ops_update_program == NULL)
55 opencl_kernel_jit(compressible_ops_update_kernel,
56 (cl_program *) &compressible_ops_update_program);
57
58 cl_kernel kernel =
59 clCreateKernel((cl_program) compressible_ops_update_program,
60 "update_uvw_kernel", &err);
61 CL_CHECK(err);
62
63 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &u));
64 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &v));
65 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &w));
66 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_x));
67 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &m_y));
68 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &m_z));
69 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &rho));
70 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &n));
71
72 const int nb = (n + 256 - 1) / 256;
73 const size_t global_item_size = 256 * nb;
74 const size_t local_item_size = 256;
75
76 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
77 NULL, &global_item_size, &local_item_size,
78 0, NULL, NULL));
79 CL_CHECK(clReleaseKernel(kernel));
80}
81
82void opencl_update_mxyz_p_ruvw(void *m_x, void *m_y, void *m_z,
83 void *p, void *ruvw,
84 void *u, void *v, void *w, void *E,
85 void *rho, real gamma, int n) {
86 cl_int err;
87
88 if (compressible_ops_update_program == NULL)
89 opencl_kernel_jit(compressible_ops_update_kernel,
90 (cl_program *) &compressible_ops_update_program);
91
92 cl_kernel kernel =
93 clCreateKernel((cl_program) compressible_ops_update_program,
94 "update_mxyz_p_ruvw_kernel", &err);
95 CL_CHECK(err);
96
97 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &m_x));
98 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &m_y));
99 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &m_z));
100 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &p));
101 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &ruvw));
102 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &u));
103 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &v));
104 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &w));
105 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &E));
106 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &rho));
107 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(real), &gamma));
108 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &n));
109
110 const int nb = (n + 256 - 1) / 256;
111 const size_t global_item_size = 256 * nb;
112 const size_t local_item_size = 256;
113
114 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
115 NULL, &global_item_size, &local_item_size,
116 0, NULL, NULL));
117 CL_CHECK(clReleaseKernel(kernel));
118}
119
120void opencl_update_e(void *E, void *p, void *ruvw, real gamma, int n) {
121
122 cl_int err;
123
124 if (compressible_ops_update_program == NULL)
125 opencl_kernel_jit(compressible_ops_update_kernel,
126 (cl_program *) &compressible_ops_update_program);
127
128 cl_kernel kernel =
129 clCreateKernel((cl_program) compressible_ops_update_program,
130 "update_e_kernel", &err);
131 CL_CHECK(err);
132
133 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &E));
134 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &p));
135 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &ruvw));
136 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), &gamma));
137 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &n));
138
139 const int nb = (n + 256 - 1) / 256;
140 const size_t global_item_size = 256 * nb;
141 const size_t local_item_size = 256;
142
143 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
144 NULL, &global_item_size, &local_item_size,
145 0, NULL, NULL));
146 CL_CHECK(clReleaseKernel(kernel));
147}
148
149void opencl_update_temperature(void *T, void *p, void *rho, real gamma,
150 int n) {
151 cl_int err;
152
153 if (compressible_ops_update_program == NULL)
154 opencl_kernel_jit(compressible_ops_update_kernel,
155 (cl_program *) &compressible_ops_update_program);
156
157 cl_kernel kernel =
158 clCreateKernel((cl_program) compressible_ops_update_program,
159 "update_temperature_kernel", &err);
160 CL_CHECK(err);
161
162 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &T));
163 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &p));
164 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &rho));
165 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(real), &gamma));
166 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &n));
167
168 const int nb = (n + 256 - 1) / 256;
169 const size_t global_item_size = 256 * nb;
170 const size_t local_item_size = 256;
171
172 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
173 NULL, &global_item_size, &local_item_size,
174 0, NULL, NULL));
175 CL_CHECK(clReleaseKernel(kernel));
176}
177
178void opencl_ns_flux_prepare(void *div_flux, void *dissipation, void *h1,
179 void *dudx, void *dudy, void *dudz,
180 void *dvdx, void *dvdy, void *dvdz,
181 void *dwdx, void *dwdy, void *dwdz,
182 void *mu, int n) {
183 cl_int err;
184
185 if (compressible_ops_update_program == NULL)
186 opencl_kernel_jit(compressible_ops_update_kernel,
187 (cl_program *) &compressible_ops_update_program);
188
189 cl_kernel kernel =
190 clCreateKernel((cl_program) compressible_ops_update_program,
191 "ns_flux_prepare_kernel", &err);
192 CL_CHECK(err);
193
194 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &div_flux));
195 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &dissipation));
196 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &h1));
197 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &dudx));
198 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &dudy));
199 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &dudz));
200 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &dvdx));
201 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &dvdy));
202 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &dvdz));
203 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &dwdx));
204 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &dwdy));
205 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &dwdz));
206 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &mu));
207 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &n));
208
209 const int nb = (n + 256 - 1) / 256;
210 const size_t global_item_size = 256 * nb;
211 const size_t local_item_size = 256;
212
213 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
214 NULL, &global_item_size, &local_item_size,
215 0, NULL, NULL));
216 CL_CHECK(clReleaseKernel(kernel));
217}
218
219void opencl_ns_flux_finalize(void *visc_m_x, void *visc_m_y,
220 void *visc_m_z, void *visc_E,
221 void *f_x, void *f_y, void *f_z,
222 void *opgrad_x, void *opgrad_y, void *opgrad_z,
223 void *u, void *v, void *w, void *B,
224 void *dissipation, int n) {
225 cl_int err;
226
227 if (compressible_ops_update_program == NULL)
228 opencl_kernel_jit(compressible_ops_update_kernel,
229 (cl_program *) &compressible_ops_update_program);
230
231 cl_kernel kernel =
232 clCreateKernel((cl_program) compressible_ops_update_program,
233 "ns_flux_finalize_kernel", &err);
234 CL_CHECK(err);
235
236 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &visc_m_x));
237 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &visc_m_y));
238 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &visc_m_z));
239 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &visc_E));
240 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &f_x));
241 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &f_y));
242 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &f_z));
243 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &opgrad_x));
244 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &opgrad_y));
245 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &opgrad_z));
246 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem), (void *) &u));
247 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_mem), (void *) &v));
248 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_mem), (void *) &w));
249 CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_mem), (void *) &B));
250 CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_mem), (void *) &dissipation));
251 CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &n));
252
253 const int nb = (n + 256 - 1) / 256;
254 const size_t global_item_size = 256 * nb;
255 const size_t local_item_size = 256;
256
257 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
258 NULL, &global_item_size, &local_item_size,
259 0, NULL, NULL));
260 CL_CHECK(clReleaseKernel(kernel));
261}
262
263void opencl_ns_flux_temperature(void *div_flux, void *h1, void *p,
264 void *rho, void *kappa, real gamma, int n) {
265 cl_int err;
266
267 if (compressible_ops_update_program == NULL)
268 opencl_kernel_jit(compressible_ops_update_kernel,
269 (cl_program *) &compressible_ops_update_program);
270
271 cl_kernel kernel =
272 clCreateKernel((cl_program) compressible_ops_update_program,
273 "ns_flux_temperature_kernel", &err);
274 CL_CHECK(err);
275
276 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &div_flux));
277 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &h1));
278 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &p));
279 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &rho));
280 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &kappa));
281 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(real), &gamma));
282 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &n));
283
284 const int nb = (n + 256 - 1) / 256;
285 const size_t global_item_size = 256 * nb;
286 const size_t local_item_size = 256;
287
288 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
289 NULL, &global_item_size, &local_item_size,
290 0, NULL, NULL));
291 CL_CHECK(clReleaseKernel(kernel));
292}
opencl_update_uvw
void opencl_update_uvw(void *u, void *v, void *w, void *m_x, void *m_y, void *m_z, void *rho, int n)
Definition compressible_ops_update.c:48
opencl_ns_flux_prepare
void opencl_ns_flux_prepare(void *div_flux, void *dissipation, void *h1, void *dudx, void *dudy, void *dudz, void *dvdx, void *dvdy, void *dvdz, void *dwdx, void *dwdy, void *dwdz, void *mu, int n)
Definition compressible_ops_update.c:178
opencl_update_temperature
void opencl_update_temperature(void *T, void *p, void *rho, real gamma, int n)
Definition compressible_ops_update.c:149
opencl_update_e
void opencl_update_e(void *E, void *p, void *ruvw, real gamma, int n)
Definition compressible_ops_update.c:120
opencl_ns_flux_temperature
void opencl_ns_flux_temperature(void *div_flux, void *h1, void *p, void *rho, void *kappa, real gamma, int n)
Definition compressible_ops_update.c:263
opencl_ns_flux_finalize
void opencl_ns_flux_finalize(void *visc_m_x, void *visc_m_y, void *visc_m_z, void *visc_E, void *f_x, void *f_y, void *f_z, void *opgrad_x, void *opgrad_y, void *opgrad_z, void *u, void *v, void *w, void *B, void *dissipation, int n)
Definition compressible_ops_update.c:219
opencl_update_mxyz_p_ruvw
void opencl_update_mxyz_p_ruvw(void *m_x, void *m_y, void *m_z, void *p, void *ruvw, void *u, void *v, void *w, void *E, void *rho, real gamma, int n)
Definition compressible_ops_update.c:82
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
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
h1
__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__ h1
Definition ax_helm_full_kernel.h:52
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
compressible_ops_update_program
void * compressible_ops_update_program