woutermans
8813ea723c
L's unfolded enc0 and dec0 are still ZLUDA's compile of NVIDIA's PTX and cost 2.95 ms of the preset's 5.86 ms per frame (RX 9070 XT, 1280x720 -> 4K). The embedding, the patch expand, the Swin block (32 channels, 2 heads, at pixel resolution) and the patch merge are now native token-lane code; the texture work around them stays ZLUDA's compile. enc0l_tail.hip cuts in at the barrier after the feature staging and dec0l_block.hip before the output code that reads the block's f16 result out of the last MMAs; make_ptx.py splices either in when D4R_TEX_SRC names it and build.sh selects both for the gfx12 native-FP8 builds. swin_tex.h carries the indexing the two tails share; swin_tok.h gains the pixel-order token layout (SWIN_TOK_NATURAL), ZLUDA's attention key order (SWIN_PV_ZLUDA_ORDER) and swin_tok_core, which the encoded tests reuse. swin_common.h takes the block weight layout (moved out of swin_block.h) and the exact and approximate reciprocal / mix forms the tails need without the HIP headers (D4R_NO_HIP_RUNTIME). The f32 sums stay in ZLUDA's order (K slots of the FP8 WMMAs, key order of the f16 attention product) and rsqrt/rcp use the hardware approximations ZLUDA lowers to, so the harness's final frame is byte-identical to the translated kernels: all four enc0 flag variants, 720p and 1080p render, 1440p and 4K output, static and moving scenes, 100 frames, and the accuracy set. With M's own sum order about one token in 20,000 differed. M's layers compile to the same instructions and its frame is byte-identical. enc0 1.69 -> 0.63 ms and dec0 1.26 -> 0.54 ms at 720p->4K, 2.73 -> 1.16 and 2.28 -> 1.01 at 1080p->4K; the preset 5.83 -> 4.06 and 8.63 -> 5.69 ms. -DD4R_SWIN_K32_MIX rounds the f16 products earlier for 0.56 and 0.49 ms and is not byte-identical, so it is not built by default. The repo's tests pass (36, 3 skipped).
51 lines
2.0 KiB
C++
51 lines
2.0 KiB
C++
// The token-lane Swin block (../m/swin_tok.h) for the native parts of ZLUDA-compiled texture kernels: preset L's
|
|
// enc0 (enc0l_tail.hip) and dec0 (dec0l_block.hip).
|
|
#pragma once
|
|
|
|
// The Swin sources are written against the HIP headers; a texture tail is built without them. These are the
|
|
// runtime names they use.
|
|
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H
|
|
#define D4R_NO_HIP_RUNTIME
|
|
#define __device__ __attribute__((device))
|
|
#define __host__ __attribute__((host))
|
|
#define __forceinline__ inline __attribute__((always_inline))
|
|
#define __shared__ __attribute__((shared))
|
|
typedef __SIZE_TYPE__ size_t;
|
|
struct d4r_tid_t
|
|
{
|
|
struct X { __attribute__((device, always_inline)) operator unsigned() const { return __builtin_amdgcn_workitem_id_x(); } } x;
|
|
struct Z { __attribute__((device, always_inline)) operator unsigned() const { return __builtin_amdgcn_workitem_id_z(); } } z;
|
|
};
|
|
static const d4r_tid_t threadIdx;
|
|
__attribute__((device, always_inline)) static inline void __syncthreads()
|
|
{
|
|
__builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
|
|
__builtin_amdgcn_s_barrier();
|
|
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
|
|
}
|
|
#endif
|
|
|
|
// 1.0 for the k32 rounding of the D4R_SWIN_K32_MIX variant (see mix2 in ../m/swin_common.h). The prep kernel stores
|
|
// it, so that it is not a constant to the compiler.
|
|
__attribute__((device, used)) float d4r_swin_one = 1.0f;
|
|
#define MIX_ONE d4r_swin_one
|
|
|
|
namespace d4r_swin
|
|
{
|
|
#define SWIN_EXACT
|
|
// The original code is ZLUDA's compile of the PTX; these make the f32 arithmetic the same instruction for
|
|
// instruction where it is not exact (see ../m/swin_tok.h, ../m/swin_common.h): the K order and the C operand of
|
|
// the FP8 WMMAs, the key order of the f16 P V product, and the hardware reciprocals.
|
|
#define SWIN_TOK_NATURAL
|
|
#ifndef D4R_SWIN_K32_MIX
|
|
#define K32_NO_MIX
|
|
#endif
|
|
#define SWIN_PV_ZLUDA_ORDER
|
|
#define SWIN_APPROX_RECIP
|
|
#ifndef D4R_FP8_WMMA
|
|
#define D4R_FP8_WMMA
|
|
#endif
|
|
#include "../m/swin_common.h"
|
|
#include "../m/swin_tok.h"
|
|
} // namespace d4r_swin
|