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 Measured a 1.6x win on GH200 and
63 a loss on MI250X/MI300A, hence the per backend default
64
65 Note NEKO_TUNE_ROUNDS and NEKO_TUNE_ITERS below are not specific to
66 this sweep: they control the sampling of every candidate the tuner
67 times, chunk sizes included. */
68#ifndef NEKO_EB_SWEEP_DEFAULT
69#define NEKO_EB_SWEEP_DEFAULT 1
70#endif
71
72#define NEKO_TUNE_INIT 1.0e30f
73#define NEKO_TUNE_WARMUP 20
74
75static int neko_eb_sweep()
76{
77 const char *v = getenv("NEKO_EB_TUNE");
78
79 if (v != NULL) {
80 return (atoi(v) != 0);
81 }
83}
84
85/*
86 * Elements per block candidate pinned by NEKO_EB, or -1 to leave it to the
87 * sweep.
88 *
89 * NEKO_AUTOTUNE selects the formulation and nothing more -- the geometry
90 * inside it is still measured unless it is pinned here -- so "unset" has to
91 * be distinguishable from candidate 0, which is a candidate like any other.
92 * Hence the -1 rather than a plain default of 0. Out of range values clamp
93 * to 0 rather than releasing the pin, as they always have.
94 */
95static int neko_eb_pin()
96{
97 const char *v = getenv("NEKO_EB");
98 int c;
99
100 if (v == NULL) {
101 return -1;
102 }
103
104 c = atoi(v);
106 c = 0;
107 }
108 return c;
109}
110
111/* Chunk candidate pinned by NEKO_CHUNKS, or -1 to sweep, see neko_eb_pin() */
112static int neko_chunks_pin()
113{
114 const char *v = getenv("NEKO_CHUNKS");
115 int c;
116
117 if (v == NULL) {
118 return -1;
119 }
120
121 c = atoi(v);
123 c = 0;
124 }
125 return c;
126}
127
129{
130 const char *v = getenv("NEKO_TUNE_ROUNDS");
131 int n = (v != NULL) ? atoi(v) : 3;
132
133 return (n < 1) ? 1 : n;
134}
135
136static int neko_tune_iters()
137{
138 const char *v = getenv("NEKO_TUNE_ITERS");
139 int n = (v != NULL) ? atoi(v) : 100;
140
141 return (n < 1) ? 1 : n;
142}
143
144/*
145 * Loop over the candidates of one formulation, binding C to each in turn.
146 *
147 * ON gates the formulation itself: hardware support, and -- when NEKO_AUTOTUNE
148 * is set -- whether it is the formulation that was asked for. PIN is the
149 * candidate that formulation's own variable forces, or -1 to measure all N of
150 * them. So NEKO_AUTOTUNE narrows the search to one kernel family without
151 * deciding the geometry within it, which is still swept and reported.
152 *
153 * An out of play formulation yields an empty range, leaving its candidates at
154 * NEKO_TUNE_INIT so it loses the comparison at the end rather than having to
155 * be excluded from it.
156 */
157#define NEKO_TUNE_FOR(C, ON, PIN, N) \
158 for (int C = (((PIN) >= 0) ? (PIN) : 0), \
159 C##_last_ = ((ON) ? (((PIN) >= 0) ? (PIN) + 1 : (N)) : 0); \
160 C < C##_last_; C++)
161
162/* One timed round of LAUNCH at candidate C, min reduced into T[C] */
163#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
164 do { \
165 float t_; \
166 cudaEventRecord(start, stream); \
167 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
168 cudaEventRecord(stop, stream); \
169 cudaEventSynchronize(stop); \
170 cudaEventElapsedTime(&t_, start, stop); \
171 if (t_ < (T)[C]) { (T)[C] = t_; } \
172 } while (0)
173
174/*
175 * Elapsed time (ms, over ITERS launches) as microseconds per call.
176 *
177 * A division rather than a constant because ITERS is NEKO_TUNE_ITERS, which
178 * is settable: the factor of 10 this used to carry is only correct at the
179 * default of 100 and silently rescales every reported time otherwise. The
180 * ranking never moves, every candidate sharing the divisor, but the number
181 * the log prints does.
182 */
183#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
184
185#define NEKO_TUNE_BEST(T, BEST, N) \
186 do { \
187 for (int c = 1; c < (N); c++) { \
188 if ((T)[c] < (T)[BEST]) { BEST = c; } \
189 } \
190 } while (0)
191
192/* Report every measured candidate of both sweeps, not just the winner */
193#define NEKO_TUNE_LOG(LX, T1, T2) \
194 do { \
195 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
196 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
197 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
198 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
199 log_message(neko_log_buf); \
200 } \
201 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
202 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
203 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
204 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
205 log_message(neko_log_buf); \
206 } \
207 } while (0)
208
209#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()