Neko 1.99.7
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 Measured a 1.6x win on GH200 and
56 a loss on MI250X/MI300A, hence the per backend default
57
58 Note NEKO_TUNE_ROUNDS and NEKO_TUNE_ITERS below are not specific to
59 this sweep: they control the sampling of every candidate the tuner
60 times, chunk sizes included. */
61#ifndef NEKO_EB_SWEEP_DEFAULT
62#define NEKO_EB_SWEEP_DEFAULT 1
63#endif
64
65#define NEKO_TUNE_INIT 1.0e30f
66#define NEKO_TUNE_WARMUP 20
67
68static int neko_eb_sweep()
69{
70 const char *v = getenv("NEKO_EB_TUNE");
71
72 if (v != NULL) {
73 return (atoi(v) != 0);
74 }
76}
77
78/* Forced candidate, used when NEKO_AUTOTUNE pins the kstep variant */
79static int neko_eb_env()
80{
81 const char *v = getenv("NEKO_EB");
82 int c = (v != NULL) ? atoi(v) : 0;
83
85 c = 0;
86 }
87 return c;
88}
89
90/* Forced chunk candidate, used when NEKO_AUTOTUNE pins the 1d variant */
91static int neko_chunks_env()
92{
93 const char *v = getenv("NEKO_CHUNKS");
94 int c = (v != NULL) ? atoi(v) : 0;
95
97 c = 0;
98 }
99 return c;
100}
101
103{
104 const char *v = getenv("NEKO_TUNE_ROUNDS");
105 int n = (v != NULL) ? atoi(v) : 3;
106
107 return (n < 1) ? 1 : n;
108}
109
110static int neko_tune_iters()
111{
112 const char *v = getenv("NEKO_TUNE_ITERS");
113 int n = (v != NULL) ? atoi(v) : 100;
114
115 return (n < 1) ? 1 : n;
116}
117
118/* One timed round of LAUNCH at candidate C, min reduced into T[C] */
119#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
120 do { \
121 float t_; \
122 cudaEventRecord(start, stream); \
123 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
124 cudaEventRecord(stop, stream); \
125 cudaEventSynchronize(stop); \
126 cudaEventElapsedTime(&t_, start, stop); \
127 if (t_ < (T)[C]) { (T)[C] = t_; } \
128 } while (0)
129
130#define NEKO_TUNE_BEST(T, BEST, N) \
131 do { \
132 for (int c = 1; c < (N); c++) { \
133 if ((T)[c] < (T)[BEST]) { BEST = c; } \
134 } \
135 } while (0)
136
137/* Report every measured candidate of both sweeps, not just the winner */
138#define NEKO_TUNE_LOG(LX, T1, T2) \
139 do { \
140 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
141 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
142 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
143 NEKO_CHUNKS_SEL(LX, c), (T1)[c] * 10.0); \
144 log_message(neko_log_buf); \
145 } \
146 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
147 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
148 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
149 NEKO_EB_SEL(LX, c), (T2)[c] * 10.0); \
150 log_message(neko_log_buf); \
151 } \
152 } while (0)
153
154#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()