Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
elem_block_tune.h
Go to the documentation of this file.
1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
3/*
4 Copyright (c) 2026, The Neko Authors
5 All rights reserved.
6
7 Redistribution and use in source and binary forms, with or without
8 modification, are permitted provided that the following conditions
9 are met:
10
11 * Redistributions of source code must retain the above copyright
12 notice, this list of conditions and the following disclaimer.
13
14 * Redistributions in binary form must reproduce the above
15 copyright notice, this list of conditions and the following
16 disclaimer in the documentation and/or other materials provided
17 with the distribution.
18
19 * Neither the name of the authors nor the names of its
20 contributors may be used to endorse or promote products derived
21 from this software without specific prior written permission.
22
23 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
24 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
25 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
26 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
27 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
28 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
29 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
30 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
31 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
32 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
33 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
34 POSSIBILITY OF SUCH DAMAGE.
35*/
36
59#include <stdlib.h>
60
61/* Sweep the elements per block candidates?
62 On by default, as on CUDA. It was off here because the blocked variants
63 measured a loss on MI250X and MI300A, where they spill -- but baking that
64 in also stopped the tuner from ever re-testing it, and the kstep family has
65 since measured badly enough on gfx90a (788 us/call at lx = 8 against 268
66 for the 1d variant) that its geometry is worth measuring per case rather
67 than assuming. Where the blocked variants still spill the tuner simply
68 rejects them; the cost is tuning time, not run time.
69
70 Note NEKO_TUNE_ROUNDS and NEKO_TUNE_ITERS below are not specific to
71 this sweep: they control the sampling of every candidate the tuner
72 times, chunk sizes included. */
73#ifndef NEKO_EB_SWEEP_DEFAULT
74#define NEKO_EB_SWEEP_DEFAULT 1
75#endif
76
77#define NEKO_TUNE_INIT 1.0e30f
78#define NEKO_TUNE_WARMUP 20
79
80static int neko_eb_sweep()
81{
82 const char *v = getenv("NEKO_EB_TUNE");
83
84 if (v != NULL) {
85 return (atoi(v) != 0);
86 }
88}
89
90/*
91 * Elements per block candidate pinned by NEKO_EB, or -1 to leave it to the
92 * sweep.
93 *
94 * NEKO_AUTOTUNE selects the formulation and nothing more -- the geometry
95 * inside it is still measured unless it is pinned here -- so "unset" has to
96 * be distinguishable from candidate 0, which is a candidate like any other.
97 * Hence the -1 rather than a plain default of 0. Out of range values clamp
98 * to 0 rather than releasing the pin, as they always have.
99 */
100static int neko_eb_pin()
101{
102 const char *v = getenv("NEKO_EB");
103 int c;
104
105 if (v == NULL) {
106 return -1;
107 }
108
109 c = atoi(v);
111 c = 0;
112 }
113 return c;
114}
115
116/* Chunk candidate pinned by NEKO_CHUNKS, or -1 to sweep, see neko_eb_pin() */
117static int neko_chunks_pin()
118{
119 const char *v = getenv("NEKO_CHUNKS");
120 int c;
121
122 if (v == NULL) {
123 return -1;
124 }
125
126 c = atoi(v);
128 c = 0;
129 }
130 return c;
131}
132
134{
135 const char *v = getenv("NEKO_TUNE_ROUNDS");
136 int n = (v != NULL) ? atoi(v) : 3;
137
138 return (n < 1) ? 1 : n;
139}
140
141static int neko_tune_iters()
142{
143 const char *v = getenv("NEKO_TUNE_ITERS");
144 int n = (v != NULL) ? atoi(v) : 100;
145
146 return (n < 1) ? 1 : n;
147}
148
149/*
150 * Loop over the candidates of one formulation, binding C to each in turn.
151 *
152 * ON gates the formulation itself: hardware support, and -- when NEKO_AUTOTUNE
153 * is set -- whether it is the formulation that was asked for. PIN is the
154 * candidate that formulation's own variable forces, or -1 to measure all N of
155 * them. So NEKO_AUTOTUNE narrows the search to one kernel family without
156 * deciding the geometry within it, which is still swept and reported.
157 *
158 * An out of play formulation yields an empty range, leaving its candidates at
159 * NEKO_TUNE_INIT so it loses the comparison at the end rather than having to
160 * be excluded from it.
161 */
162#define NEKO_TUNE_FOR(C, ON, PIN, N) \
163 for (int C = (((PIN) >= 0) ? (PIN) : 0), \
164 C##_last_ = ((ON) ? (((PIN) >= 0) ? (PIN) + 1 : (N)) : 0); \
165 C < C##_last_; C++)
166
167/* One timed round of LAUNCH at candidate C, min reduced into T[C] */
168#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
169 do { \
170 float t_; \
171 HIP_CHECK(hipEventRecord(start, stream)); \
172 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
173 HIP_CHECK(hipEventRecord(stop, stream)); \
174 HIP_CHECK(hipEventSynchronize(stop)); \
175 HIP_CHECK(hipEventElapsedTime(&t_, start, stop)); \
176 if (t_ < (T)[C]) { (T)[C] = t_; } \
177 } while (0)
178
179/*
180 * Elapsed time (ms, over ITERS launches) as microseconds per call.
181 *
182 * A division rather than a constant because ITERS is NEKO_TUNE_ITERS, which
183 * is settable: the factor of 10 this used to carry is only correct at the
184 * default of 100 and silently rescales every reported time otherwise. The
185 * ranking never moves, every candidate sharing the divisor, but the number
186 * the log prints does.
187 */
188#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
189
190#define NEKO_TUNE_BEST(T, BEST, N) \
191 do { \
192 for (int c = 1; c < (N); c++) { \
193 if ((T)[c] < (T)[BEST]) { BEST = c; } \
194 } \
195 } while (0)
196
197/* Report every measured candidate of one sweep, not just the winner. Split
198 per formulation because the vector operator has no 1d variant and so wants
199 the kstep half on its own */
200#define NEKO_TUNE_LOG_1D(LX, T1) \
201 do { \
202 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
203 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
204 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
205 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
206 log_message(neko_log_buf); \
207 } \
208 } while (0)
209
210#define NEKO_TUNE_LOG_KSTEP(LX, T2) \
211 do { \
212 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
213 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
214 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
215 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
216 log_message(neko_log_buf); \
217 } \
218 } while (0)
219
220/* Both sweeps, for the operators that have a 1d variant */
221#define NEKO_TUNE_LOG(LX, T1, T2) \
222 do { \
223 NEKO_TUNE_LOG_1D(LX, T1); \
224 NEKO_TUNE_LOG_KSTEP(LX, T2); \
225 } while (0)
226
227#endif // __MATH_ELEM_BLOCK_TUNE_H__
__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)
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
#define NEKO_EB_SWEEP_DEFAULT
static int neko_chunks_pin()
static int neko_tune_rounds()
static int neko_eb_pin()
static int neko_tune_iters()
static int neko_eb_sweep()