Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
wave.h
Go to the documentation of this file.
1#ifndef __MATH_HIP_WAVE_H__
2#define __MATH_HIP_WAVE_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#ifdef NEKO_WAVE_SIZE
60
61/* Set on the command line, hence identical in both passes */
62#define NEKO_WAVE_SIZE_UNIFORM NEKO_WAVE_SIZE
63
64#else
65
66/*
67 * An explicit wavefront macro is tried first: it is the only form that tracks
68 * an explicit -mwavefrontsize flag. Both spellings have to be checked. The
69 * bare __AMDGCN_WAVEFRONT_SIZE is the original and is what every ROCm up to
70 * the 6.x series defines; __AMDGCN_WAVEFRONT_SIZE__ was added later, and ROCm
71 * 6.3 deprecated both without providing a compile time replacement. They are
72 * expanded here rather than at the point of use, so that the whole build reads
73 * them exactly once.
74 *
75 * That one read is still a diagnostic, because the deprecation is a
76 * `#pragma clang deprecated` and fires on expansion, including inside `#if`.
77 * It is silenced immediately around the two tests and nowhere else: the
78 * replacement the deprecation points at, __builtin_amdgcn_wavefrontsize(), is
79 * folded in the backend and so cannot be seen by the preprocessor, which is
80 * the only place a value is needed here. `defined()` does not expand its
81 * operand, so the guards themselves are silent and only the value tests have
82 * to be wrapped. When a compile time replacement does appear, this whole block
83 * -- pragmas included -- is what it replaces.
84 *
85 * The gfx family macros are the fallback for a toolchain that defines neither.
86 * gfx6 through gfx9 (GCN, Vega, CDNA) are 64 lanes; gfx10 and later (RDNA)
87 * default to 32.
88 *
89 * A target that matches nothing at all is a build error rather than a silent
90 * 64: on RDNA that would not fail, because __shfl_down past the end of a
91 * wavefront returns the lane's own value, so the reduction would quietly
92 * double every partial sum instead.
93 */
94# if defined(__clang__)
95# pragma clang diagnostic push
96# pragma clang diagnostic ignored "-Wdeprecated-pragma"
97# endif
98
99# if defined(__AMDGCN_WAVEFRONT_SIZE__)
100# if __AMDGCN_WAVEFRONT_SIZE__ == 32
101# define NEKO_WAVE_SIZE 32
102# else
103# define NEKO_WAVE_SIZE 64
104# endif
105# elif defined(__AMDGCN_WAVEFRONT_SIZE)
106# if __AMDGCN_WAVEFRONT_SIZE == 32
107# define NEKO_WAVE_SIZE 32
108# else
109# define NEKO_WAVE_SIZE 64
110# endif
111# elif defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__)
112# define NEKO_WAVE_SIZE 32
113# elif defined(__GFX6__) || defined(__GFX7__) || defined(__GFX8__) || \
114 defined(__GFX9__)
115# define NEKO_WAVE_SIZE 64
116# elif defined(__HIP_DEVICE_COMPILE__)
117# error "Unknown AMD wavefront width, rebuild with -DNEKO_WAVE_SIZE=32 or 64"
118# else
119# define NEKO_WAVE_SIZE 64 /* host pass, see the note above */
120# endif
121
122# if defined(__clang__)
123# pragma clang diagnostic pop
124# endif
125
126/*
127 * The pass independent width. Auto detection cannot be used here, so this
128 * stays at the CDNA 64 unless the width is pinned on the command line. It only
129 * feeds tuning heuristics (elem_block.h), never correctness, and the blocked
130 * candidates it sizes are measured off by default on this backend anyway.
131 */
132#define NEKO_WAVE_SIZE_UNIFORM 64
133
134#endif /* NEKO_WAVE_SIZE */
135
136#if (NEKO_WAVE_SIZE != 32) && (NEKO_WAVE_SIZE != 64)
137#error "NEKO_WAVE_SIZE must be 32 or 64"
138#endif
139
140/*
141 * The two stage block reductions stage one partial per wavefront in a flat
142 * __shared__ T [64]. Every reducing kernel is launched at the 1024 thread
143 * maximum, so that is 16 entries at wave64 and 32 at wave32: the existing 64
144 * covers both widths and is left alone.
145 */
146
147#endif // __MATH_HIP_WAVE_H__