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