elem_block.h

Go to the documentation of this file.

#ifndef __MATH_ELEM_BLOCK_H__
#define __MATH_ELEM_BLOCK_H__
/*
 Copyright (c) 2026, The Neko Authors
 All rights reserved.
 
 Redistribution and use in source and binary forms, with or without
 modification, are permitted provided that the following conditions
 are met:
 
   * Redistributions of source code must retain the above copyright
     notice, this list of conditions and the following disclaimer.
 
   * Redistributions in binary form must reproduce the above
     copyright notice, this list of conditions and the following
     disclaimer in the documentation and/or other materials provided
     with the distribution.
 
   * Neither the name of the authors nor the names of its
     contributors may be used to endorse or promote products derived
     from this software without specific prior written permission.
 
 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
 POSSIBILITY OF SUCH DAMAGE.
*/
 
#include "wave.h"
 
#ifndef NEKO_EB_WAVE
#define NEKO_EB_WAVE NEKO_WAVE_SIZE_UNIFORM
#endif
 
/* LDS per CU, and the per workgroup maximum, on CDNA */
#ifndef NEKO_EB_MAX_LDS
#define NEKO_EB_MAX_LDS 65536
#endif
 
#define NEKO_EB_CANDIDATES 3
 
template< int LX, int C >
struct elem_block {
  static const int waves = 2 << C;
  static const int fit = (NEKO_EB_WAVE * waves)/(LX * LX);
  static const int value = (C == 0) ? 1 : (fit > 0 ? fit : 1);
};
 
/*
 * Launch bounds.
 *
 * HIP's second argument is MIN_WARPS_PER_EXECUTION_UNIT -- minimum wavefronts
 * resident per SIMD -- and not CUDA's min blocks per SM. The first argument
 * must track the real block size: launching more threads than declared is a
 * launch failure, not a slowdown.
 *
 * The second argument defaults to 1, i.e. no constraint. It carried 3
 * historically, which holds the compiler to a 512/3 VGPR budget on every
 * kstep launch. The kstep kernels keep ru[LX] and rw[LX] -- 16 doubles, 32
 * VGPRs -- live across a fully unrolled k loop that also issues seven metric
 * loads per step, and the vector variants keep six such arrays, ~96 VGPRs,
 * before any addressing. That is the shape that spills to scratch under a
 * cap.
 *
 * Measured on MI250X, lx = 8, 13824 elements, on two builds differing only
 * in this setting: kstep eb=1 runs 1317 us with the cap at 3 and 456 us
 * without it, eb=4 4824 against 2254. Over the same pair the 1d kernel,
 * which declares no launch bounds at all, moves 449.82 -> 455.53, and the
 * MFMA kernels, which carry a bare __launch_bounds__ with no second
 * argument, move less than 4%. Only the kernels this macro touches moved.
 *
 * Note what the fix does not buy: kstep eb=1 uncapped only draws level with
 * the 1d kernel (456 against 450), because both are at the DRAM roof. It
 * matters for the vector operator, which has no 1d variant at all, and for
 * not paying for spilled candidates on every tune.
 *
 * The CUDA header records the same failure mode from its own second argument
 * -- "squeezed registers to 80 at lx >= 14 and made cdtp spill 736 bytes per
 * thread" -- and dropped it; see cuda/elem_block.h.
 *
 * Set -DNEKO_EB_MIN_WAVES_EU=3 to restore the old behaviour, and check
 * ScratchSize with -Rpass-analysis=kernel-resource-usage when changing it.
 */
#ifndef NEKO_EB_MIN_WAVES_EU
#define NEKO_EB_MIN_WAVES_EU 1
#endif
 
/* A minimum of one wave per SIMD is no constraint at all, so emit the
   single argument form rather than a bound that only looks like one. */
#if NEKO_EB_MIN_WAVES_EU > 1
#define NEKO_EB_BOUNDS(NT) __launch_bounds__((NT), NEKO_EB_MIN_WAVES_EU)
#else
#define NEKO_EB_BOUNDS(NT) __launch_bounds__((NT))
#endif
 
#ifndef NEKO_EB_WAVE_BARRIER
#define NEKO_EB_WAVE_BARRIER 0
#endif
 
/*
 * Is a kstep slice wavefront contained, so that a wave barrier covers it?
 *
 * Deliberately NOT keyed on NEKO_EB_WAVE. That macro is the command line
 * width, and it sizes kernel template arguments and launch geometry, where
 * the host and device passes have to agree on one number; getting it wrong
 * there costs occupancy, not correctness. Here it would cost correctness, so
 * this asks the target instead of being told.
 *
 * __builtin_amdgcn_wavefrontsize() is the supported spelling.
 * __AMDGCN_WAVEFRONT_SIZE__ is deprecated -- clang warns that compile-time
 * constant access to the wavefront size is going away -- so there is no
 * static_assert to be had, and wanting one was the wrong instinct: under a
 * fixed --offload-arch the builtin folds to a constant and the branch below
 * disappears, and under a multi target fatbin it correctly resolves per
 * target, which a build time check could not do at all.
 */
__device__ __host__ constexpr bool neko_eb_slice_in_wave(unsigned wave, int lx)
{
  return ((unsigned) (lx * lx) <= wave) &&
         ((wave % (unsigned) (lx * lx)) == 0);
}
 
template< int LX >
__device__ __forceinline__ void neko_eb_kstep_barrier()
{
#if NEKO_EB_WAVE_BARRIER && defined(__HIP_DEVICE_COMPILE__)
  /* Uniform across the whole block, so the __syncthreads() arm is never
     reached by only part of it */
  if (neko_eb_slice_in_wave(__builtin_amdgcn_wavefrontsize(), LX)) {
    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront");
    __builtin_amdgcn_wave_barrier();
    __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront");
  } else {
    __syncthreads();
  }
#else
  __syncthreads();
#endif
}
 
/*
 * Launch geometry helpers. NELV is the element count; the grid is sized so
 * the tail block is partially filled and clamps, rather than the kernel
 * returning early -- these kernels have __syncthreads() in the k loop, so an
 * early return would leave the barrier unmatched.
 */
#define NEKO_EB(LX, C) (elem_block<LX, C>::value)
#define NEKO_EB_NTHRDS(LX, C) dim3((LX), (LX), NEKO_EB(LX, C))
#define NEKO_EB_NBLCKS(NELV, LX, C)                                           \
  dim3(((NELV) + NEKO_EB(LX, C) - 1)/NEKO_EB(LX, C), 1, 1)
#define NEKO_EB_SEL(LX, SEL)                                                  \
  ((SEL) == 0 ? NEKO_EB(LX, 0) : (SEL) == 1 ? NEKO_EB(LX, 1) : NEKO_EB(LX, 2))
 
#define NEKO_CHUNKS_CANDIDATES 4
 
/*
 * Candidate C selects 1024 >> C, i.e. 1024, 512, 256, 128. Candidate 0 is the
 * historical value, so it is always available as an A/B baseline.
 *
 * A heuristic is deliberately avoided here. The best chunk depends on how the
 * shared memory footprint (which grows as LX^3) caps the resident block count
 * on the target device, and that differs between architectures -- 164 kB per
 * SM on A100, 228 kB on H100, 64 kB of LDS per CU on CDNA. Picking by thread
 * utilisation alone gets it wrong from LX = 6 upward, because once the block
 * count is capped by shared memory a smaller block simply wastes thread slots.
 * So all four are instantiated and the tuner measures them.
 */
template< int LX, int C >
struct chunk_block {
  static const int raw = 1024 >> C;
  static const int value = (raw >= LX * LX) ? raw : 1024;
};
 
#define NEKO_CHUNKS(LX, C) (chunk_block<LX, C>::value)
#define NEKO_CHUNKS_NTHRDS(LX, C) dim3(NEKO_CHUNKS(LX, C), 1, 1)
#define NEKO_CHUNKS_SEL(LX, SEL)                                              \
  ((SEL) == 0 ? NEKO_CHUNKS(LX, 0) : (SEL) == 1 ? NEKO_CHUNKS(LX, 1) :        \
   (SEL) == 2 ? NEKO_CHUNKS(LX, 2) : NEKO_CHUNKS(LX, 3))
 
#endif // __MATH_ELEM_BLOCK_H__