Neko 1.99.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
euler_res.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2025, The Neko Authors
3 All rights reserved.
4 Redistribution and use in source and binary forms, with or without
5 modification, are permitted provided that the following conditions
6 are met:
7 * Redistributions of source code must retain the above copyright
8 notice, this list of conditions and the following disclaimer.
9 * Redistributions in binary form must reproduce the above
10 copyright notice, this list of conditions and the following
11 disclaimer in the documentation and/or other materials provided
12 with the distribution.
13 * Neither the name of the authors nor the names of its
14 contributors may be used to endorse or promote products derived
15 from this software without specific prior written permission.
16 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
17 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
18 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
19 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
20 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
21 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
22 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
24 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
25 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
26 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
27 POSSIBILITY OF SUCH DAMAGE.
28*/
29
30#ifdef __APPLE__
31#include <OpenCL/cl.h>
32#else
33#include <CL/cl.h>
34#endif
35
36#include <stdio.h>
37#include <device/device_config.h>
38#include <device/opencl/jit.h>
39#include <device/opencl/prgm_lib.h>
40#include <device/opencl/check.h>
41
42#include "euler_res_kernel.cl.h"
43
44void euler_res_part_visc_opencl(void *rhs_u, void *Binv, void *lap_sol,
45 void *h, real *c_avisc, int *n) {
46 cl_int err;
47
48 if (euler_res_program == NULL)
49 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
50
51 cl_kernel kernel = clCreateKernel(euler_res_program,
52 "euler_res_part_visc_kernel", &err);
53 CL_CHECK(err);
54
55 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &rhs_u));
56 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &Binv));
57 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &lap_sol));
58 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &h));
59 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(real), c_avisc));
60 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), n));
61
62 const int nb = ((*n) + 256 - 1) / 256;
63 const size_t global_item_size = 256 * nb;
64 const size_t local_item_size = 256;
65
66 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
67 NULL, &global_item_size, &local_item_size,
68 0, NULL, NULL));
69}
70
71void euler_res_part_mx_flux_opencl(void *f_x, void *f_y, void *f_z,
72 void *m_x, void *m_y, void *m_z,
73 void *rho_field, void *p, int *n) {
74 cl_int err;
75
76 if (euler_res_program == NULL)
77 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
78
79 cl_kernel kernel = clCreateKernel(euler_res_program,
80 "euler_res_part_mx_flux_kernel", &err);
81 CL_CHECK(err);
82
83 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &f_x));
84 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &f_y));
85 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &f_z));
86 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_x));
87 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &m_y));
88 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &m_z));
89 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &rho_field));
90 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &p));
91 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), n));
92
93 const int nb = ((*n) + 256 - 1) / 256;
94 const size_t global_item_size = 256 * nb;
95 const size_t local_item_size = 256;
96
97 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
98 NULL, &global_item_size, &local_item_size,
99 0, NULL, NULL));
100}
101
102void euler_res_part_my_flux_opencl(void *f_x, void *f_y, void *f_z,
103 void *m_x, void *m_y, void *m_z,
104 void *rho_field, void *p, int *n) {
105 cl_int err;
106
107 if (euler_res_program == NULL)
108 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
109
110 cl_kernel kernel = clCreateKernel(euler_res_program,
111 "euler_res_part_my_flux_kernel", &err);
112 CL_CHECK(err);
113
114 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &f_x));
115 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &f_y));
116 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &f_z));
117 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_x));
118 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &m_y));
119 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &m_z));
120 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &rho_field));
121 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &p));
122 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), n));
123
124 const int nb = ((*n) + 256 - 1) / 256;
125 const size_t global_item_size = 256 * nb;
126 const size_t local_item_size = 256;
127
128 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
129 NULL, &global_item_size, &local_item_size,
130 0, NULL, NULL));
131}
132
133void euler_res_part_mz_flux_opencl(void *f_x, void *f_y, void *f_z,
134 void *m_x, void *m_y, void *m_z,
135 void *rho_field, void *p, int *n) {
136 cl_int err;
137
138 if (euler_res_program == NULL)
139 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
140
141 cl_kernel kernel = clCreateKernel(euler_res_program,
142 "euler_res_part_mz_flux_kernel", &err);
143 CL_CHECK(err);
144
145 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &f_x));
146 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &f_y));
147 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &f_z));
148 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_x));
149 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &m_y));
150 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &m_z));
151 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &rho_field));
152 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &p));
153 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), n));
154
155 const int nb = ((*n) + 256 - 1) / 256;
156 const size_t global_item_size = 256 * nb;
157 const size_t local_item_size = 256;
158
159 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
160 NULL, &global_item_size, &local_item_size,
161 0, NULL, NULL));
162}
163
164void euler_res_part_E_flux_opencl(void *f_x, void *f_y, void *f_z,
165 void *m_x, void *m_y, void *m_z,
166 void *rho_field, void *p, void * E, int *n) {
167 cl_int err;
168
169 if (euler_res_program == NULL)
170 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
171
172 cl_kernel kernel = clCreateKernel(euler_res_program,
173 "euler_res_part_E_flux_kernel", &err);
174 CL_CHECK(err);
175
176 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &f_x));
177 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &f_y));
178 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &f_z));
179 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_x));
180 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &m_y));
181 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &m_z));
182 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &rho_field));
183 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &E));
184 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &p));
185 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), n));
186
187 const int nb = ((*n) + 256 - 1) / 256;
188 const size_t global_item_size = 256 * nb;
189 const size_t local_item_size = 256;
190
191 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
192 NULL, &global_item_size, &local_item_size,
193 0, NULL, NULL));
194}
195
196void euler_res_part_coef_mult_opencl(void *rhs_rho, void *rhs_m_x,
197 void *rhs_m_y, void *rhs_m_z,
198 void *rhs_E, void *mult, int *n) {
199 cl_int err;
200
201 if (euler_res_program == NULL)
202 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
203
204 cl_kernel kernel = clCreateKernel(euler_res_program,
205 "euler_res_part_coef_mult_kernel", &err);
206 CL_CHECK(err);
207
208 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &rhs_rho));
209 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &rhs_m_x));
210 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &rhs_m_y));
211 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &rhs_m_z));
212 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &rhs_E));
213 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &mult));
214 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), n));
215
216 const int nb = ((*n) + 256 - 1) / 256;
217 const size_t global_item_size = 256 * nb;
218 const size_t local_item_size = 256;
219
220 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
221 NULL, &global_item_size, &local_item_size,
222 0, NULL, NULL));
223}
224
225void euler_res_part_rk_sum_opencl(void *rho, void *m_x, void *m_y, void *m_z,
226 void *E, void *k_rho_i, void *k_m_x_i,
227 void *k_m_y_i, void *k_m_z_i, void *k_E_i,
228 real *dt, real *c, int *n) {
229 cl_int err;
230
231 if (euler_res_program == NULL)
232 opencl_kernel_jit(euler_res_kernel, (cl_program *) &euler_res_program);
233
234 cl_kernel kernel = clCreateKernel(euler_res_program,
235 "euler_res_part_rk_sum_kernel", &err);
236 CL_CHECK(err);
237
238 CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &rho));
239 CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &m_x));
240 CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &m_y));
241 CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &m_z));
242 CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &E));
243 CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &k_rho_i));
244 CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &k_m_x_i));
245 CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &k_m_y_i));
246 CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &k_m_z_i));
247 CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_mem), (void *) &k_E_i));
248 CL_CHECK(clSetKernelArg(kernel, 10, sizeof(real), dt));
249 CL_CHECK(clSetKernelArg(kernel, 11, sizeof(real), c));
250 CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), n));
251
252 const int nb = ((*n) + 256 - 1) / 256;
253 const size_t global_item_size = 256 * nb;
254 const size_t local_item_size = 256;
255
256 CL_CHECK(clEnqueueNDRangeKernel((cl_command_queue) glb_cmd_queue, kernel, 1,
257 NULL, &global_item_size, &local_item_size,
258 0, NULL, NULL));
259}
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
euler_res_part_coef_mult_opencl
void euler_res_part_coef_mult_opencl(void *rhs_rho, void *rhs_m_x, void *rhs_m_y, void *rhs_m_z, void *rhs_E, void *mult, int *n)
Definition euler_res.c:196
euler_res_part_E_flux_opencl
void euler_res_part_E_flux_opencl(void *f_x, void *f_y, void *f_z, void *m_x, void *m_y, void *m_z, void *rho_field, void *p, void *E, int *n)
Definition euler_res.c:164
euler_res_part_my_flux_opencl
void euler_res_part_my_flux_opencl(void *f_x, void *f_y, void *f_z, void *m_x, void *m_y, void *m_z, void *rho_field, void *p, int *n)
Definition euler_res.c:102
euler_res_part_mx_flux_opencl
void euler_res_part_mx_flux_opencl(void *f_x, void *f_y, void *f_z, void *m_x, void *m_y, void *m_z, void *rho_field, void *p, int *n)
Definition euler_res.c:71
euler_res_part_mz_flux_opencl
void euler_res_part_mz_flux_opencl(void *f_x, void *f_y, void *f_z, void *m_x, void *m_y, void *m_z, void *rho_field, void *p, int *n)
Definition euler_res.c:133
euler_res_part_rk_sum_opencl
void euler_res_part_rk_sum_opencl(void *rho, void *m_x, void *m_y, void *m_z, void *E, void *k_rho_i, void *k_m_x_i, void *k_m_y_i, void *k_m_z_i, void *k_E_i, real *dt, real *c, int *n)
Definition euler_res.c:225
euler_res_part_visc_opencl
void euler_res_part_visc_opencl(void *rhs_u, void *Binv, void *lap_sol, void *h, real *c_avisc, int *n)
Definition euler_res.c:44
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
euler_res_program
void * euler_res_program