Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
constrain_mixed_bc_kernel.h
Go to the documentation of this file.
1/*
2 Copyright (c) 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#ifndef __BC_CONSTRAIN_MIXED_BC_KERNEL__
36#define __BC_CONSTRAIN_MIXED_BC_KERNEL__
37
38template< typename T >
39__global__ void constrain_mixed_bc_zero_kernel(
40 const int * __restrict__ mixed_msk,
41 T * __restrict__ x,
42 T * __restrict__ y,
43 T * __restrict__ z,
44 const int constraint_n,
45 const int constraint_t1,
46 const int constraint_t2,
47 const T * __restrict__ n,
48 const T * __restrict__ t1,
49 const T * __restrict__ t2,
50 const int m) {
51
52 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
53 const int str = blockDim.x * gridDim.x;
54
55 for (int i = idx; i < m; i += str) {
56 const int k = mixed_msk[i];
57 const int off = 3 * i;
58
59 T u1 = x[k];
60 T u2 = y[k];
61 T u3 = z[k];
62
63 T uloc_n = u1 * n[off] + u2 * n[off + 1] + u3 * n[off + 2];
64 T uloc_t1 = u1 * t1[off] + u2 * t1[off + 1] + u3 * t1[off + 2];
65 T uloc_t2 = u1 * t2[off] + u2 * t2[off + 1] + u3 * t2[off + 2];
66
67 if (constraint_n != 0)
68 uloc_n = 0;
69 if (constraint_t1 != 0)
70 uloc_t1 = 0;
71 if (constraint_t2 != 0)
72 uloc_t2 = 0;
73
74 x[k] = uloc_n * n[off] + uloc_t1 * t1[off] + uloc_t2 * t2[off];
75 y[k] = uloc_n * n[off + 1] + uloc_t1 * t1[off + 1] +
76 uloc_t2 * t2[off + 1];
77 z[k] = uloc_n * n[off + 2] + uloc_t1 * t1[off + 2] +
78 uloc_t2 * t2[off + 2];
79 }
80}
81
82template< typename T >
83__global__ void constrain_mixed_bc_set_kernel(
84 const int * __restrict__ mixed_msk,
85 T * __restrict__ x,
86 T * __restrict__ y,
87 T * __restrict__ z,
88 const int constraint_n,
89 const int constraint_t1,
90 const int constraint_t2,
91 const T * __restrict__ n,
92 const T * __restrict__ t1,
93 const T * __restrict__ t2,
94 const T * __restrict__ values_n,
95 const T * __restrict__ values_t1,
96 const T * __restrict__ values_t2,
97 const int m) {
98
99 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
100 const int str = blockDim.x * gridDim.x;
101
102 for (int i = idx; i < m; i += str) {
103 const int k = mixed_msk[i];
104 const int off = 3 * i;
105
106 T u1 = x[k];
107 T u2 = y[k];
108 T u3 = z[k];
109
110 T uloc_n = u1 * n[off] + u2 * n[off + 1] + u3 * n[off + 2];
111 T uloc_t1 = u1 * t1[off] + u2 * t1[off + 1] + u3 * t1[off + 2];
112 T uloc_t2 = u1 * t2[off] + u2 * t2[off + 1] + u3 * t2[off + 2];
113
114 T glb_1 = values_n[i];
115 T glb_2 = values_t1[i];
116 T glb_3 = values_t2[i];
117 T glb_n = glb_1 * n[off] + glb_2 * n[off + 1] + glb_3 * n[off + 2];
118 T glb_t1 = glb_1 * t1[off] + glb_2 * t1[off + 1] + glb_3 * t1[off + 2];
119 T glb_t2 = glb_1 * t2[off] + glb_2 * t2[off + 1] + glb_3 * t2[off + 2];
120
121 if (constraint_n != 0)
122 uloc_n = glb_n;
123 if (constraint_t1 != 0)
124 uloc_t1 = glb_t1;
125 if (constraint_t2 != 0)
126 uloc_t2 = glb_t2;
127
128 x[k] = uloc_n * n[off] + uloc_t1 * t1[off] + uloc_t2 * t2[off];
129 y[k] = uloc_n * n[off + 1] + uloc_t1 * t1[off + 1] +
130 uloc_t2 * t2[off + 1];
131 z[k] = uloc_n * n[off + 2] + uloc_t1 * t1[off + 2] +
132 uloc_t2 * t2[off + 2];
133 }
134}
135
136template< typename T >
137__global__ void constrain_mixed_bc_set_const_kernel(
138 const int * __restrict__ mixed_msk,
139 T * __restrict__ x,
140 T * __restrict__ y,
141 T * __restrict__ z,
142 const int constraint_n,
143 const int constraint_t1,
144 const int constraint_t2,
145 const T * __restrict__ n,
146 const T * __restrict__ t1,
147 const T * __restrict__ t2,
148 const T value_n,
149 const T value_t1,
150 const T value_t2,
151 const int m) {
152
153 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
154 const int str = blockDim.x * gridDim.x;
155
156 for (int i = idx; i < m; i += str) {
157 const int k = mixed_msk[i];
158 const int off = 3 * i;
159
160 T u1 = x[k];
161 T u2 = y[k];
162 T u3 = z[k];
163
164 T uloc_n = u1 * n[off] + u2 * n[off + 1] + u3 * n[off + 2];
165 T uloc_t1 = u1 * t1[off] + u2 * t1[off + 1] + u3 * t1[off + 2];
166 T uloc_t2 = u1 * t2[off] + u2 * t2[off + 1] + u3 * t2[off + 2];
167
168 T glb_n = value_n * n[off] + value_t1 * n[off + 1] +
169 value_t2 * n[off + 2];
170 T glb_t1 = value_n * t1[off] + value_t1 * t1[off + 1] +
171 value_t2 * t1[off + 2];
172 T glb_t2 = value_n * t2[off] + value_t1 * t2[off + 1] +
173 value_t2 * t2[off + 2];
174
175 if (constraint_n != 0)
176 uloc_n = glb_n;
177 if (constraint_t1 != 0)
178 uloc_t1 = glb_t1;
179 if (constraint_t2 != 0)
180 uloc_t2 = glb_t2;
181
182 x[k] = uloc_n * n[off] + uloc_t1 * t1[off] + uloc_t2 * t2[off];
183 y[k] = uloc_n * n[off + 1] + uloc_t1 * t1[off + 1] +
184 uloc_t2 * t2[off + 1];
185 z[k] = uloc_n * n[off + 2] + uloc_t1 * t1[off + 2] +
186 uloc_t2 * t2[off + 2];
187 }
188}
189
190#endif
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
i
const int i
Definition ax_helm_full_kernel.h:95
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:110
constrain_mixed_bc_zero_kernel
__global__ void constrain_mixed_bc_zero_kernel(const int *__restrict__ mixed_msk, T *__restrict__ x, T *__restrict__ y, T *__restrict__ z, const int constraint_n, const int constraint_t1, const int constraint_t2, const T *__restrict__ n, const T *__restrict__ t1, const T *__restrict__ t2, const int m)
Definition constrain_mixed_bc_kernel.h:39
constrain_mixed_bc_set_kernel
__global__ void constrain_mixed_bc_set_kernel(const int *__restrict__ mixed_msk, T *__restrict__ x, T *__restrict__ y, T *__restrict__ z, const int constraint_n, const int constraint_t1, const int constraint_t2, const T *__restrict__ n, const T *__restrict__ t1, const T *__restrict__ t2, const T *__restrict__ values_n, const T *__restrict__ values_t1, const T *__restrict__ values_t2, const int m)
Definition constrain_mixed_bc_kernel.h:83
constrain_mixed_bc_set_const_kernel
__global__ void constrain_mixed_bc_set_const_kernel(const int *__restrict__ mixed_msk, T *__restrict__ x, T *__restrict__ y, T *__restrict__ z, const int constraint_n, const int constraint_t1, const int constraint_t2, const T *__restrict__ n, const T *__restrict__ t1, const T *__restrict__ t2, const T value_n, const T value_t1, const T value_t2, const int m)
Definition constrain_mixed_bc_kernel.h:137