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
52#include <stdlib.h>
53
54/* Sweep the elements per block candidates?
55 On by default, as on CUDA. It was off here because the blocked variants
56 measured a loss on MI250X and MI300A, where they spill -- but baking that
57 in also stopped the tuner from ever re-testing it, and the kstep family has
58 since measured badly enough on gfx90a (788 us/call at lx = 8 against 268
59 for the 1d variant) that its geometry is worth measuring per case rather
60 than assuming. Where the blocked variants still spill the tuner simply
61 rejects them; the cost is tuning time, not run time.
62
63 Note NEKO_TUNE_ROUNDS and NEKO_TUNE_ITERS below are not specific to
64 this sweep: they control the sampling of every candidate the tuner
65 times, chunk sizes included. */
66#ifndef NEKO_EB_SWEEP_DEFAULT
67#define NEKO_EB_SWEEP_DEFAULT 1
68#endif
69
70#define NEKO_TUNE_INIT 1.0e30f
71#define NEKO_TUNE_WARMUP 20
72
73static int neko_eb_sweep()
74{
75 const char *v = getenv("NEKO_EB_TUNE");
76
77 if (v != NULL) {
78 return (atoi(v) != 0);
79 }
81}
82
83/* Forced candidate, used when NEKO_AUTOTUNE pins the kstep variant */
84static int neko_eb_env()
85{
86 const char *v = getenv("NEKO_EB");
87 int c = (v != NULL) ? atoi(v) : 0;
88
90 c = 0;
91 }
92 return c;
93}
94
95/* Forced chunk candidate, used when NEKO_AUTOTUNE pins the 1d variant */
96static int neko_chunks_env()
97{
98 const char *v = getenv("NEKO_CHUNKS");
99 int c = (v != NULL) ? atoi(v) : 0;
100
102 c = 0;
103 }
104 return c;
105}
106
108{
109 const char *v = getenv("NEKO_TUNE_ROUNDS");
110 int n = (v != NULL) ? atoi(v) : 3;
111
112 return (n < 1) ? 1 : n;
113}
114
115static int neko_tune_iters()
116{
117 const char *v = getenv("NEKO_TUNE_ITERS");
118 int n = (v != NULL) ? atoi(v) : 100;
119
120 return (n < 1) ? 1 : n;
121}
122
123/* One timed round of LAUNCH at candidate C, min reduced into T[C] */
124#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
125 do { \
126 float t_; \
127 HIP_CHECK(hipEventRecord(start, stream)); \
128 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
129 HIP_CHECK(hipEventRecord(stop, stream)); \
130 HIP_CHECK(hipEventSynchronize(stop)); \
131 HIP_CHECK(hipEventElapsedTime(&t_, start, stop)); \
132 if (t_ < (T)[C]) { (T)[C] = t_; } \
133 } while (0)
134
135/*
136 * Elapsed time (ms, over ITERS launches) as microseconds per call.
137 *
138 * A division rather than a constant because ITERS is NEKO_TUNE_ITERS, which
139 * is settable: the factor of 10 this used to carry is only correct at the
140 * default of 100 and silently rescales every reported time otherwise. The
141 * ranking never moves, every candidate sharing the divisor, but the number
142 * the log prints does.
143 */
144#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
145
146#define NEKO_TUNE_BEST(T, BEST, N) \
147 do { \
148 for (int c = 1; c < (N); c++) { \
149 if ((T)[c] < (T)[BEST]) { BEST = c; } \
150 } \
151 } while (0)
152
153/* Report every measured candidate of both sweeps, not just the winner */
154#define NEKO_TUNE_LOG(LX, T1, T2) \
155 do { \
156 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
157 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
158 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
159 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
160 log_message(neko_log_buf); \
161 } \
162 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
163 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
164 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
165 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
166 log_message(neko_log_buf); \
167 } \
168 } while (0)
169
170#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_eb_env()
static int neko_tune_rounds()
static int neko_tune_iters()
static int neko_chunks_env()
static int neko_eb_sweep()