From c0aa47ec110ad7e02fd3c9f28f02be1bd381a41b Mon Sep 17 00:00:00 2001 From: wangyuascend-spec Date: Mon, 7 Sep 2026 11:41:35 +0800 Subject: [PATCH 1/2] feat(normalization): move 4PE solutions into solution tree --- benchmark/one-level-arch/compile_all.sh | 5 +- .../kernels/multi_thread/README.md | 2 +- .../normalization/rms_norm_binary/SKILL.md | 2 +- .../one-level-arch/kernels/solution/.gitkeep | 0 .../group_norm_grad/group_norm_grad_pto.hpp | 432 ++++++++++++++++++ .../group_norm_grad_1d_pto.hpp | 339 ++++++++++++++ .../normalization/rms_norm/rms_norm_pto.hpp | 150 ++++++ .../normalization/rms_norm_binary/SKILL.md | 191 ++++++++ .../rms_norm_binary/rms_norm_binary_pto.hpp | 260 +++++++++++ .../normalization/group_norm_grad/Makefile | 29 -- .../normalization/group_norm_grad/compile.all | 14 - .../group_norm_grad/src/group_norm_grad.cpp | 91 ---- .../src/group_norm_grad_static.cpp | 84 ---- .../normalization/group_norm_grad_1d/Makefile | 30 -- .../group_norm_grad_1d/compile.all | 14 - .../src/group_norm_grad_1d.cpp | 81 ---- .../src/group_norm_grad_1d_static.cpp | 74 --- .../kernel/normalization/rms_norm/Makefile | 30 -- .../kernel/normalization/rms_norm/compile.all | 21 - .../normalization/rms_norm/src/rms_norm.cpp | 65 --- .../rms_norm/src/rms_norm_static.cpp | 72 --- .../normalization/rms_norm_binary/Makefile | 21 - .../normalization/rms_norm_binary/compile.all | 12 - .../rms_norm_binary/src/rms_norm_binary.cpp | 71 --- .../src/rms_norm_binary_static.cpp | 64 --- .../one-level-arch/test/solution/.gitkeep | 0 .../normalization/group_norm_grad/Makefile | 17 + .../normalization/group_norm_grad/compile.all | 5 + .../src/gen_group_norm_grad_data.py | 16 +- .../group_norm_grad/src/group_norm_grad.cpp | 128 ++++++ .../src/group_norm_grad_data_compare.py | 12 +- .../normalization/group_norm_grad_1d/Makefile | 16 + .../group_norm_grad_1d/compile.all | 5 + .../src/gen_group_norm_grad_1d_data.py | 10 +- .../src/group_norm_grad_1d.cpp | 123 +++++ .../src/group_norm_grad_1d_data_compare.py | 10 +- .../solution/normalization/rms_norm/Makefile | 14 + .../normalization/rms_norm/compile.all | 5 + .../rms_norm/src/data/build_data_obj.sh | 0 .../rms_norm/src/gen_rms_norm_data.py | 12 +- .../normalization/rms_norm/src/rms_norm.cpp | 95 ++++ .../rms_norm/src/rms_norm_data_compare.py | 4 +- .../rms_norm/src/run_precision_check.py | 6 +- .../normalization/rms_norm_binary/Makefile | 14 + .../normalization/rms_norm_binary/compile.all | 5 + .../src/gen_rms_norm_binary_data.py | 14 +- .../rms_norm_binary/src/rms_norm_binary.cpp | 109 +++++ .../src/rms_norm_binary_data_compare.py | 4 +- 48 files changed, 1962 insertions(+), 816 deletions(-) delete mode 100644 benchmark/one-level-arch/kernels/solution/.gitkeep create mode 100644 benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad/group_norm_grad_pto.hpp create mode 100644 benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad_1d/group_norm_grad_1d_pto.hpp create mode 100644 benchmark/one-level-arch/kernels/solution/normalization/rms_norm/rms_norm_pto.hpp create mode 100644 benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/SKILL.md create mode 100644 benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/Makefile delete mode 100755 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_static.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/Makefile delete mode 100755 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_static.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm/Makefile delete mode 100755 benchmark/one-level-arch/test/kernel/normalization/rms_norm/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_static.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/Makefile delete mode 100755 benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_static.cpp delete mode 100644 benchmark/one-level-arch/test/solution/.gitkeep create mode 100644 benchmark/one-level-arch/test/solution/normalization/group_norm_grad/Makefile create mode 100755 benchmark/one-level-arch/test/solution/normalization/group_norm_grad/compile.all rename benchmark/one-level-arch/test/{kernel => solution}/normalization/group_norm_grad/src/gen_group_norm_grad_data.py (94%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad.cpp rename benchmark/one-level-arch/test/{kernel => solution}/normalization/group_norm_grad/src/group_norm_grad_data_compare.py (91%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/Makefile create mode 100755 benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/compile.all rename benchmark/one-level-arch/test/{kernel => solution}/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py (96%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp rename benchmark/one-level-arch/test/{kernel => solution}/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py (91%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/rms_norm/Makefile create mode 100755 benchmark/one-level-arch/test/solution/normalization/rms_norm/compile.all rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm/src/data/build_data_obj.sh (100%) rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm/src/gen_rms_norm_data.py (93%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm.cpp rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm/src/rms_norm_data_compare.py (96%) rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm/src/run_precision_check.py (92%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/Makefile create mode 100755 benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/compile.all rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py (91%) create mode 100644 benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary.cpp rename benchmark/one-level-arch/test/{kernel => solution}/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py (96%) diff --git a/benchmark/one-level-arch/compile_all.sh b/benchmark/one-level-arch/compile_all.sh index 1f00beba..b0988b95 100755 --- a/benchmark/one-level-arch/compile_all.sh +++ b/benchmark/one-level-arch/compile_all.sh @@ -59,8 +59,11 @@ compile_operator "$REPO_ROOT/test/kernel/reduction/reducesum_row" "reducesum_row compile_operator "$REPO_ROOT/test/kernel/control" "control" compile_operator "$REPO_ROOT/test/kernel/fa" "fa" compile_operator "$REPO_ROOT/test/kernel/sort" "sort" -compile_operator "$REPO_ROOT/test/kernel/deepseek" "deepseek" compile_operator "$REPO_ROOT/test/kernel/flashMLA" "flashMLA" +compile_operator "$REPO_ROOT/test/solution/normalization/rms_norm" "solution/normalization/rms_norm" +compile_operator "$REPO_ROOT/test/solution/normalization/rms_norm_binary" "solution/normalization/rms_norm_binary" +compile_operator "$REPO_ROOT/test/solution/normalization/group_norm_grad" "solution/normalization/group_norm_grad" +compile_operator "$REPO_ROOT/test/solution/normalization/group_norm_grad_1d" "solution/normalization/group_norm_grad_1d" compile_operator "$REPO_ROOT/test/kernel/multi_thread/vec" "multi_thread/vec" compile_operator "$REPO_ROOT/test/kernel/multi_thread/broadcast" "multi_thread/broadcast" compile_operator "$REPO_ROOT/test/kernel/multi_thread/concat" "multi_thread/concat" diff --git a/benchmark/one-level-arch/kernels/multi_thread/README.md b/benchmark/one-level-arch/kernels/multi_thread/README.md index 782786ac..0078f2fc 100644 --- a/benchmark/one-level-arch/kernels/multi_thread/README.md +++ b/benchmark/one-level-arch/kernels/multi_thread/README.md @@ -43,7 +43,7 @@ Kernel and test paths mirror the single-PE tree. For example: | `kernels/single_thread/gather/gather.hpp` | `kernels/multi_thread/gather/gather.hpp` | | `test/kernel/gather/` | `test/kernel/multi_thread/gather/` | | `test/kernel/element_wise/gelu/` | `test/kernel/multi_thread/element_wise/gelu/` | -| `test/kernel/normalization/rms_norm_binary/` | `test/kernel/multi_thread/normalization/rms_norm_binary/` | +| `test/solution/normalization/rms_norm_binary/` | `test/kernel/multi_thread/normalization/rms_norm_binary/` | Each operator directory has its own `Makefile`, `compile.all`, and `src/` instead of sharing a mixed test source. One model failure therefore does not diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md b/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md index 8220c0e1..df25db21 100644 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md +++ b/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md @@ -41,7 +41,7 @@ Current default test shape: **`[1, 8192]`**, `tile_r=1024` → **`Rb=8`**, fp16. |------|------| | Kernel | `$ROOT/SuperNPUBench/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp` | | Reference (single-tile) | `.../kernels/single_thread/normalization/rms_norm/rms_norm.hpp` | -| Testcase | `.../test/kernel/normalization/rms_norm_binary/` | +| Testcase | `.../test/solution/normalization/rms_norm_binary/` | | Host entry | `.../rms_norm_binary/src/rms_norm_binary.cpp` | | Gen golden | `.../rms_norm_binary/src/gen_rms_norm_binary_data.py` | | Compare | `.../rms_norm_binary/src/rms_norm_binary_data_compare.py` | diff --git a/benchmark/one-level-arch/kernels/solution/.gitkeep b/benchmark/one-level-arch/kernels/solution/.gitkeep deleted file mode 100644 index e69de29b..00000000 diff --git a/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad/group_norm_grad_pto.hpp b/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad/group_norm_grad_pto.hpp new file mode 100644 index 00000000..25cb28c7 --- /dev/null +++ b/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad/group_norm_grad_pto.hpp @@ -0,0 +1,432 @@ +// ============================================================================= +// group_norm_grad_pto.hpp — GroupNorm backward, HxW > 1 (one-level PTO) +// ============================================================================= +// +// Matches PyTorch GroupNormBackwardKernelImplInternal +// (aten/src/ATen/native/cuda/group_norm_kernel.cu): +// 1) Spatial reduce ds/db = Σ_hw (dY*X), Σ_hw dY per (n,c) +// 2) Fused c2/c3 from ds/db over channels → workspace +// 3) dX = (rstd*gamma)*dY + c2*X + c3 +// 4) dgamma / dbeta from ds/db over N +// +// Layout: X/dY/dX [N,C,HxW]; mean/rstd [N,G] fp32; gamma/dgamma/dbeta [C] +// workspace float[2*N*C+2*N*G] +// All pointers required (dy,x,mean,rstd,gamma,dx,dgamma,dbeta,workspace). +// tiling[5] = {N, C, G, HxW, tile_hw} +// tile_hw <= 0 → min(HxW, tCap); spatial R-split when HxW > tile_hw. +// dX still requires HxW <= tCap (one spatial tile; RF-limited). +// +// tCap: logical tile >= 512B (TileOP IsValidActiveSize / TSize=1..7). +// fp16 → Cols>=256; fp32 data tiles → Cols>=128. Reduction/broadcast +// tile_v follows the hardware contract and uses physical Columns=1. +// Cols=1024 like rms_norm overflows Tile RF here. +// +// Torch CUDA launch 总览 (NVIDIA, warp=32): +// Step1 ComputeInternalGradientsCUDAKernel +// grid=N*C, block=(HxW<512)?32:512 +// Step2 ComputeBackwardFusedParamsCUDAKernel +// grid=dim3(N,G), block=(D<512)?32:512 +// Step3 dX gpu_kernel (+ optional c1) +// block=128, vt=4(fp16)/2(fp32), grid=ceil(numel/(128*vt)) +// Step4 GammaBetaBackwardCUDAKernel1/2 +// N<=128: grid=ceil(C/256), block=256 +// N>128: grid=ceil(C/32), block=dim3(32,16) +// ============================================================================= +#ifndef SUPERNPU_GROUP_NORM_GRAD_PTO_HPP +#define SUPERNPU_GROUP_NORM_GRAD_PTO_HPP + +#include + +#include + +namespace gn_grad { + + +inline int64_t workspace_elems(int64_t N, int64_t C, int64_t G) { + return 2 * N * C + 2 * N * G; +} + +constexpr int kMaxPeCount = 64; +static volatile uint32_t kPeBarrier[kMaxPeCount] = {}; + +__attribute__((noinline)) inline uint32_t read_pe_id() { + return get_thread_idx(); +} + +template +__attribute__((noinline)) void pe_barrier(uint32_t phase) { + static_assert(peNum > 0 && peNum <= kMaxPeCount); + if constexpr (peNum > 1) { + const uint32_t pe = read_pe_id(); + kPeBarrier[pe] = phase; + for (int participant = 0; participant < peNum; ++participant) { + while (kPeBarrier[participant] < phase) { + } + } + } +} + +// --------------------------------------------------------------------------- +// Step 1: spatial reduce for one (n, c) → ds[nc], db[nc] (HxW R-split) +// +// Torch: ComputeInternalGradientsCUDAKernel +// grid = N * C // 每个 (n,c) 一个 block；本函数 = 其中一个 block +// block = (HxW < 512) ? 32 : 512 +// 线程: threadIdx.x 沿 hw 做 grid-stride + warp/block reduce +// → ds[n,c]=Σ dY*X, db[n,c]=Σ dY +// --------------------------------------------------------------------------- +template +inline void spatial_reduce_nc(dtype *dy, dtype *x, float *ds, float *db, + int64_t N, int64_t C, int64_t HxW, + int64_t tile_hw, int64_t n, int64_t c) { + const int64_t nc = n * C + c; + const int64_t base = nc * HxW; + + gm_f gds(ds + nc, static_cast(N * C), 1); + gm_f gdb(db + nc, static_cast(N * C), 1); + + tile_v ds_acc(1); + tile_v db_acc(1); + tile_v cur(1); + TEXPANDS(ds_acc, 0.0f); + TEXPANDS(db_acc, 0.0f); + + // Torch: for (hw = threadIdx.x; hw < HxW; hw += blockDim.x) + BlockReduce + // PTO: tile 覆盖一段 HxW，TROWSUM 代替 block 内线程归约 + for (int64_t hw0 = 0; hw0 < HxW; hw0 += tile_hw) { + const size_t vh = static_cast( + (hw0 + tile_hw <= HxW) ? tile_hw : (HxW - hw0)); + const int64_t offset = base + hw0; + + gm_h gdy(dy + offset, static_cast(N * C), static_cast(HxW)); + gm_h gx(x + offset, static_cast(N * C), static_cast(HxW)); + + tile_h h0(1, vh); + tile_f x_f(1, vh); + tile_f dy_f(1, vh); + tile_f prod(1, vh); + + TLOAD(h0, gx); + TCVT(x_f, h0); + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TMUL(prod, dy_f, x_f); + TROWSUM(cur, prod); + TADD(ds_acc, ds_acc, cur); + TROWSUM(cur, dy_f); + TADD(db_acc, db_acc, cur); + } + + TSTORE(gds, ds_acc); + TSTORE(gdb, db_acc); +} + +// --------------------------------------------------------------------------- +// Step 2: fused c2/c3 for one (n, g) → c2[ng], c3[ng] +// +// Torch: ComputeBackwardFusedParamsCUDAKernel +// grid = dim3(N, G) // blockIdx.x=n, blockIdx.y=g；本函数 = 其中一个 +// block = (D < 512) ? 32 : 512 +// 线程: threadIdx.x 沿 group 内通道 i∈[0,D) stride，再 block reduce +// → c2,c3 每 (n,g) 各一个标量 +// --------------------------------------------------------------------------- +template +inline void fused_params_group(dtype *gamma, float *mean, float *rstd, + float *ds, float *db, float *c2_buf, + float *c3_buf, int64_t N, int64_t C, int64_t G, + int64_t D, int64_t n, int64_t g, float s) { + const int64_t ng = n * G + g; + const int64_t c0 = g * D; + const size_t active_d = static_cast(D); + + gm_f gmean(mean + ng, static_cast(N * G), 1); + gm_f grstd(rstd + ng, static_cast(N * G), 1); + gm_f gds(ds + n * C + c0, 1, static_cast(C)); + gm_f gdb(db + n * C + c0, 1, static_cast(C)); + gm_f gc2(c2_buf + ng, static_cast(N * G), 1); + gm_f gc3(c3_buf + ng, static_cast(N * G), 1); + + tile_f ds_f(1, active_d); + tile_f db_f(1, active_d); + tile_f gamma_f(1, active_d); + tile_f t0(1, active_d); + tile_h h0(1, active_d); + tile_v mean_t(1); + tile_v rstd_t(1); + tile_v sum1(1); + tile_v sum2(1); + tile_v c2(1); + tile_v c3(1); + + TLOAD(ds_f, gds); + TLOAD(db_f, gdb); + TLOAD(mean_t, gmean); + TLOAD(rstd_t, grstd); + + { + gm_h gg(gamma + c0, 1, static_cast(C)); + TLOAD(h0, gg); + TCVT(gamma_f, h0); + } + + // Torch: threads 各算 ds*gamma / db*gamma 再 reduce → sum1/sum2 + TMUL(t0, ds_f, gamma_f); + TROWSUM(sum1, t0); + TMUL(t0, db_f, gamma_f); + TROWSUM(sum2, t0); + + // c2/c3 由 block 内 thread 0（归约后）写出；此处标量 tile 完成同样公式 + TMUL(c2, sum2, mean_t); + TSUB(c2, c2, sum1); + TMUL(c3, rstd_t, rstd_t); + TMUL(c3, c3, rstd_t); + TMUL(c2, c2, c3); + TMULS(c2, c2, s); + + TMUL(c3, c2, mean_t); + TMULS(c3, c3, -1.0f); + TMUL(sum1, sum2, rstd_t); + TMULS(sum1, sum1, s); + TSUB(c3, c3, sum1); + + TSTORE(gc2, c2); + TSTORE(gc3, c3); +} + +// --------------------------------------------------------------------------- +// Step 3: dX for one (n, c) using stored c2/c3 and rstd*gamma +// +// Torch: gpu_kernel 元素级 (可选先算 c1) +// block = 128 +// vt = 4 (fp16/bf16) / 2 (fp32+) +// grid = ceil(numel / (128 * vt)) // numel = N*C*HxW +// 线程: 线性下标覆盖全部元素；c2/c3 按 (n,g) 广播 +// 本函数一次处理一个 (n,c) 的整段 HxW（Tile 覆盖空间维） +// --------------------------------------------------------------------------- +template +inline void dx_nc(dtype *dy, dtype *x, dtype *gamma, float *rstd, float *c2_buf, + float *c3_buf, dtype *dx, int64_t N, int64_t C, int64_t G, + int64_t D, int64_t HxW, int64_t tile_hw, int64_t n, + int64_t c) { + const int64_t g = c / D; + const int64_t ng = n * G + g; + gm_f grstd(rstd + ng, static_cast(N * G), 1); + gm_f gc2(c2_buf + ng, static_cast(N * G), 1); + gm_f gc3(c3_buf + ng, static_cast(N * G), 1); + + tile_v rstd_t(1); + tile_v c1(1); + tile_v c2(1); + tile_v c3(1); + + TLOAD(rstd_t, grstd); + TLOAD(c2, gc2); + TLOAD(c3, gc3); + + // Torch 可选 c1 预计算同为 gpu_kernel block=128；此处 c1 = rstd*gamma[c] + // TCVT must keep matching logical shapes (PTO 0.58). tile_h Cols=tCap + // (fp16→256) while tile_v Cols=128, so convert via tile_f then TROWSUM. + { + gm_h gg(gamma + c, 1, 1); + tile_h hg(1, 1); + tile_f gf(1, 1); + tile_v gv(1); + TLOAD(hg, gg); + TCVT(gf, hg); + TROWSUM(gv, gf); + TMUL(c1, gv, rstd_t); + } + + const int64_t base = (n * C + c) * HxW; + for (int64_t hw0 = 0; hw0 < HxW; hw0 += tile_hw) { + const size_t active_hw = static_cast( + (hw0 + tile_hw <= HxW) ? tile_hw : (HxW - hw0)); + const int64_t offset = base + hw0; + gm_h gdy(dy + offset, static_cast(N * C), static_cast(HxW)); + gm_h gx(x + offset, static_cast(N * C), static_cast(HxW)); + gm_h gdx(dx + offset, static_cast(N * C), static_cast(HxW)); + tile_h h0(1, active_hw); + tile_f x_f(1, active_hw); + tile_f dy_f(1, active_hw); + tile_f dx_f(1, active_hw); + tile_f tmp(1, active_hw); + TLOAD(h0, gx); + TCVT(x_f, h0); + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TROWEXPANDMUL(dx_f, dy_f, c1); + TROWEXPANDMUL(tmp, x_f, c2); + TADD(dx_f, dx_f, tmp); + TROWEXPANDADD(dx_f, dx_f, c3); + TCVT(h0, dx_f); + TSTORE(gdx, h0); + } +} + +// --------------------------------------------------------------------------- +// Step 4a: dbeta — dbeta[c] = Σ_n db[n,c] +// +// Torch: GammaBetaBackwardCUDAKernel1/2（与 dgamma 同一次 launch） +// N<=128: grid=ceil(C/256), block=256；每线程一个 c，循环 n +// N>128: grid=ceil(C/32), block=dim3(32,16) +// 本函数按 group 一次写 D 个通道（对 N 串行累加） +// --------------------------------------------------------------------------- +template +inline void dbeta_group(float *db, dtype *dbeta, int64_t N, int64_t C, + int64_t D, int64_t g) { + const int64_t c0 = g * D; + const size_t active_d = static_cast(D); + + tile_f acc(1, active_d); + tile_f cur(1, active_d); + tile_h h0(1, active_d); + TEXPANDS(acc, 0.0f); + + for (int64_t n = 0; n < N; ++n) { + gm_f gdb(db + n * C + c0, 1, static_cast(C)); + TLOAD(cur, gdb); + TADD(acc, acc, cur); + } + + gm_h gout(dbeta + c0, 1, static_cast(C)); + TCVT(h0, acc); + TSTORE(gout, h0); +} + +// --------------------------------------------------------------------------- +// Step 4b: dgamma — dgamma[c] = Σ_n (ds - db*mean)*rstd +// +// Torch: 与 dbeta 同 Kernel1/2 launch（见上） +// --------------------------------------------------------------------------- +template +inline void dgamma_group(float *ds, float *db, float *mean, float *rstd, + dtype *dgamma, int64_t N, int64_t C, int64_t G, + int64_t D, int64_t g) { + const int64_t c0 = g * D; + const size_t active_d = static_cast(D); + + tile_f acc(1, active_d); + tile_f ds_f(1, active_d); + tile_f db_f(1, active_d); + tile_f t0(1, active_d); + tile_h h0(1, active_d); + tile_v mean_t(1); + tile_v rstd_t(1); + TEXPANDS(acc, 0.0f); + + for (int64_t n = 0; n < N; ++n) { + const int64_t ng = n * G + g; + gm_f gds(ds + n * C + c0, 1, static_cast(C)); + gm_f gdb(db + n * C + c0, 1, static_cast(C)); + gm_f gmean(mean + ng, static_cast(N * G), 1); + gm_f grstd(rstd + ng, static_cast(N * G), 1); + + TLOAD(ds_f, gds); + TLOAD(db_f, gdb); + TLOAD(mean_t, gmean); + TLOAD(rstd_t, grstd); + + TROWEXPANDMUL(t0, db_f, mean_t); + TSUB(t0, ds_f, t0); + TROWEXPANDMUL(t0, t0, rstd_t); + TADD(acc, acc, t0); + } + + gm_h gout(dgamma + c0, 1, static_cast(C)); + TCVT(h0, acc); + TSTORE(gout, h0); +} + +} // namespace gn_grad + +// tiling: [N, C, G, HxW, tile_hw] +// workspace: float[2*N*C + 2*N*G] +// +// 入口循环 ↔ Torch 各 kernel 的 grid 遍历： +// for n,c spatial_reduce ↔ grid = N*C +// for n,g fused_params ↔ grid = dim3(N,G) +// for n,c dx_nc ↔ numel 上 gpu_kernel 线性网格 +// for g dbeta/dgamma ↔ 按通道写回（Kernel1/2） +template +void group_norm_grad(dtype *dy, dtype *x, float *mean, float *rstd, + dtype *gamma, const int64_t *tiling, dtype *dx, + dtype *dgamma, dtype *dbeta, float *workspace) { + static_assert(peNum == 4, "normalization kernels support only 4PE"); + // One 8192-element spatial tile uses 32 KiB for fp32, within the 64 KiB TileOP limit. + // Larger HxW values are handled by the tile_hw loops. + constexpr int64_t tCap = 8192; + constexpr int64_t tV = 1; // row-reduction/broadcast physical Columns=1 + + const int64_t N = tiling[0]; + const int64_t C = tiling[1]; + const int64_t G = tiling[2]; + const int64_t HxW = tiling[3]; + const int64_t tile_hw = + tiling[4] > 0 ? tiling[4] : (HxW < tCap ? HxW : tCap); + const uint32_t tid = gn_grad::read_pe_id(); + if (N <= 0 || C <= 0 || G <= 0 || HxW <= 0 || C % G != 0) { + return; + } + const int64_t D = C / G; + if (tile_hw <= 0 || tile_hw > tCap || D > tCap || + tid >= static_cast(peNum)) { + return; + } + + float *ds = workspace; + float *db = workspace + N * C; + float *c2_buf = workspace + 2 * N * C; + float *c3_buf = c2_buf + N * G; + using gm_h = global_tensor>; + using gm_f = global_tensor>; + using tile_h = + Tile; + using tile_f = + Tile; + using tile_v = + Tile; + + const float s = 1.0f / static_cast(D * HxW); + + for (int64_t nc = tid; nc < N * C; nc += peNum) { + const int64_t n = nc / C; + const int64_t c = nc % C; + gn_grad::spatial_reduce_nc(dy, x, ds, db, N, C, HxW, + tile_hw, n, c); + } + gn_grad::pe_barrier(1); + + for (int64_t ng = tid; ng < N * G; ng += peNum) { + const int64_t n = ng / G; + const int64_t g = ng % G; + gn_grad::fused_params_group(gamma, mean, rstd, ds, db, + c2_buf, c3_buf, N, C, G, D, n, g, + s); + } + gn_grad::pe_barrier(2); + + for (int64_t nc = tid; nc < N * C; nc += peNum) { + const int64_t n = nc / C; + const int64_t c = nc % C; + gn_grad::dx_nc( + dy, x, gamma, rstd, c2_buf, c3_buf, dx, N, C, G, D, HxW, + tile_hw, n, c); + } + + for (int64_t g = tid; g < G; g += peNum) { + gn_grad::dbeta_group(db, dbeta, N, + C, D, g); + gn_grad::dgamma_group( + ds, db, mean, rstd, dgamma, N, C, G, D, g); + } + gn_grad::pe_barrier(3); +} + +#endif // SUPERNPU_GROUP_NORM_GRAD_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad_1d/group_norm_grad_1d_pto.hpp b/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad_1d/group_norm_grad_1d_pto.hpp new file mode 100644 index 00000000..866a47a5 --- /dev/null +++ b/benchmark/one-level-arch/kernels/solution/normalization/group_norm_grad_1d/group_norm_grad_1d_pto.hpp @@ -0,0 +1,339 @@ +// ============================================================================= +// group_norm_grad_1d_pto.hpp — GroupNorm backward, HxW == 1 (one-level PTO) +// ============================================================================= +// +// Matches PyTorch GroupNorm1dBackward +// (aten/src/ATen/native/cuda/group_norm_kernel.cu): +// Stage A1 reduce → c2/c3 (per n,g) +// Stage A2 dX = (rstd*gamma)*dY + c2*X + c3 +// Stage B dgamma / dbeta +// +// Layout: X/dY/dX [N,C]; mean/rstd [N,G] fp32; gamma/dgamma/dbeta [C] +// All pointers required (dy,x,mean,rstd,gamma,dx,dgamma,dbeta). +// +// tiling[4] = {N, C, G, tile_d} +// tile_d <= 0 → min(D, tD); channel R-split on dgamma/dbeta when D > tile_d. +// Stage A requires D <= tD (one tile). +// +// Tile capacity: logical tile >= 512B (TileOP IsValidActiveSize / TSize=1..7). +// fp16 → Cols>=256; fp32 data tiles → Cols>=128. Reduction/broadcast +// tile_v follows the hardware contract and uses physical Columns=1. +// Reduce and dX are separate passes so large tiles do not stay live across both. +// +// Torch CUDA launch 总览 (NVIDIA, warp=32; HxW==1 特化): +// A1 Compute1dBackwardFusedParamsCUDAKernel +// grid=dim3(N,G), block=(D<512)?32:512 +// A2 dX gpu_kernel +// block=128, vt=4(fp16)/2(fp32), grid=ceil(N*C/(128*vt)) +// B GammaBeta1dBackwardCUDAKernel1/2 +// N<=128: grid=ceil(C/256), block=256 +// N>128: grid=ceil(C/32), block=dim3(32,16) +// ============================================================================= +#ifndef SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP +#define SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP + +#include + +#include + +namespace gn_grad_1d { + + +// --------------------------------------------------------------------------- +// Stage A1: channel reduce → c2/c3 for one (n, g) +// scratch[2] = {c2, c3} +// +// Torch: Compute1dBackwardFusedParamsCUDAKernel +// grid = dim3(N, G) // blockIdx.x=n, blockIdx.y=g；本函数 = 其中一个 +// block = (D < 512) ? 32 : 512 +// 线程: threadIdx.x 沿 i∈[0,D) stride，读 dY/X/gamma 累加后 block reduce +// → sum1=Σ dY*X*gamma, sum2=Σ dY*gamma → c2,c3 +// --------------------------------------------------------------------------- +template +inline void fused_params_group(dtype *dy, dtype *x, float *mean, float *rstd, + dtype *gamma, float *scratch, int64_t N, + int64_t C, int64_t G, int64_t D, int64_t n, + int64_t g, float s) { + const int64_t ng = n * G + g; + const int64_t c0 = g * D; + const int64_t offset = n * C + c0; + const size_t active_d = static_cast(D); + + gm_h gdy(dy + offset, static_cast(N), static_cast(C)); + gm_h gx(x + offset, static_cast(N), static_cast(C)); + gm_f gmean(mean + ng, static_cast(N * G), 1); + gm_f grstd(rstd + ng, static_cast(N * G), 1); + gm_f gc2(scratch + 0, 1, 1); + gm_f gc3(scratch + 1, 1, 1); + + tile_h h0(1, active_d); + tile_h h1(1, active_d); + tile_f x_f(1, active_d); + tile_f dy_f(1, active_d); + tile_f t0(1, active_d); + tile_f t1(1, active_d); + tile_v mean_t(1); + tile_v rstd_t(1); + tile_v sum1(1); + tile_v sum2(1); + tile_v c2(1); + tile_v c3(1); + + // Torch: 各 thread 读本组一段通道；PTO 一次 Tile 覆盖整组 D + TLOAD(h0, gx); + TCVT(x_f, h0); + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TLOAD(mean_t, gmean); + TLOAD(rstd_t, grstd); + + { + gm_h gg(gamma + c0, 1, static_cast(C)); + TLOAD(h1, gg); + TCVT(t0, h1); // gamma + } + + // sum2 = Σ dy*gamma ; sum1 = Σ dy*gamma*x ↔ thread 局部累加 + BlockReduce + TMUL(t1, dy_f, t0); + TROWSUM(sum2, t1); + TMUL(t1, t1, x_f); + TROWSUM(sum1, t1); + + // c2 = (sum2*mean - sum1) * rstd^3 * s （归约后标量，通常 thread0 写） + TMUL(c2, sum2, mean_t); + TSUB(c2, c2, sum1); + TMUL(c3, rstd_t, rstd_t); + TMUL(c3, c3, rstd_t); + TMUL(c2, c2, c3); + TMULS(c2, c2, s); + + // c3 = -c2*mean - sum2*rstd*s + TMUL(c3, c2, mean_t); + TMULS(c3, c3, -1.0f); + TMUL(sum1, sum2, rstd_t); + TMULS(sum1, sum1, s); + TSUB(c3, c3, sum1); + + TSTORE(gc2, c2); + TSTORE(gc3, c3); +} + +// --------------------------------------------------------------------------- +// Stage A2: dX for one (n, g) from spilled c2/c3 +// +// Torch: gpu_kernel 元素级 +// block = 128 +// vt = 4 (fp16/bf16) / 2 (fp32+) +// grid = ceil(N*C / (128*vt)) +// 线程: 线性下标覆盖 [N,C]；c2/c3 按 (n,g) 广播到组内通道 +// 本函数一次写完一组 D 个通道（HxW=1） +// --------------------------------------------------------------------------- +template +inline void dx_group(dtype *dy, dtype *x, float *rstd, dtype *gamma, + float *scratch, dtype *dx, int64_t N, int64_t C, + int64_t G, int64_t D, int64_t n, int64_t g) { + const int64_t ng = n * G + g; + const int64_t c0 = g * D; + const int64_t offset = n * C + c0; + const size_t active_d = static_cast(D); + + gm_h gdy(dy + offset, static_cast(N), static_cast(C)); + gm_h gx(x + offset, static_cast(N), static_cast(C)); + gm_h gdx(dx + offset, static_cast(N), static_cast(C)); + gm_f grstd(rstd + ng, static_cast(N * G), 1); + gm_f gc2(scratch + 0, 1, 1); + gm_f gc3(scratch + 1, 1, 1); + + tile_h h0(1, active_d); + tile_h h1(1, active_d); + tile_f x_f(1, active_d); + tile_f dy_f(1, active_d); + tile_f t0(1, active_d); + tile_f t1(1, active_d); + tile_v rstd_t(1); + tile_v c2(1); + tile_v c3(1); + + TLOAD(h0, gx); + TCVT(x_f, h0); + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TLOAD(rstd_t, grstd); + TLOAD(c2, gc2); + TLOAD(c3, gc3); + + { + gm_h gg(gamma + c0, 1, static_cast(C)); + TLOAD(h1, gg); + TCVT(t0, h1); // gamma + } + + // dX = (rstd*gamma)*dY + c2*X + c3 + TROWEXPANDMUL(t1, t0, rstd_t); + TMUL(t1, t1, dy_f); + TROWEXPANDMUL(t0, x_f, c2); + TADD(t1, t1, t0); + TROWEXPANDADD(t1, t1, c3); + + TCVT(h0, t1); + TSTORE(gdx, h0); +} + +// --------------------------------------------------------------------------- +// Stage B: dbeta — dbeta[c] = Σ_n dY[n,c] +// +// Torch: GammaBeta1dBackwardCUDAKernel1/2（与 dgamma 同 launch） +// N<=128: grid=ceil(C/256), block=256；每线程一个 c，循环 n +// N>128: grid=ceil(C/32), block=dim3(32,16) +// --------------------------------------------------------------------------- +template +inline void dbeta_group(dtype *dy, dtype *dbeta, int64_t N, int64_t C, + int64_t D, int64_t tile_d, int64_t g) { + const int64_t c0 = g * D; + + for (int64_t d0 = 0; d0 < D; d0 += tile_d) { + const size_t vd = static_cast( + (d0 + tile_d <= D) ? tile_d : (D - d0)); + + tile_h h0(1, vd); + tile_f dy_f(1, vd); + tile_f acc(1, vd); + TEXPANDS(acc, 0.0f); + + // Torch Kernel1: 单线程 for(n) 累加；此处 Tile 一次累加一组通道 + for (int64_t n = 0; n < N; ++n) { + gm_h gdy(dy + n * C + c0 + d0, static_cast(N), + static_cast(C)); + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TADD(acc, acc, dy_f); + } + + gm_h gdb(dbeta + c0 + d0, 1, static_cast(C)); + TCVT(h0, acc); + TSTORE(gdb, h0); + } +} + +// --------------------------------------------------------------------------- +// Stage B: dgamma — dgamma[c] = Σ_n dY*(X-mean)*rstd +// +// Torch: 与 dbeta 同 Kernel1/2 launch（见上） +// --------------------------------------------------------------------------- +template +inline void dgamma_group(dtype *dy, dtype *x, float *mean, float *rstd, + dtype *dgamma, int64_t N, int64_t C, int64_t G, + int64_t D, int64_t tile_d, int64_t g) { + const int64_t c0 = g * D; + + for (int64_t d0 = 0; d0 < D; d0 += tile_d) { + const size_t vd = static_cast( + (d0 + tile_d <= D) ? tile_d : (D - d0)); + + tile_h h0(1, vd); + tile_f dy_f(1, vd); + tile_f x_f(1, vd); + tile_f t0(1, vd); + tile_f acc(1, vd); + tile_v mean_t(1); + tile_v rstd_t(1); + TEXPANDS(acc, 0.0f); + + for (int64_t n = 0; n < N; ++n) { + const int64_t ng = n * G + g; + const int64_t offset = n * C + c0 + d0; + + gm_h gdy(dy + offset, static_cast(N), static_cast(C)); + gm_h gx(x + offset, static_cast(N), static_cast(C)); + gm_f gmean(mean + ng, static_cast(N * G), 1); + gm_f grstd(rstd + ng, static_cast(N * G), 1); + + TLOAD(h0, gdy); + TCVT(dy_f, h0); + TLOAD(h0, gx); + TCVT(x_f, h0); + TLOAD(mean_t, gmean); + TLOAD(rstd_t, grstd); + + TROWEXPANDMUL(t0, x_f, rstd_t); + TMUL(t0, t0, dy_f); + TROWEXPANDMUL(x_f, dy_f, mean_t); + TROWEXPANDMUL(x_f, x_f, rstd_t); + TSUB(t0, t0, x_f); + TADD(acc, acc, t0); + } + + gm_h gdg(dgamma + c0 + d0, 1, static_cast(C)); + TCVT(h0, acc); + TSTORE(gdg, h0); + } +} + +} // namespace gn_grad_1d + +// tiling: [N, C, G, tile_d] +// +// 入口循环 ↔ Torch grid： +// for n,g fused_params + dx_group ↔ grid=dim3(N,G) 再接 numel 上 gpu_kernel +// for g dbeta/dgamma ↔ Kernel1/2 按通道写回 +template +void group_norm_grad_1d(dtype *dy, dtype *x, float *mean, float *rstd, + dtype *gamma, const int64_t *tiling, dtype *dx, + dtype *dgamma, dtype *dbeta) { + static_assert(peNum == 4, "normalization kernels support only 4PE"); + // Capacity in elements: every Tile buffer >= 512B (dtype strip + float strip). + constexpr int64_t tDDtype = + (512 + static_cast(sizeof(dtype)) - 1) / + static_cast(sizeof(dtype)); + constexpr int64_t tD = tDDtype > 128 ? tDDtype : 128; + constexpr int64_t tV = 1; // row-reduction/broadcast physical Columns=1 + + const int64_t N = tiling[0]; + const int64_t C = tiling[1]; + const int64_t G = tiling[2]; + const uint32_t tid = get_thread_idx(); + if (N <= 0 || C <= 0 || G <= 0 || C % G != 0 || + tid >= static_cast(peNum)) { + return; + } + const int64_t D = C / G; + const int64_t tile_d = tiling[3] > 0 ? tiling[3] : (D < tD ? D : tD); + if (tile_d <= 0 || tile_d > tD || D > tD) { + return; + } + + using gm_h = global_tensor>; + using gm_f = global_tensor>; + using tile_h = + Tile; + using tile_f = + Tile; + using tile_v = + Tile; + + const float s = 1.0f / static_cast(D); + float scratch[2]; // c2, c3 for one (n,g) + + for (int64_t ng = tid; ng < N * G; ng += peNum) { + const int64_t n = ng / G; + const int64_t g = ng % G; + gn_grad_1d::fused_params_group(dy, x, mean, rstd, gamma, + scratch, N, C, G, D, n, g, s); + gn_grad_1d::dx_group( + dy, x, rstd, gamma, scratch, dx, N, C, G, D, n, g); + } + + for (int64_t g = tid; g < G; g += peNum) { + gn_grad_1d::dbeta_group( + dy, dbeta, N, C, D, tile_d, g); + gn_grad_1d::dgamma_group( + dy, x, mean, rstd, dgamma, N, C, G, D, tile_d, g); + } +} + +#endif // SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/solution/normalization/rms_norm/rms_norm_pto.hpp b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm/rms_norm_pto.hpp new file mode 100644 index 00000000..4d46f8d4 --- /dev/null +++ b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm/rms_norm_pto.hpp @@ -0,0 +1,150 @@ +// ============================================================================= +// rms_norm_pto.hpp — RMSNorm (one-level PTO) +// ============================================================================= +// +// Shape dims: A (outer / row), R (reduce / col). +// +// out[a] = x[a] * rsqrt(mean(x[a]^2) + eps) +// +// Entry: +// rms_norm(x, tiling, out, eps); +// peNum defaults to 1; PE partitioning stays inside the kernel. +// tiling[4] = {g_a, g_r, tile_a, tile_r} (int64_t) +// tile_r <= 0 means use g_r (full-row tile). +// +// Pipeline (fp16 in/out, fp32 compute): +// TLOAD → TCVT → TMUL(x,x) → TROWSUM → TMULS(1/g_r) → TADDS(eps) +// → Newton rsqrt → TROWEXPANDMUL → TCVT → TSTORE +// +// Dynamic ValidRow/ValidCol: Tile Valid = -1, ctor passes runtime values. +// Full A tiles in the main loop; trailing rows handled separately. +// ============================================================================= +#ifndef SUPERNPU_RMS_NORM_PTO_HPP +#define SUPERNPU_RMS_NORM_PTO_HPP + +#include + +#include + +namespace rms_detail { + + +template +inline void rsqrt_newton(TileVec &out, TileVec &a) { + auto body = [&](auto &x, auto &t1, auto &t2) { + TRECIP(x, a); + for (int64_t i = 0; i < 4; ++i) { + TMUL(t1, x, x); + TMUL(t2, t1, a); + TMULS(t2, t2, -0.5f); + TADDS(t2, t2, 1.5f); + TMUL(x, x, t2); + } + TMULS(out, x, 1.0f); + }; + if constexpr (TileVec::ValidRow > 0) { + TileVec x, t1, t2; + body(x, t1, t2); + } else { + const size_t vr = static_cast(a.GetValidRow()); + TileVec x(vr), t1(vr), t2(vr); + body(x, t1, t2); + } +} + +template +inline void rms_norm_tile(dtype *x, dtype *out, int64_t gA, int64_t gR, + int64_t a_off, int64_t active_a, int64_t active_r, + float inv_r, float eps) { + const int64_t offset = a_off * gR; + gm_t gi(x + offset, static_cast(gA), static_cast(gR)); + gm_t go(out + offset, static_cast(gA), static_cast(gR)); + + tile_h src_h(static_cast(active_a), + static_cast(active_r)); + tile_h dst_h(static_cast(active_a), + static_cast(active_r)); + tile_f src(static_cast(active_a), + static_cast(active_r)); + tile_f squared(static_cast(active_a), + static_cast(active_r)); + tile_f dst(static_cast(active_a), + static_cast(active_r)); + tile_v sqrsum(static_cast(active_a)); + tile_v mean(static_cast(active_a)); + tile_v denom(static_cast(active_a)); + tile_v rms(static_cast(active_a)); + + TLOAD(src_h, gi); + TCVT(src, src_h); + TMUL(squared, src, src); + TROWSUM(sqrsum, squared); + TMULS(mean, sqrsum, inv_r); + TADDS(denom, mean, eps); + rsqrt_newton(rms, denom); + TROWEXPANDMUL(dst, src, rms); + TCVT(dst_h, dst); + TSTORE(go, dst_h); +} + +} // namespace rms_detail + +// tiling: [g_a, g_r, tile_a, tile_r] +template +void rms_norm(dtype *x, const int64_t *tiling, dtype *out, float eps = 1e-6f) { + static_assert(peNum == 4, "normalization kernels support only 4PE"); + + // Physical capacity (Rows×Cols); Valid comes from tiling (tile_a,tile_r). + // Size must cover ValidRow×ValidCol; SoftCore should not require Rows≥ValidRow. + constexpr int64_t tA = 1; + constexpr int64_t tR = 8192; + + const int64_t globalA = tiling[0]; + const int64_t gR = tiling[1]; + const int64_t tile_a = tiling[2] > 0 ? tiling[2] : tA; + const int64_t tile_r = tiling[3] > 0 ? tiling[3] : gR; + const uint32_t tid = get_thread_idx(); + + if (globalA <= 0 || gR <= 0 || tile_a <= 0 || tile_r <= 0 || + tile_r > tR || tid >= static_cast(peNum)) { + return; + } + + // Ceil partition: the first PEs take rows_per_pe rows and the last + // active PE owns the remainder. For M=333 and 4PE: 84, 84, 84, 81. + const int64_t rows_per_pe = (globalA + peNum - 1) / peNum; + const int64_t pe_start = static_cast(tid) * rows_per_pe; + if (pe_start >= globalA) { + return; + } + const int64_t remaining = globalA - pe_start; + const int64_t peA = + remaining < rows_per_pe ? remaining : rows_per_pe; + if (peA < tile_a) { + return; + } + const int64_t pe_offset = pe_start * gR; + x += pe_offset; + out += pe_offset; + + using gm_t = global_tensor>; + using tile_h = Tile; + using tile_f = Tile; + // Row-reduction output and row-broadcast input use physical Columns=1. + using tile_v = Tile; + + const float inv_r = 1.0f / static_cast(gR); + + // Full A tiles; peel the last iteration for the trailing block. + int64_t ia = 0; + for (; ia + tile_a < peA; ia += tile_a) { + rms_detail::rms_norm_tile( + x, out, peA, gR, ia, tile_a, tile_r, inv_r, eps); + } + // Tail (or sole) block: ValidRow = remaining rows along A. + rms_detail::rms_norm_tile( + x, out, peA, gR, ia, peA - ia, tile_r, inv_r, eps); +} + +#endif // SUPERNPU_RMS_NORM_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/SKILL.md b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/SKILL.md new file mode 100644 index 00000000..1df27a5b --- /dev/null +++ b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/SKILL.md @@ -0,0 +1,191 @@ +--- +name: rms-norm-binary +description: >- + Build, run, and debug the one-level rms_norm_binary kernel (R-split RMSNorm) + with SuperNPUBench + run_op.py + gfrun/gfsim precision checks. Use when + editing rms_norm_binary_pto.hpp, rms_norm_binary tests, workspace/GetCacheId + reduce, TADD cross-tile sum, or verifying [16,16384] fp16 binary RMSNorm. + Shape dims are A (outer) and R (reduce): g_a/g_r, tile_a/tile_r, tA/tR. +--- + +# rms_norm_binary — kernel & verification + +Default root: `/home/wangyu/Code/SuperScalar`（下文 `$ROOT`）。 + +## Shape naming + +| Old | New | Meaning | +|-----|-----|---------| +| M / `g_m` / `tM` / `tile_m` | **A** / `g_a` / `tA` / `tile_a` | outer / row | +| N / `g_n` / `tN` / `tile_n` | **R** / `g_r` / `tR` / `tile_r` | reduce / col | +| `Nb` | `Rb` | `# R-tiles = ceil(g_r / tile_r)` | + +`tiling[4] = {g_a, g_r, tile_a, tile_r}`. + +## What it is + +One-level PTO RMSNorm when `g_r > tile_r`: + +```text +out[a] = x[a] * rsqrt(mean(x[a]^2) + eps) +``` + +R is split into `Rb = ceil(g_r / tile_r)` tiles. Each tile does local +`TROWSUM(x^2)`, then tiles are reduced to a full-row sum. + +Current default test shape: **`[1, 8192]`**, `tile_r=1024` → **`Rb=8`**, fp16. + +## Key paths + +| Role | Path | +|------|------| +| Kernel | `$ROOT/SuperNPUBench/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp` | +| Reference (single-tile) | `.../kernels/solution/normalization/rms_norm/rms_norm_pto.hpp` | +| Testcase | `.../test/solution/normalization/rms_norm_binary/` | +| Host entry | `.../rms_norm_binary/src/rms_norm_binary.cpp` | +| Gen golden | `.../rms_norm_binary/src/gen_rms_norm_binary_data.py` | +| Compare | `.../rms_norm_binary/src/rms_norm_binary_data_compare.py` | +| Runner | `$ROOT/run_op.py` preset `rms_norm_binary` | +| Toolchain | `$ROOT/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin` → `COMPILER_DIR` | +| Sims | `$ROOT/SuperScalarModel/bin/gfrun`, `gfsim` | +| Related skill | `.../kernels/reduction/binary-accumulation-cache-id/SKILL.md` | + +ELF after build: + +```text +.../output/solution/normalization/rms_norm_binary/elf/ + solution_normalization_rms_norm_binary_rms_norm_binary_DType__half.elf +``` + +Compare dir (precision): + +```text +.../compare/solution_normalization_rms_norm_binary_rms_norm_binary_DType__half/ + input.bin golden.bin output.bin tiling_info.bin +``` + +## Kernel pipeline (current) + +File: `rms_norm_binary_pto.hpp`. **No `rms_norm_dyn_ops.hpp`.** TEPL style like +`rms_norm_pto.hpp`. + +```text +Pass1: + TLOAD(sum ← zeros) + per R-tile: TLOAD → TCVT → TMUL(x,x) → TROWSUM → TADD(sum, sum, cur) + +Pass1.5: + TMULS(mean, sum, 1/g_r) → TADDS(eps) → Newton rsqrt → rms + +Pass2 (per R-tile): + TLOAD → TCVT → TROWEXPANDMUL(x, rms) → TCVT → TSTORE +``` + +Important implementation notes: + +1. **Cross-tile sum is streaming** (`sum += cur`), not GetCacheId carry-merge. +2. Zero-init `sum` outside the R loop; uniform `TADD` inside (no first-tile branch). +3. `tile_v`: physical `Columns=1` and static `Valid=1,1`, matching the + hardware contract for `TROWSUM` output and later binary TEPL operations. +4. `workspace` stores one FP32 reduction value per row and cache level. Its + layout is `[kMaxLevels, g_a, kWsCols]` with `kWsCols=1`. +5. Do **not** put early-return parameter checks in the kernel (caller owns tiling). + +## How to verify + +```bash +export COMPILER_DIR=$ROOT/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin +cd $ROOT +python3 run_op.py rms_norm_binary +``` + +What `run_op.py` does: + +1. `gen_rms_norm_binary_data.py` → write `input.bin` / `golden.bin` / tiling +2. `make TESTCASE=rms_norm_binary DType=__half res_check=on` → ELF with I/O +3. `gfrun` functional sim → writes `output.bin` +4. `rms_norm_binary_data_compare.py` → atol/rtol vs golden +5. Rebuild **without** `res_check` (res_check ELF often crashes `gfsim`) +6. `gfsim` performance / cycle sim +7. After compile: write `.diss` via `llvm-objdump -dl` (disable: `--no-diss`) + +Useful flags: + +```bash +python3 run_op.py rms_norm_binary --func-only # gfrun + precision only +python3 run_op.py rms_norm_binary --perf-only # gfsim only +python3 run_op.py rms_norm_binary --compile-only +python3 run_op.py rms_norm_binary --no-check-precision +python3 run_op.py rms_norm_binary --skip-compile +python3 run_op.py rms_norm_binary --no-diss +``` + +Manual make (same case): + +```bash +cd $ROOT/SuperNPUBench/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary +make TESTCASE=rms_norm_binary DType=__half COMPILER_DIR=$COMPILER_DIR +# or: bash compile.all +``` + +## Expected results (as of current kernel) + +| Step | Typical result | +|------|----------------| +| Compile | OK | +| gfrun | **PASS** (`Success to Reach the End`) | +| Precision | **PASS** (`max_abs` often `0.0` on `[1,8192]`) | +| gfsim | **FAIL** | + +### gfsim failure (known) + +Symptom: + +```text +TMOV ... INVALID +FATAL: gfsim received signal 11 +# or: Bank store offset out of range! +``` + +Cause: compiler-inserted `TMOV` (tile rename) for small `tile_v` kept live across +R-split loops; timing sim corrupts tile metadata. **Not a golden mismatch** — +gfrun + compare already pass. + +Baseline `rms_norm` (no cross-tile accumulate / second R loop) usually **PASS**es gfsim. + +## Testcase layout + +`rms_norm_binary.cpp` defaults: + +```cpp +G_A=1, G_R=8192, TILE_A=1, TILE_R=1024 +workspace_buf[K_MAX_LEVELS * G_A * K_WS_COLS] // K_WS_COLS=1 +``` + +Precision scripts default shape `--g-r 8192`, `--tile-r 1024`. + +`run_op.py` preset name is exactly **`rms_norm_binary`** (no size suffix). + +## Agent checklist when changing the kernel + +1. Keep compute TEPL-only; do not reintroduce `rms_norm_dyn_ops.hpp` unless asked. +2. Keep `tile_v` physical Columns and Valid Columns both equal to 1 so + `TROWSUM`, workspace reload, and binary TEPL operands have identical layout. +3. Avoid taking addresses of `tile_v` / large pointer arrays of tiles (Liveouts / + illegal spill). +4. Do not mix `TROWSUM` u-reg lineage with `TLOAD` of small reduce tiles in the + same hot path without verifying Match Instruction / gfsim. +5. After edits: `python3 run_op.py rms_norm_binary` (or `--func-only` if only + checking correctness). +6. If implementing true GetCacheId carry + workspace reload, also read + `binary-accumulation-cache-id/SKILL.md` and expect toolchain/sim constraints + above. + +## Anti-patterns + +- Naming the run_op preset `rms_norm_binary_1x8192` / `..._1x32768` — canonical + name is `rms_norm_binary`. +- Treating gfsim FAIL as a precision bug when gfrun+compare already PASS. +- Putting workspace spill between `rsqrt` and `TROWEXPANDMUL` (clobbers `rms`). +- Using TEPL `TADD` with `Valid=-1` (`Match Instruction Error`). +- Reverting shape names to M/N — use **A/R** consistently. diff --git a/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp new file mode 100644 index 00000000..af8724bd --- /dev/null +++ b/benchmark/one-level-arch/kernels/solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp @@ -0,0 +1,260 @@ +// ============================================================================= +// rms_norm_binary_pto.hpp — RMSNorm for g_r > tile_r (R-split) +// ============================================================================= +// +// tiling[5] = {g_a, g_r, tile_a, tile_r, pow_r} +// +// 每块 RowSum 后立刻 UpdateCache（workspace = cacheBuffer），对齐 AscendC： +// DataCopy(aReg, src); +// for (j = 0; j < cid; ++j) { +// DataCopy(bReg, cache + j * stride); +// Add(aReg, aReg, bReg); +// } +// DataCopy(cache + cid * stride, aReg); +// cid = GetCacheId(idx) = ctz(idx+1) +// sum = cache[GetCacheId(r-1)] （r 为 2^k） +// +// workspace: [0, kMaxLevels) cache 档 +// ============================================================================= +#ifndef SUPERNPU_RMS_NORM_BINARY_PTO_HPP +#define SUPERNPU_RMS_NORM_BINARY_PTO_HPP + +#include + +#include + +namespace rms_bin { + +// Row-reduction results have physical Columns=1. Workspace cache entries +// must preserve that layout so TLOAD and TADD match the TROWSUM output. +constexpr int kWsCols = 1; +constexpr int kMaxLevels = 6; + + +inline int64_t GetCacheId(int64_t idx) { + return static_cast( + __builtin_ctzll(static_cast(idx + 1))); +} + +template +inline void rsqrt_newton(TileVec &out, TileVec &a) { + TileVec x, t1, t2; + TRECIP(x, a); + for (int64_t i = 0; i < 4; ++i) { + TMUL(t1, x, x); + TMUL(t2, t1, a); + TMULS(t2, t2, -0.5f); + TADDS(t2, t2, 1.5f); + TMUL(x, x, t2); + } + TMULS(out, x, 1.0f); +} + +} // namespace rms_bin + +template +void rms_norm_binary(dtype *x, const int64_t *tiling, dtype *out, + float *workspace, float eps = 1e-6f) { + static_assert(peNum == 4, "normalization kernels support only 4PE"); + constexpr int64_t tA = 1; + constexpr int64_t tR = 8192; + + const int64_t globalA = tiling[0]; + const int64_t gR = tiling[1]; + const int64_t tile_r = tiling[3] > 0 ? tiling[3] : tR; + const int64_t powR = tiling[4]; + const uint32_t tid = get_thread_idx(); + + if (globalA <= 0 || gR <= 1 || tile_r <= 0 || tile_r > tR || + powR <= 0 || powR >= gR || gR > 2 * powR || + tid >= static_cast(peNum)) { + return; + } + + // Ceil partition: for M=333 and 4PE, rows are split 84, 84, 84, 81. + const int64_t rows_per_pe = (globalA + peNum - 1) / peNum; + const int64_t pe_start = static_cast(tid) * rows_per_pe; + if (pe_start >= globalA) { + return; + } + const int64_t remaining = globalA - pe_start; + const int64_t gA = + remaining < rows_per_pe ? remaining : rows_per_pe; + if (gA < tA) { + return; + } + + const int64_t pe_offset = pe_start * gR; + x += pe_offset; + out += pe_offset; + // Workspace is level-major: [level][global row]. + workspace += pe_start * rms_bin::kWsCols; + + const int64_t remR = gR - powR; + const int64_t headR = powR - remR; + const int64_t n_rem_full = remR / tile_r; + const int64_t rem_tail = remR - n_rem_full * tile_r; + const int64_t n_head_full = headR / tile_r; + const int64_t head_tail = headR - n_head_full * tile_r; + const int64_t n_full = gR / tile_r; + const int64_t tail_r = gR - n_full * tile_r; + const float inv_r = 1.0f / static_cast(gR); + + using gm_t = global_tensor>; + using gm_f = global_tensor>; + using tile_h = Tile; + using tile_f = Tile; + using tile_v = Tile; + + for (int64_t ia = 0; ia < gA; ++ia) { + constexpr size_t active_a = 1; + const size_t full_r = static_cast(tile_r); + + tile_v cur, buf, sum, mean, denom, rms, zero; + TEXPANDS(zero, 0.0f); + + float *cache = workspace + ia * rms_bin::kWsCols; + const int64_t stride = globalA * rms_bin::kWsCols; + + for (int64_t lv = 0; lv < rms_bin::kMaxLevels; ++lv) { + gm_f go(cache + lv * stride, 1, rms_bin::kWsCols); + TSTORE(go, zero); + } + + int64_t r = 0; + + // UpdateCache（AscendC 同构） +#define RMS_BIN_UPDATE_CACHE() \ + do { \ + const uint16_t cid = \ + static_cast(rms_bin::GetCacheId(r)); \ + for (uint16_t j = 0; j < cid; ++j) { \ + gm_f gj(cache + static_cast(j) * stride, 1, \ + rms_bin::kWsCols); \ + TLOAD(buf, gj); \ + TADD(cur, cur, buf); \ + } \ + gm_f gc(cache + static_cast(cid) * stride, 1, \ + rms_bin::kWsCols); \ + TSTORE(gc, cur); \ + ++r; \ + } while (0) + + for (int64_t tr = 0; tr < n_rem_full; ++tr) { + const int64_t offset = ia * gR + tr * tile_r; + gm_t gi0(x + offset, static_cast(gA), static_cast(gR)); + gm_t gi1(x + offset + powR, static_cast(gA), + static_cast(gR)); + tile_h src0_h(active_a, full_r); + tile_h src1_h(active_a, full_r); + tile_f src0(active_a, full_r); + tile_f src1(active_a, full_r); + tile_f sq0(active_a, full_r); + tile_f sq1(active_a, full_r); + + TLOAD(src0_h, gi0); + TLOAD(src1_h, gi1); + TCVT(src0, src0_h); + TCVT(src1, src1_h); + TMUL(sq0, src0, src0); + TMUL(sq1, src1, src1); + TADD(sq0, sq0, sq1); + TROWSUM(cur, sq0); + RMS_BIN_UPDATE_CACHE(); + } + + if (rem_tail > 0) { + const int64_t offset = ia * gR + n_rem_full * tile_r; + const size_t ar = static_cast(rem_tail); + gm_t gi0(x + offset, static_cast(gA), static_cast(gR)); + gm_t gi1(x + offset + powR, static_cast(gA), + static_cast(gR)); + tile_h src0_h(active_a, ar); + tile_h src1_h(active_a, ar); + tile_f src0(active_a, ar); + tile_f src1(active_a, ar); + tile_f sq0(active_a, ar); + tile_f sq1(active_a, ar); + + TLOAD(src0_h, gi0); + TLOAD(src1_h, gi1); + TCVT(src0, src0_h); + TCVT(src1, src1_h); + TMUL(sq0, src0, src0); + TMUL(sq1, src1, src1); + TADD(sq0, sq0, sq1); + TROWSUM(cur, sq0); + RMS_BIN_UPDATE_CACHE(); + } + + for (int64_t tr = 0; tr < n_head_full; ++tr) { + const int64_t offset = ia * gR + remR + tr * tile_r; + gm_t gi(x + offset, static_cast(gA), static_cast(gR)); + tile_h src_h(active_a, full_r); + tile_f src(active_a, full_r); + tile_f sq(active_a, full_r); + TLOAD(src_h, gi); + TCVT(src, src_h); + TMUL(sq, src, src); + TROWSUM(cur, sq); + RMS_BIN_UPDATE_CACHE(); + } + if (head_tail > 0) { + const int64_t offset = ia * gR + remR + n_head_full * tile_r; + const size_t ar = static_cast(head_tail); + gm_t gi(x + offset, static_cast(gA), static_cast(gR)); + tile_h src_h(active_a, ar); + tile_f src(active_a, ar); + tile_f sq(active_a, ar); + TLOAD(src_h, gi); + TCVT(src, src_h); + TMUL(sq, src, src); + TROWSUM(cur, sq); + RMS_BIN_UPDATE_CACHE(); + } +#undef RMS_BIN_UPDATE_CACHE + + { + const int64_t rid = r > 0 ? rms_bin::GetCacheId(r - 1) : 0; + gm_f gr(cache + rid * stride, 1, rms_bin::kWsCols); + TLOAD(sum, gr); + } + + TMULS(mean, sum, inv_r); + TADDS(denom, mean, eps); + rms_bin::rsqrt_newton(rms, denom); + + for (int64_t tr = 0; tr < n_full; ++tr) { + const int64_t offset = ia * gR + tr * tile_r; + gm_t gi(x + offset, static_cast(gA), static_cast(gR)); + gm_t go(out + offset, static_cast(gA), static_cast(gR)); + tile_h src_h(active_a, full_r); + tile_h dst_h(active_a, full_r); + tile_f src(active_a, full_r); + tile_f dst(active_a, full_r); + TLOAD(src_h, gi); + TCVT(src, src_h); + TROWEXPANDMUL(dst, src, rms); + TCVT(dst_h, dst); + TSTORE(go, dst_h); + } + if (tail_r > 0) { + const int64_t offset = ia * gR + n_full * tile_r; + const size_t ar = static_cast(tail_r); + gm_t gi(x + offset, static_cast(gA), static_cast(gR)); + gm_t go(out + offset, static_cast(gA), static_cast(gR)); + tile_h src_h(active_a, ar); + tile_h dst_h(active_a, ar); + tile_f src(active_a, ar); + tile_f dst(active_a, ar); + TLOAD(src_h, gi); + TCVT(src, src_h); + TROWEXPANDMUL(dst, src, rms); + TCVT(dst_h, dst); + TSTORE(go, dst_h); + } + } +} + +#endif // SUPERNPU_RMS_NORM_BINARY_PTO_HPP diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/Makefile b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/Makefile deleted file mode 100644 index efdd2fb9..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/Makefile +++ /dev/null @@ -1,29 +0,0 @@ -N_BATCH ?= 2 -C_CH ?= 16 -G_GRP ?= 4 -HxW_SZ ?= 16 -TILE_HW ?= 8 -DType ?= __half - -ifeq ($(TESTCASE), group_norm_grad) -DEFINES += -DDType=$(DType) -DEFINES += -DN_BATCH=$(N_BATCH) -DEFINES += -DC_CH=$(C_CH) -DEFINES += -DG_GRP=$(G_GRP) -DEFINES += -DHxW_SZ=$(HxW_SZ) -DEFINES += -DTILE_HW=$(TILE_HW) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP)_HxW$(HxW_SZ).elf -endif - -ifeq ($(TESTCASE), group_norm_grad_static) -DEFINES += -DDType=$(DType) -DEFINES += -DN_BATCH=$(N_BATCH) -DEFINES += -DC_CH=$(C_CH) -DEFINES += -DG_GRP=$(G_GRP) -DEFINES += -DHxW_SZ=$(HxW_SZ) -DEFINES += -DTILE_HW=$(TILE_HW) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP)_HxW$(HxW_SZ)_tHW$(TILE_HW).elf -endif - -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp -include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/compile.all b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/compile.all deleted file mode 100755 index 82e05bf3..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/compile.all +++ /dev/null @@ -1,14 +0,0 @@ -#!/bin/bash -# GroupNorm Grad HxW>1: fp16; N=2 C=16 G=4 HxW=16 - -set -euo pipefail - -: "${COMPILER_DIR:?Set COMPILER_DIR to the linx_blockisa_llvm_musl bin directory}" - -make TESTCASE=group_norm_grad COMPILER_DIR="$COMPILER_DIR" \ - DType=__half N_BATCH=2 C_CH=16 G_GRP=4 HxW_SZ=16 TILE_HW=8 diss - -# Static Valid; same N/C/G/HxW/tile_hw -make TESTCASE=group_norm_grad_static COMPILER_DIR="$COMPILER_DIR" \ - DType=__half N_BATCH=2 C_CH=16 G_GRP=4 HxW_SZ=16 TILE_HW=8 diss - diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad.cpp b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad.cpp deleted file mode 100644 index 2c1a6788..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad.cpp +++ /dev/null @@ -1,91 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/group_norm_grad/group_norm_grad.hpp" - -#ifndef DType -#define DType __half -#endif - -// Default: HxW>1, N=2, C=16, G=4 → D=4, HxW=16 -#ifndef N_BATCH -#define N_BATCH 2 -#endif -#ifndef C_CH -#define C_CH 16 -#endif -#ifndef G_GRP -#define G_GRP 4 -#endif -#ifndef HxW_SZ -#define HxW_SZ 16 -#endif -#ifndef TILE_HW -#define TILE_HW 8 -#endif - -int main() { - using dtype = DType; - - // tiling: {N, C, G, HxW, tile_hw} - int64_t tiling_info[5] = {N_BATCH, C_CH, G_GRP, HxW_SZ, TILE_HW}; - - const int64_t N = tiling_info[0]; - const int64_t C = tiling_info[1]; - const int64_t G = tiling_info[2]; - const int64_t HxW = tiling_info[3]; - - constexpr int64_t kElems = N_BATCH * C_CH * HxW_SZ; - constexpr int64_t kWs = - 2 * N_BATCH * C_CH + 2 * N_BATCH * G_GRP; // ds+db+c2+c3 - - dtype dy_buf[kElems]; - dtype x_buf[kElems]; - float mean_buf[N_BATCH * G_GRP]; - float rstd_buf[N_BATCH * G_GRP]; - dtype gamma_buf[C_CH]; - dtype dx_buf[kElems]; - dtype dgamma_buf[C_CH]; - dtype dbeta_buf[C_CH]; - float workspace_buf[kWs]; - - dtype *dy = dy_buf; - dtype *x = x_buf; - float *mean = mean_buf; - float *rstd = rstd_buf; - dtype *gamma = gamma_buf; - dtype *dx = dx_buf; - dtype *dgamma = dgamma_buf; - dtype *dbeta = dbeta_buf; - float *workspace = workspace_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, - static_cast(kElems) * sizeof(dtype)); - readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, - static_cast(kElems) * sizeof(dtype)); - readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, - static_cast(N) * G * sizeof(float)); - readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, - static_cast(N) * G * sizeof(float)); - readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, - static_cast(C) * sizeof(dtype)); -#endif - - group_norm_grad(dy, x, mean, rstd, gamma, tiling_info, dx, dgamma, - dbeta, workspace); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, - static_cast(kElems) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, - static_cast(C) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, - static_cast(C) * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_static.cpp b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_static.cpp deleted file mode 100644 index 451fea10..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_static.cpp +++ /dev/null @@ -1,84 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/group_norm_grad/group_norm_grad.hpp" - -#ifndef DType -#define DType __half -#endif - -// Same as dynamic group_norm_grad.cpp: N=2 C=16 G=4 HxW=16 tile_hw=8 -#ifndef N_BATCH -#define N_BATCH 2 -#endif -#ifndef C_CH -#define C_CH 16 -#endif -#ifndef G_GRP -#define G_GRP 4 -#endif -#ifndef HxW_SZ -#define HxW_SZ 16 -#endif -#ifndef TILE_HW -#define TILE_HW 8 -#endif - -int main() { - using dtype = DType; - - constexpr int64_t kElems = static_cast(N_BATCH) * C_CH * HxW_SZ; - constexpr int64_t kWs = - 2 * static_cast(N_BATCH) * C_CH + - 2 * static_cast(N_BATCH) * G_GRP; - - dtype dy_buf[kElems]; - dtype x_buf[kElems]; - float mean_buf[N_BATCH * G_GRP]; - float rstd_buf[N_BATCH * G_GRP]; - dtype gamma_buf[C_CH]; - dtype dx_buf[kElems]; - dtype dgamma_buf[C_CH]; - dtype dbeta_buf[C_CH]; - float workspace_buf[kWs]; - - dtype *dy = dy_buf; - dtype *x = x_buf; - float *mean = mean_buf; - float *rstd = rstd_buf; - dtype *gamma = gamma_buf; - dtype *dx = dx_buf; - dtype *dgamma = dgamma_buf; - dtype *dbeta = dbeta_buf; - float *workspace = workspace_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, - static_cast(kElems) * sizeof(dtype)); - readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, - static_cast(kElems) * sizeof(dtype)); - readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, - static_cast(N_BATCH) * G_GRP * sizeof(float)); - readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, - static_cast(N_BATCH) * G_GRP * sizeof(float)); - readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, - static_cast(C_CH) * sizeof(dtype)); -#endif - - group_norm_grad( - dy, x, mean, rstd, gamma, dx, dgamma, dbeta, workspace); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, - static_cast(kElems) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, - static_cast(C_CH) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, - static_cast(C_CH) * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/Makefile b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/Makefile deleted file mode 100644 index c995ac55..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/Makefile +++ /dev/null @@ -1,30 +0,0 @@ -N_BATCH ?= 8 -C_CH ?= 64 -G_GRP ?= 8 -TILE_D ?= -1 -DType ?= __half - -ifeq ($(TESTCASE), group_norm_grad_1d) -DEFINES += -DDType=$(DType) -DEFINES += -DN_BATCH=$(N_BATCH) -DEFINES += -DC_CH=$(C_CH) -DEFINES += -DG_GRP=$(G_GRP) -DEFINES += -DTILE_D=$(TILE_D) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP).elf -endif - -ifeq ($(TESTCASE), group_norm_grad_1d_static) -GN1D_TILE_D := $(TILE_D) -ifeq ($(GN1D_TILE_D),-1) -GN1D_TILE_D := 8 -endif -DEFINES += -DDType=$(DType) -DEFINES += -DN_BATCH=$(N_BATCH) -DEFINES += -DC_CH=$(C_CH) -DEFINES += -DG_GRP=$(G_GRP) -DEFINES += -DTILE_D=$(GN1D_TILE_D) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP)_tD$(GN1D_TILE_D).elf -endif - -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp -include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/compile.all b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/compile.all deleted file mode 100755 index 273cf201..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/compile.all +++ /dev/null @@ -1,14 +0,0 @@ -#!/bin/bash -# GroupNorm Grad HxW==1: fp16; shape N=8 C=64 G=8 (D=8) - -set -euo pipefail - -: "${COMPILER_DIR:?Set COMPILER_DIR to the linx_blockisa_llvm_musl bin directory}" - -make TESTCASE=group_norm_grad_1d COMPILER_DIR="$COMPILER_DIR" \ - DType=__half N_BATCH=8 C_CH=64 G_GRP=8 TILE_D=-1 diss - -# Static Valid; tile_d=-1 → D=C/G=8 -make TESTCASE=group_norm_grad_1d_static COMPILER_DIR="$COMPILER_DIR" \ - DType=__half N_BATCH=8 C_CH=64 G_GRP=8 TILE_D=8 diss - diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp deleted file mode 100644 index 85116203..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp +++ /dev/null @@ -1,81 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp" - -#ifndef DType -#define DType __half -#endif - -// Default shape: HxW==1, N=8, C=64, G=8 → D=8 -#ifndef N_BATCH -#define N_BATCH 8 -#endif -#ifndef C_CH -#define C_CH 64 -#endif -#ifndef G_GRP -#define G_GRP 8 -#endif -#ifndef TILE_D -#define TILE_D -1 -#endif - -int main() { - using dtype = DType; - - // tiling: {N, C, G, tile_d} - int64_t tiling_info[4] = {N_BATCH, C_CH, G_GRP, TILE_D}; - - const int64_t N = tiling_info[0]; - const int64_t C = tiling_info[1]; - const int64_t G = tiling_info[2]; - - dtype dy_buf[N_BATCH * C_CH]; - dtype x_buf[N_BATCH * C_CH]; - float mean_buf[N_BATCH * G_GRP]; - float rstd_buf[N_BATCH * G_GRP]; - dtype gamma_buf[C_CH]; - dtype dx_buf[N_BATCH * C_CH]; - dtype dgamma_buf[C_CH]; - dtype dbeta_buf[C_CH]; - - dtype *dy = dy_buf; - dtype *x = x_buf; - float *mean = mean_buf; - float *rstd = rstd_buf; - dtype *gamma = gamma_buf; - dtype *dx = dx_buf; - dtype *dgamma = dgamma_buf; - dtype *dbeta = dbeta_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, - static_cast(N) * C * sizeof(dtype)); - readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, - static_cast(N) * C * sizeof(dtype)); - readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, - static_cast(N) * G * sizeof(float)); - readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, - static_cast(N) * G * sizeof(float)); - readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, - static_cast(C) * sizeof(dtype)); -#endif - - group_norm_grad_1d(dy, x, mean, rstd, gamma, tiling_info, dx, - dgamma, dbeta); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, - static_cast(N) * C * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, - static_cast(C) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, - static_cast(C) * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_static.cpp b/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_static.cpp deleted file mode 100644 index 7e89f488..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_static.cpp +++ /dev/null @@ -1,74 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp" - -#ifndef DType -#define DType __half -#endif - -// Same as dynamic group_norm_grad_1d.cpp: N=8 C=64 G=8, tile_d=-1 → D=8 -#ifndef N_BATCH -#define N_BATCH 8 -#endif -#ifndef C_CH -#define C_CH 64 -#endif -#ifndef G_GRP -#define G_GRP 8 -#endif -#ifndef TILE_D -#define TILE_D 8 -#endif - -int main() { - using dtype = DType; - - dtype dy_buf[N_BATCH * C_CH]; - dtype x_buf[N_BATCH * C_CH]; - float mean_buf[N_BATCH * G_GRP]; - float rstd_buf[N_BATCH * G_GRP]; - dtype gamma_buf[C_CH]; - dtype dx_buf[N_BATCH * C_CH]; - dtype dgamma_buf[C_CH]; - dtype dbeta_buf[C_CH]; - - dtype *dy = dy_buf; - dtype *x = x_buf; - float *mean = mean_buf; - float *rstd = rstd_buf; - dtype *gamma = gamma_buf; - dtype *dx = dx_buf; - dtype *dgamma = dgamma_buf; - dtype *dbeta = dbeta_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, - static_cast(N_BATCH) * C_CH * sizeof(dtype)); - readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, - static_cast(N_BATCH) * C_CH * sizeof(dtype)); - readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, - static_cast(N_BATCH) * G_GRP * sizeof(float)); - readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, - static_cast(N_BATCH) * G_GRP * sizeof(float)); - readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, - static_cast(C_CH) * sizeof(dtype)); -#endif - - group_norm_grad_1d( - dy, x, mean, rstd, gamma, dx, dgamma, dbeta); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, - static_cast(N_BATCH) * C_CH * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, - static_cast(C_CH) * sizeof(dtype)); - writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, - static_cast(C_CH) * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/Makefile b/benchmark/one-level-arch/test/kernel/normalization/rms_norm/Makefile deleted file mode 100644 index 6647e936..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/Makefile +++ /dev/null @@ -1,30 +0,0 @@ -DType ?= __half -G_A ?= 16 -G_R ?= 512 -TILE_A ?= 2 -TILE_R ?= 512 -# 1 = single PE (default). 4 = split A across PEs via get_thread_idx(); -# gfrun needs -s softcore.multiThreadNum=4 (same as multi_thread/vec). -PE_NUM ?= 1 - -ifeq ($(TESTCASE), rms_norm) -DEFINES += -DDType=$(DType) -DPE_NUM=$(PE_NUM) -ifeq ($(PE_NUM),1) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType).elf -else -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_PE$(PE_NUM).elf -endif -endif - -ifeq ($(TESTCASE), rms_norm_static) -DEFINES += -DDType=$(DType) -DPE_NUM=$(PE_NUM) -DEFINES += -DG_A=$(G_A) -DG_R=$(G_R) -DTILE_A=$(TILE_A) -DTILE_R=$(TILE_R) -ifeq ($(PE_NUM),1) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_tA$(TILE_A)_tR$(TILE_R).elf -else -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_tA$(TILE_A)_tR$(TILE_R)_PE$(PE_NUM).elf -endif -endif - -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp -include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/compile.all b/benchmark/one-level-arch/test/kernel/normalization/rms_norm/compile.all deleted file mode 100755 index 11c3a606..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/compile.all +++ /dev/null @@ -1,21 +0,0 @@ -#!/bin/bash -# RMSNorm one-level: fp16 -# rms_norm.cpp tiling_info = {16,512,1,-1} (g_a,g_r,tile_a,tile_r; tile_r<0 => full row) -# rms_norm_static: G_A=16 G_R=512 TILE_A=2 TILE_R=512 -# PE_NUM=4 splits A (16/4=4 rows/PE). Run those ELFs with: -# gfrun -f -s softcore.multiThreadNum=4 - -set -euo pipefail - -: "${COMPILER_DIR:?Set COMPILER_DIR to the linx_blockisa_llvm_musl bin directory}" - -make TESTCASE=rms_norm COMPILER_DIR="$COMPILER_DIR" DType=__half diss - -# Static Valid (compile-time G_A/G_R/TILE_A/TILE_R) -make TESTCASE=rms_norm_static COMPILER_DIR="$COMPILER_DIR" DType=__half \ - G_A=16 G_R=512 TILE_A=2 TILE_R=512 diss - -# 4PE: split A (16/4=4 rows per PE) -make TESTCASE=rms_norm COMPILER_DIR="$COMPILER_DIR" DType=__half PE_NUM=4 diss -make TESTCASE=rms_norm_static COMPILER_DIR="$COMPILER_DIR" DType=__half \ - G_A=16 G_R=512 TILE_A=2 TILE_R=512 PE_NUM=4 diss diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm.cpp b/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm.cpp deleted file mode 100644 index af9135d9..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm.cpp +++ /dev/null @@ -1,65 +0,0 @@ -#include - -#include - -#include "fileop.h" - -#ifndef PE_NUM -#define PE_NUM 1 -#endif - -#if PE_NUM == 1 -#include "single_thread/normalization/rms_norm/rms_norm.hpp" -#else -#include "multi_thread/normalization/rms_norm/rms_norm.hpp" -#endif - -#ifndef DType -#define DType __half -#endif - -#ifndef EPS -#define EPS 1e-6f -#endif - -int main() { - using dtype = DType; - - static_assert(PE_NUM > 0, "PE_NUM must be positive"); - static_assert(16 % PE_NUM == 0, "g_a=16 must be divisible by PE_NUM"); - - // tiling_info: {g_a, g_r, tile_a, tile_r} — host-visible full A. - // 4PE: split A across PEs like multi_thread/vec (get_thread_idx). - int64_t tiling_info[4] = {16, 512, 1, -1}; - - const int64_t g_a = tiling_info[0]; - const int64_t g_r = tiling_info[1]; - constexpr int64_t pe_a = 16 / PE_NUM; - - dtype input_buf[16 * 512]; - dtype output_buf[16 * 512]; - dtype *input = input_buf; - dtype *output = output_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, - static_cast(g_a) * g_r * sizeof(dtype)); -#endif - -#if PE_NUM == 1 - rms_norm(input, tiling_info, output, EPS); -#else - const uint32_t tid = get_thread_idx(); - int64_t tiling_pe[4] = {pe_a, g_r, tiling_info[2], tiling_info[3]}; - rms_norm(input + tid * pe_a * g_r, tiling_pe, - output + tid * pe_a * g_r, EPS); -#endif - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, - static_cast(g_a) * g_r * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_static.cpp b/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_static.cpp deleted file mode 100644 index 4ffaa376..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_static.cpp +++ /dev/null @@ -1,72 +0,0 @@ -#include - -#include - -#include "fileop.h" - -#ifndef PE_NUM -#define PE_NUM 1 -#endif - -#if PE_NUM == 1 -#include "single_thread/normalization/rms_norm/rms_norm.hpp" -#else -#include "multi_thread/normalization/rms_norm/rms_norm.hpp" -#endif - -#ifndef DType -#define DType __half -#endif - -#ifndef EPS -#define EPS 1e-6f -#endif - -// Same as dynamic rms_norm.cpp tiling_info {16,512,2,512} -#ifndef G_A -#define G_A 16 -#endif -#ifndef G_R -#define G_R 512 -#endif -#ifndef TILE_A -#define TILE_A 2 -#endif -#ifndef TILE_R -#define TILE_R 512 -#endif - -int main() { - using dtype = DType; - - static_assert(PE_NUM > 0, "PE_NUM must be positive"); - static_assert(G_A % PE_NUM == 0, "G_A must be divisible by PE_NUM"); - static_assert((G_A / PE_NUM) >= TILE_A, "PE-local G_A must cover one tile_a"); - - constexpr int pe_a = G_A / PE_NUM; - - dtype input_buf[G_A * G_R]; - dtype output_buf[G_A * G_R]; - dtype *input = input_buf; - dtype *output = output_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, - static_cast(G_A) * G_R * sizeof(dtype)); -#endif - -#if PE_NUM == 1 - rms_norm(input, output, EPS); -#else - // Full [G_A, G_R] buffers; kernel splits A with get_thread_idx(). - rms_norm(input, output, EPS); -#endif - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, - static_cast(G_A) * G_R * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/Makefile b/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/Makefile deleted file mode 100644 index 073485b0..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/Makefile +++ /dev/null @@ -1,21 +0,0 @@ -DType ?= __half -G_A ?= 1 -G_R ?= 8192 -TILE_A ?= 1 -TILE_R ?= 1024 -POW_R ?= 4096 - -ifeq ($(TESTCASE), rms_norm_binary) -DEFINES += -DDType=$(DType) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType).elf -endif - -ifeq ($(TESTCASE), rms_norm_binary_static) -DEFINES += -DDType=$(DType) -DEFINES += -DG_A=$(G_A) -DG_R=$(G_R) -DTILE_A=$(TILE_A) -DTILE_R=$(TILE_R) -DEFINES += -DPOW_R=$(POW_R) -TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_tA$(TILE_A)_tR$(TILE_R)_powR$(POW_R).elf -endif - -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp -include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/compile.all b/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/compile.all deleted file mode 100755 index d0339e67..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/compile.all +++ /dev/null @@ -1,12 +0,0 @@ -#!/bin/bash -# RMSNorm binary-accumulate: [1,8192] fp16, tile_r=1024, pow_r=4096 -set -euo pipefail - -: "${COMPILER_DIR:?Set COMPILER_DIR to the linx_blockisa_llvm_musl bin directory}" - -make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half diss - -# Static Valid; same [1,8192] tile_r=1024 pow_r=4096 -make TESTCASE=rms_norm_binary_static COMPILER_DIR="$COMPILER_DIR" DType=__half \ - G_A=1 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096 diss - diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary.cpp b/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary.cpp deleted file mode 100644 index 7ec5fff7..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary.cpp +++ /dev/null @@ -1,71 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp" - -#ifndef DType -#define DType __half -#endif - -#ifndef EPS -#define EPS 1e-6f -#endif - -// Spec: shape [g_a, g_r] = [1, 8192], fp16 -// tiling: tile_a=1, tile_r=1024, pow_r=4096 -// pow_r is 2^n and pow_r < g_r <= 2*pow_r -#ifndef G_A -#define G_A 1 -#endif -#ifndef G_R -#define G_R 8192 -#endif -#ifndef TILE_A -#define TILE_A 1 -#endif -#ifndef TILE_R -#define TILE_R 1024 -#endif -#ifndef POW_R -#define POW_R 4096 -#endif -// Must match rms_bin::kWsCols / kMaxLevels -#ifndef K_WS_COLS -#define K_WS_COLS 128 -#endif -#ifndef K_MAX_LEVELS -#define K_MAX_LEVELS 6 -#endif - -int main() { - using dtype = DType; - - int64_t tiling_info[5] = {G_A, G_R, TILE_A, TILE_R, POW_R}; - - const int64_t g_a = tiling_info[0]; - const int64_t g_r = tiling_info[1]; - - dtype input_buf[G_A * G_R]; - dtype output_buf[G_A * G_R]; - float workspace_buf[K_MAX_LEVELS * G_A * K_WS_COLS]; - dtype *input = input_buf; - dtype *output = output_buf; - float *workspace = workspace_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, - static_cast(g_a) * g_r * sizeof(dtype)); -#endif - - rms_norm_binary(input, tiling_info, output, workspace, EPS); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, - static_cast(g_a) * g_r * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_static.cpp b/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_static.cpp deleted file mode 100644 index c722118b..00000000 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_static.cpp +++ /dev/null @@ -1,64 +0,0 @@ -#include - -#include - -#include "fileop.h" -#include "single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp" - -#ifndef DType -#define DType __half -#endif - -#ifndef EPS -#define EPS 1e-6f -#endif - -// Same as dynamic rms_norm_binary.cpp: [1,8192], tile_r=1024, pow_r=4096 -#ifndef G_A -#define G_A 1 -#endif -#ifndef G_R -#define G_R 8192 -#endif -#ifndef TILE_A -#define TILE_A 1 -#endif -#ifndef TILE_R -#define TILE_R 1024 -#endif -#ifndef POW_R -#define POW_R 4096 -#endif -#ifndef K_WS_COLS -#define K_WS_COLS 128 -#endif -#ifndef K_MAX_LEVELS -#define K_MAX_LEVELS 6 -#endif - -int main() { - using dtype = DType; - - dtype input_buf[G_A * G_R]; - dtype output_buf[G_A * G_R]; - float workspace_buf[K_MAX_LEVELS * G_A * K_WS_COLS]; - dtype *input = input_buf; - dtype *output = output_buf; - float *workspace = workspace_buf; - -#ifdef RES_CHECK -#ifndef CHK_DIR -#error "CHK_DIR must be set when RES_CHECK is enabled" -#endif - readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, - static_cast(G_A) * G_R * sizeof(dtype)); -#endif - - rms_norm_binary(input, output, - workspace, EPS); - -#ifdef RES_CHECK - writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, - static_cast(G_A) * G_R * sizeof(dtype)); -#endif -} diff --git a/benchmark/one-level-arch/test/solution/.gitkeep b/benchmark/one-level-arch/test/solution/.gitkeep deleted file mode 100644 index e69de29b..00000000 diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/Makefile b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/Makefile new file mode 100644 index 00000000..66de46c3 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/Makefile @@ -0,0 +1,17 @@ +N_BATCH ?= 32 +C_CH ?= 16 +G_GRP ?= 8 +HxW_SZ ?= 8192 +DType ?= __half +PE_NUM ?= 4 +ifneq ($(PE_NUM),4) +$(error group_norm_grad supports only the dynamic 4PE testcase) +endif +START_FILE = $(ROOT)/test/common/_start_multipe.s +ifeq ($(TESTCASE), group_norm_grad) +DEFINES += -DDType=$(DType) -DN_BATCH=$(N_BATCH) -DC_CH=$(C_CH) +DEFINES += -DG_GRP=$(G_GRP) -DHxW_SZ=$(HxW_SZ) -DPE_NUM=$(PE_NUM) +TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP)_HxW$(HxW_SZ)_PE$(PE_NUM).elf +endif +SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp +include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/compile.all b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/compile.all new file mode 100755 index 00000000..2b20ca7b --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/compile.all @@ -0,0 +1,5 @@ +#!/bin/bash +# Dynamic 4PE GroupNormGrad N=32,C=16,G=8,HxW=8192; tiling is calculated from runtime shape. +set -euo pipefail +: "${COMPILER_DIR:?Set COMPILER_DIR to the compiler bin directory}" +make TESTCASE=group_norm_grad COMPILER_DIR="$COMPILER_DIR" DType=__half N_BATCH=32 C_CH=16 G_GRP=8 HxW_SZ=8192 PE_NUM=4 diss diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/gen_group_norm_grad_data.py b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/gen_group_norm_grad_data.py similarity index 94% rename from benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/gen_group_norm_grad_data.py rename to benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/gen_group_norm_grad_data.py index f1de96d7..50c845f7 100644 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/gen_group_norm_grad_data.py +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/gen_group_norm_grad_data.py @@ -9,7 +9,7 @@ golden_dx / golden_dgamma / golden_dbeta : float16 Math matches PyTorch GroupNormBackward (HxW>1): spatial ds/db then fused c2/c3. -Default: N=2, C=16, G=4, HxW=16 (D=4). +Default: N=32, C=16, G=8, HxW=8192 (D=2), tile_hw=8192. """ from __future__ import annotations @@ -22,10 +22,10 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_group_norm_grad_group_norm_grad" - "_DType__half_N2_C16_G4_HxW16" + / "solution_normalization_group_norm_grad_group_norm_grad" + "_DType__half_N32_C16_G8_HxW8192_PE4" ) @@ -228,11 +228,11 @@ def gen_all( def main() -> None: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument("--n", type=int, default=2) + parser.add_argument("--n", type=int, default=32) parser.add_argument("--c", type=int, default=16) - parser.add_argument("--g", type=int, default=4) - parser.add_argument("--hxw", type=int, default=16) - parser.add_argument("--tile-hw", type=int, default=8) + parser.add_argument("--g", type=int, default=8) + parser.add_argument("--hxw", type=int, default=8192) + parser.add_argument("--tile-hw", type=int, default=8192) parser.add_argument("--eps", type=float, default=1e-5) parser.add_argument("--seed", type=int, default=123) parser.add_argument("-o", "--out-dir", type=Path, default=DEFAULT_CMP_DIR) diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad.cpp b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad.cpp new file mode 100644 index 00000000..dcc0bf09 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad.cpp @@ -0,0 +1,128 @@ +#include + +#include + +#include "fileop.h" +#include "solution/normalization/group_norm_grad/group_norm_grad_pto.hpp" + +#ifndef DType +#define DType __half +#endif + +// Dynamic 4PE validation: N=32, C=16, G=8, HxW=8192. +#ifndef N_BATCH +#define N_BATCH 32 +#endif +#ifndef C_CH +#define C_CH 16 +#endif +#ifndef G_GRP +#define G_GRP 8 +#endif +#ifndef HxW_SZ +#define HxW_SZ 8192 +#endif +#ifndef PE_NUM +#define PE_NUM 1 +#endif + +namespace { +template +constexpr int64_t group_norm_tile_hw(int64_t spatial_size) { + constexpr int64_t kTileCapacity = 8192; + return spatial_size < kTileCapacity ? spatial_size : kTileCapacity; +} +} // namespace + +#ifdef RES_CHECK +namespace { +volatile uint32_t input_ready = 0; +volatile uint32_t kernel_done[PE_NUM] = {}; +volatile uint32_t output_written = 0; +} // namespace +#endif + +int main() { + using dtype = DType; + + // tiling: {N, C, G, HxW, tile_hw} + constexpr int64_t kTileHw = group_norm_tile_hw(HxW_SZ); + static_assert(N_BATCH > 0 && C_CH > 0 && G_GRP > 0 && HxW_SZ > 0); + static_assert(C_CH % G_GRP == 0 && kTileHw > 0); + int64_t tiling_info[5] = {N_BATCH, C_CH, G_GRP, HxW_SZ, kTileHw}; + + const int64_t N = tiling_info[0]; + const int64_t C = tiling_info[1]; + const int64_t G = tiling_info[2]; + const int64_t HxW = tiling_info[3]; + + constexpr int64_t kElems = N_BATCH * C_CH * HxW_SZ; + constexpr int64_t kWs = + 2 * N_BATCH * C_CH + 2 * N_BATCH * G_GRP; + + static dtype dy_buf[kElems]; + static dtype x_buf[kElems]; + static float mean_buf[N_BATCH * G_GRP]; + static float rstd_buf[N_BATCH * G_GRP]; + static dtype gamma_buf[C_CH]; + static dtype dx_buf[kElems]; + static dtype dgamma_buf[C_CH]; + static dtype dbeta_buf[C_CH]; + static float workspace_buf[kWs]; + + dtype *dy = dy_buf; + dtype *x = x_buf; + float *mean = mean_buf; + float *rstd = rstd_buf; + dtype *gamma = gamma_buf; + dtype *dx = dx_buf; + dtype *dgamma = dgamma_buf; + dtype *dbeta = dbeta_buf; + float *workspace = workspace_buf; + +#ifdef RES_CHECK +#ifndef CHK_DIR +#error "CHK_DIR must be set when RES_CHECK is enabled" +#endif + const uint32_t tid = gn_grad::read_pe_id(); + if (tid == 0) { + readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, + static_cast(kElems) * sizeof(dtype)); + readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, + static_cast(kElems) * sizeof(dtype)); + readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, + static_cast(N) * G * sizeof(float)); + readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, + static_cast(N) * G * sizeof(float)); + readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, + static_cast(C) * sizeof(dtype)); + input_ready = 1; + } else { + while (input_ready == 0) { + } + } +#endif + + group_norm_grad(dy, x, mean, rstd, gamma, tiling_info, dx, + dgamma, dbeta, workspace); + +#ifdef RES_CHECK + kernel_done[tid] = 1; + if (tid == 0) { + for (int pe = 0; pe < PE_NUM; ++pe) { + while (kernel_done[pe] == 0) { + } + } + writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, + static_cast(kElems) * sizeof(dtype)); + writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, + static_cast(C) * sizeof(dtype)); + writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, + static_cast(C) * sizeof(dtype)); + output_written = 1; + } else { + while (output_written == 0) { + } + } +#endif +} diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_data_compare.py b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad_data_compare.py similarity index 91% rename from benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_data_compare.py rename to benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad_data_compare.py index 16f3c70a..3d71c62b 100644 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad/src/group_norm_grad_data_compare.py +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad/src/group_norm_grad_data_compare.py @@ -12,10 +12,10 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_group_norm_grad_group_norm_grad" - "_DType__half_N2_C16_G4_HxW16" + / "solution_normalization_group_norm_grad_group_norm_grad" + "_DType__half_N32_C16_G8_HxW8192_PE4" ) @@ -77,11 +77,13 @@ def compare_one( rel_diffs.append(d / denom) mse = se / n if n else 0.0 + ref_mean_square = sum(b * b for b in ref) / n if n else 0.0 + mse_limit = mse_tol * max(ref_mean_square, 1.0) max_abs = max(abs_diffs) if abs_diffs else 0.0 mean_abs = sum(abs_diffs) / n if n else 0.0 max_rel = max(rel_diffs) if rel_diffs else 0.0 within = all((d <= atol) or (r <= rtol) for d, r in zip(abs_diffs, rel_diffs)) - mse_ok = mse <= mse_tol + mse_ok = mse <= mse_limit status = "pass" if within and mse_ok else "fail" return { "name": name, @@ -90,6 +92,8 @@ def compare_one( "max_abs": max_abs, "mean_abs": mean_abs, "mse": mse, + "ref_mean_square": ref_mean_square, + "mse_limit": mse_limit, "max_rel": max_rel, "atol": atol, "rtol": rtol, diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/Makefile b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/Makefile new file mode 100644 index 00000000..c2a27c79 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/Makefile @@ -0,0 +1,16 @@ +N_BATCH ?= 512 +C_CH ?= 64 +G_GRP ?= 8 +DType ?= __half +PE_NUM ?= 4 +ifneq ($(PE_NUM),4) +$(error group_norm_grad_1d supports only the dynamic 4PE testcase) +endif +START_FILE = $(ROOT)/test/common/_start_multipe.s +ifeq ($(TESTCASE), group_norm_grad_1d) +DEFINES += -DDType=$(DType) -DN_BATCH=$(N_BATCH) -DC_CH=$(C_CH) +DEFINES += -DG_GRP=$(G_GRP) -DPE_NUM=$(PE_NUM) +TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_N$(N_BATCH)_C$(C_CH)_G$(G_GRP)_PE$(PE_NUM).elf +endif +SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp +include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/compile.all b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/compile.all new file mode 100755 index 00000000..6a83733d --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/compile.all @@ -0,0 +1,5 @@ +#!/bin/bash +# Dynamic 4PE GroupNormGrad1D N=512,C=64,G=8,D=8; tiling is calculated from runtime shape. +set -euo pipefail +: "${COMPILER_DIR:?Set COMPILER_DIR to the compiler bin directory}" +make TESTCASE=group_norm_grad_1d COMPILER_DIR="$COMPILER_DIR" DType=__half N_BATCH=512 C_CH=64 G_GRP=8 PE_NUM=4 diss diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py similarity index 96% rename from benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py rename to benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py index 702f0d91..72c636eb 100644 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/gen_group_norm_grad_1d_data.py @@ -9,7 +9,7 @@ golden_dx.bin / golden_dgamma.bin / golden_dbeta.bin : float16 Math matches PyTorch GroupNorm1dBackward (fp32 accumulate, cast to fp16). -Default: N=8, C=64, G=8 (D=8), tile_d=-1. +Default: N=512, C=64, G=8 (D=8), tile_d=8. """ from __future__ import annotations @@ -22,9 +22,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_group_norm_grad_1d_group_norm_grad_1d_DType__half_N8_C64_G8" + / "solution_normalization_group_norm_grad_1d_group_norm_grad_1d_DType__half_N512_C64_G8_PE4" ) @@ -204,10 +204,10 @@ def gen_all( def main() -> None: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument("--n", type=int, default=8) + parser.add_argument("--n", type=int, default=512) parser.add_argument("--c", type=int, default=64) parser.add_argument("--g", type=int, default=8) - parser.add_argument("--tile-d", type=int, default=-1) + parser.add_argument("--tile-d", type=int, default=8) parser.add_argument("--eps", type=float, default=1e-5) parser.add_argument("--seed", type=int, default=123) parser.add_argument("-o", "--out-dir", type=Path, default=DEFAULT_CMP_DIR) diff --git a/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp new file mode 100644 index 00000000..06562d22 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d.cpp @@ -0,0 +1,123 @@ +#include + +#include + +#include "fileop.h" +#include "solution/normalization/group_norm_grad_1d/group_norm_grad_1d_pto.hpp" + +#ifndef DType +#define DType __half +#endif + +// Dynamic 4PE validation: HxW==1, N=512, C=64, G=8, D=8. +#ifndef N_BATCH +#define N_BATCH 512 +#endif +#ifndef C_CH +#define C_CH 64 +#endif +#ifndef G_GRP +#define G_GRP 8 +#endif +#ifndef PE_NUM +#define PE_NUM 1 +#endif + +namespace { +template +constexpr int64_t group_norm_1d_tile_d(int64_t channels, int64_t groups) { + constexpr int64_t kDtypeCapacity = + (512 + static_cast(sizeof(dtype)) - 1) / + static_cast(sizeof(dtype)); + constexpr int64_t kTileCapacity = + kDtypeCapacity > 128 ? kDtypeCapacity : 128; + const int64_t group_width = channels / groups; + return group_width < kTileCapacity ? group_width : kTileCapacity; +} +} // namespace + +#ifdef RES_CHECK +namespace { +volatile uint32_t input_ready = 0; +volatile uint32_t kernel_done[PE_NUM] = {}; +volatile uint32_t output_written = 0; +} // namespace +#endif + +int main() { + using dtype = DType; + + // tiling: {N, C, G, tile_d} + constexpr int64_t kTileD = group_norm_1d_tile_d(C_CH, G_GRP); + static_assert(N_BATCH > 0 && C_CH > 0 && G_GRP > 0); + static_assert(C_CH % G_GRP == 0 && kTileD > 0); + int64_t tiling_info[4] = {N_BATCH, C_CH, G_GRP, kTileD}; + + const int64_t N = tiling_info[0]; + const int64_t C = tiling_info[1]; + const int64_t G = tiling_info[2]; + + static dtype dy_buf[N_BATCH * C_CH]; + static dtype x_buf[N_BATCH * C_CH]; + static float mean_buf[N_BATCH * G_GRP]; + static float rstd_buf[N_BATCH * G_GRP]; + static dtype gamma_buf[C_CH]; + static dtype dx_buf[N_BATCH * C_CH]; + static dtype dgamma_buf[C_CH]; + static dtype dbeta_buf[C_CH]; + + dtype *dy = dy_buf; + dtype *x = x_buf; + float *mean = mean_buf; + float *rstd = rstd_buf; + dtype *gamma = gamma_buf; + dtype *dx = dx_buf; + dtype *dgamma = dgamma_buf; + dtype *dbeta = dbeta_buf; + +#ifdef RES_CHECK +#ifndef CHK_DIR +#error "CHK_DIR must be set when RES_CHECK is enabled" +#endif + const uint32_t tid = get_thread_idx(); + if (tid == 0) { + readBinaryFile(CHK_DIR "/dy.bin", (uint8_t *)dy, + static_cast(N) * C * sizeof(dtype)); + readBinaryFile(CHK_DIR "/x.bin", (uint8_t *)x, + static_cast(N) * C * sizeof(dtype)); + readBinaryFile(CHK_DIR "/mean.bin", (uint8_t *)mean, + static_cast(N) * G * sizeof(float)); + readBinaryFile(CHK_DIR "/rstd.bin", (uint8_t *)rstd, + static_cast(N) * G * sizeof(float)); + readBinaryFile(CHK_DIR "/gamma.bin", (uint8_t *)gamma, + static_cast(C) * sizeof(dtype)); + input_ready = 1; + } else { + while (input_ready == 0) { + } + } +#endif + + group_norm_grad_1d(dy, x, mean, rstd, gamma, tiling_info, + dx, dgamma, dbeta); + +#ifdef RES_CHECK + kernel_done[tid] = 1; + if (tid == 0) { + for (int pe = 0; pe < PE_NUM; ++pe) { + while (kernel_done[pe] == 0) { + } + } + writeBinaryFile(CHK_DIR "/dx.bin", (uint8_t *)dx, + static_cast(N) * C * sizeof(dtype)); + writeBinaryFile(CHK_DIR "/dgamma.bin", (uint8_t *)dgamma, + static_cast(C) * sizeof(dtype)); + writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, + static_cast(C) * sizeof(dtype)); + output_written = 1; + } else { + while (output_written == 0) { + } + } +#endif +} diff --git a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py similarity index 91% rename from benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py rename to benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py index 9e3bf08b..75511586 100644 --- a/benchmark/one-level-arch/test/kernel/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py +++ b/benchmark/one-level-arch/test/solution/normalization/group_norm_grad_1d/src/group_norm_grad_1d_data_compare.py @@ -15,9 +15,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_group_norm_grad_1d_group_norm_grad_1d_DType__half_N8_C64_G8" + / "solution_normalization_group_norm_grad_1d_group_norm_grad_1d_DType__half_N512_C64_G8_PE4" ) @@ -79,11 +79,13 @@ def compare_one( rel_diffs.append(d / denom) mse = se / n if n else 0.0 + ref_mean_square = sum(b * b for b in ref) / n if n else 0.0 + mse_limit = mse_tol * max(ref_mean_square, 1.0) max_abs = max(abs_diffs) if abs_diffs else 0.0 mean_abs = sum(abs_diffs) / n if n else 0.0 max_rel = max(rel_diffs) if rel_diffs else 0.0 within = all((d <= atol) or (r <= rtol) for d, r in zip(abs_diffs, rel_diffs)) - mse_ok = mse <= mse_tol + mse_ok = mse <= mse_limit status = "pass" if within and mse_ok else "fail" return { "name": name, @@ -92,6 +94,8 @@ def compare_one( "max_abs": max_abs, "mean_abs": mean_abs, "mse": mse, + "ref_mean_square": ref_mean_square, + "mse_limit": mse_limit, "max_rel": max_rel, "atol": atol, "rtol": rtol, diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm/Makefile b/benchmark/one-level-arch/test/solution/normalization/rms_norm/Makefile new file mode 100644 index 00000000..2b164b7c --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/Makefile @@ -0,0 +1,14 @@ +DType ?= __half +G_A ?= 512 +G_R ?= 8192 +PE_NUM ?= 4 +ifneq ($(PE_NUM),4) +$(error rms_norm supports only the dynamic 4PE testcase) +endif +START_FILE = $(ROOT)/test/common/_start_multipe.s +ifeq ($(TESTCASE), rms_norm) +DEFINES += -DDType=$(DType) -DPE_NUM=$(PE_NUM) -DG_A=$(G_A) -DG_R=$(G_R) +TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_PE$(PE_NUM).elf +endif +SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp +include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm/compile.all b/benchmark/one-level-arch/test/solution/normalization/rms_norm/compile.all new file mode 100755 index 00000000..46c391b2 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/compile.all @@ -0,0 +1,5 @@ +#!/bin/bash +# Dynamic 4PE RMSNorm [512,8192]; tiling is calculated from runtime shape. +set -euo pipefail +: "${COMPILER_DIR:?Set COMPILER_DIR to the compiler bin directory}" +make TESTCASE=rms_norm COMPILER_DIR="$COMPILER_DIR" DType=__half G_A=512 G_R=8192 PE_NUM=4 diss diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/data/build_data_obj.sh b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/data/build_data_obj.sh similarity index 100% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/data/build_data_obj.sh rename to benchmark/one-level-arch/test/solution/normalization/rms_norm/src/data/build_data_obj.sh diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/gen_rms_norm_data.py b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/gen_rms_norm_data.py similarity index 93% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/gen_rms_norm_data.py rename to benchmark/one-level-arch/test/solution/normalization/rms_norm/src/gen_rms_norm_data.py index 24d63b2c..b20d8bab 100755 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/gen_rms_norm_data.py +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/gen_rms_norm_data.py @@ -6,7 +6,7 @@ golden.bin : same shape float16, out = x * rsqrt(mean(x^2)+eps) (fp32 compute then cast to fp16, matching kernel pipeline) -Default: g_a=16, g_r=512, tile=(1,512), eps=1e-6. +Default: g_a=512, g_r=8192, tile=(1,8192), eps=1e-6. """ from __future__ import annotations @@ -19,9 +19,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_rms_norm_rms_norm_DType__half" + / "solution_normalization_rms_norm_rms_norm_DType__half_gA512_gR8192_PE4" ) DATA_DIR = SCRIPT_DIR / "data" @@ -140,10 +140,10 @@ def gen_all( def main() -> None: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument("--g-a", type=int, default=16) - parser.add_argument("--g-r", type=int, default=512) + parser.add_argument("--g-a", type=int, default=512) + parser.add_argument("--g-r", type=int, default=8192) parser.add_argument("--tile-a", type=int, default=1) - parser.add_argument("--tile-r", type=int, default=512) + parser.add_argument("--tile-r", type=int, default=8192) parser.add_argument("--eps", type=float, default=1e-6) parser.add_argument("--seed", type=int, default=123) parser.add_argument("-o", "--out-dir", type=Path, default=DEFAULT_CMP_DIR) diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm.cpp b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm.cpp new file mode 100644 index 00000000..e65c6127 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm.cpp @@ -0,0 +1,95 @@ +#include + +#include + +#include "fileop.h" +#include "solution/normalization/rms_norm/rms_norm_pto.hpp" + +#ifndef DType +#define DType __half +#endif + +#ifndef EPS +#define EPS 1e-6f +#endif + +#ifndef PE_NUM +#define PE_NUM 1 +#endif +#ifndef G_A +#define G_A 512 +#endif +#ifndef G_R +#define G_R 8192 +#endif + +namespace { +constexpr int64_t rms_tile_a(int64_t global_a, int64_t pe_num) { + return global_a > 0 && pe_num > 0 ? 1 : 0; +} +constexpr int64_t rms_tile_r(int64_t reduce_size) { + constexpr int64_t kMaxTileR = 8192; + return reduce_size < kMaxTileR ? reduce_size : kMaxTileR; +} +} // namespace + +#ifdef RES_CHECK +namespace { +volatile uint32_t input_ready = 0; +volatile uint32_t kernel_done[PE_NUM] = {}; +volatile uint32_t output_written = 0; +} // namespace +#endif + +int main() { + using dtype = DType; + + // tiling_info is always the host-visible full shape. PE partitioning is + // entirely owned by the kernel. + constexpr int64_t kTileA = rms_tile_a(G_A, PE_NUM); + constexpr int64_t kTileR = rms_tile_r(G_R); + static_assert(G_A > 0 && G_R > 0); + static_assert(kTileA > 0 && kTileR == G_R); + int64_t tiling_info[4] = {G_A, G_R, kTileA, kTileR}; + + const int64_t g_a = tiling_info[0]; + const int64_t g_r = tiling_info[1]; + + static dtype input_buf[G_A * G_R]; + static dtype output_buf[G_A * G_R]; + dtype *input = input_buf; + dtype *output = output_buf; + +#ifdef RES_CHECK +#ifndef CHK_DIR +#error "CHK_DIR must be set when RES_CHECK is enabled" +#endif + const uint32_t tid = get_thread_idx(); + if (tid == 0) { + readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, + static_cast(g_a) * g_r * sizeof(dtype)); + input_ready = 1; + } else { + while (input_ready == 0) { + } + } +#endif + + rms_norm(input, tiling_info, output, EPS); + +#ifdef RES_CHECK + kernel_done[tid] = 1; + if (tid == 0) { + for (int pe = 0; pe < PE_NUM; ++pe) { + while (kernel_done[pe] == 0) { + } + } + writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, + static_cast(g_a) * g_r * sizeof(dtype)); + output_written = 1; + } else { + while (output_written == 0) { + } + } +#endif +} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_data_compare.py b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm_data_compare.py similarity index 96% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_data_compare.py rename to benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm_data_compare.py index 4d95daf5..f534ad5b 100755 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/rms_norm_data_compare.py +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/rms_norm_data_compare.py @@ -12,9 +12,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_rms_norm_rms_norm_DType__half" + / "solution_normalization_rms_norm_rms_norm_DType__half_gA512_gR8192_PE4" ) diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/run_precision_check.py b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/run_precision_check.py similarity index 92% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/run_precision_check.py rename to benchmark/one-level-arch/test/solution/normalization/rms_norm/src/run_precision_check.py index 517b1a26..99036ae1 100755 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm/src/run_precision_check.py +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm/src/run_precision_check.py @@ -18,8 +18,8 @@ CASE_DIR = SCRIPT_DIR.parent # .../rms_norm ONE_LEVEL = CASE_DIR.parents[3] # .../one-level-arch SUPERSCALAR = ONE_LEVEL.parents[2] # .../SuperScalar (workspace) -ELF_NAME = "kernel_normalization_rms_norm_rms_norm_DType__half.elf" -ELF_PATH = ONE_LEVEL / "output" / "kernel" / "normalization" / "rms_norm" / "elf" / ELF_NAME +ELF_NAME = "solution_normalization_rms_norm_rms_norm_DType__half_gA512_gR8192_PE4.elf" +ELF_PATH = ONE_LEVEL / "output" / "solution" / "normalization" / "rms_norm" / "elf" / ELF_NAME CMP_DIR = ONE_LEVEL / "compare" / ELF_NAME.replace(".elf", "") DEFAULT_COMPILER = ( SUPERSCALAR / "linx-toolchain-build" / "output" / "linx_blockisa_llvm_musl" / "bin" @@ -90,7 +90,7 @@ def main() -> int: # Ensure compare dir exists for host I/O under RES_CHECK. CMP_DIR.mkdir(parents=True, exist_ok=True) run( - [str(args.gfrun), "-f", str(ELF_PATH), "-t", "1"], + [str(args.gfrun), "-f", str(ELF_PATH), "-s", "softcore.multiThreadNum=4"], cwd=args.gfrun.parent.parent, # SuperScalarModel ) diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/Makefile b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/Makefile new file mode 100644 index 00000000..7d7c394e --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/Makefile @@ -0,0 +1,14 @@ +DType ?= __half +G_A ?= 16 +G_R ?= 16384 +PE_NUM ?= 4 +ifneq ($(PE_NUM),4) +$(error rms_norm_binary supports only the dynamic 4PE testcase) +endif +START_FILE = $(ROOT)/test/common/_start_multipe.s +ifeq ($(TESTCASE), rms_norm_binary) +DEFINES += -DDType=$(DType) -DG_A=$(G_A) -DG_R=$(G_R) -DPE_NUM=$(PE_NUM) +TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_PE$(PE_NUM).elf +endif +SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/$(TESTCASE).cpp +include ../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/compile.all b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/compile.all new file mode 100755 index 00000000..bb7818e3 --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/compile.all @@ -0,0 +1,5 @@ +#!/bin/bash +# Dynamic 4PE RMSNorm binary [16,16384]; tiling is calculated from runtime shape. +set -euo pipefail +: "${COMPILER_DIR:?Set COMPILER_DIR to the compiler bin directory}" +make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half G_A=16 G_R=16384 PE_NUM=4 diss diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py similarity index 91% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py rename to benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py index 2e80dc89..1bfa558d 100755 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/gen_rms_norm_binary_data.py @@ -1,5 +1,5 @@ #!/usr/bin/env python3 -"""Generate rms_norm_binary host bins for [1, 8192] fp16. +"""Generate rms_norm_binary host bins for [16, 16384] fp16. tiling_info.bin : 5 x int64 LE = (g_a, g_r, tile_a, tile_r, pow_r) input.bin : g_a * g_r x float16 @@ -18,9 +18,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half" + / "solution_normalization_rms_norm_binary_rms_norm_binary_DType__half_gA16_gR16384_PE4" ) DATA_DIR = SCRIPT_DIR / "data" @@ -121,11 +121,11 @@ def gen_all( def main() -> None: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument("--g-a", type=int, default=1) - parser.add_argument("--g-r", type=int, default=8192) + parser.add_argument("--g-a", type=int, default=16) + parser.add_argument("--g-r", type=int, default=16384) parser.add_argument("--tile-a", type=int, default=1) - parser.add_argument("--tile-r", type=int, default=1024) - parser.add_argument("--pow-r", type=int, default=4096) + parser.add_argument("--tile-r", type=int, default=8192) + parser.add_argument("--pow-r", type=int, default=8192) parser.add_argument("--eps", type=float, default=1e-6) parser.add_argument("--seed", type=int, default=123) parser.add_argument("-o", "--out-dir", type=Path, default=DEFAULT_CMP_DIR) diff --git a/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary.cpp b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary.cpp new file mode 100644 index 00000000..c32d380d --- /dev/null +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary.cpp @@ -0,0 +1,109 @@ +#include + +#include + +#include "fileop.h" +#include "solution/normalization/rms_norm_binary/rms_norm_binary_pto.hpp" + +#ifndef DType +#define DType __half +#endif + +#ifndef EPS +#define EPS 1e-6f +#endif + +#ifndef PE_NUM +#define PE_NUM 1 +#endif + +// Dynamic 4PE validation shape: [16, 16384], fp16. +#ifndef G_A +#define G_A 16 +#endif +#ifndef G_R +#define G_R 16384 +#endif +// Must match rms_bin::kWsCols / kMaxLevels +#ifndef K_WS_COLS +#define K_WS_COLS 1 +#endif +#ifndef K_MAX_LEVELS +#define K_MAX_LEVELS 6 +#endif + +namespace { +constexpr int64_t floor_power_of_two(int64_t value) { + int64_t result = 1; + while (result <= value / 2) { + result *= 2; + } + return result; +} +constexpr int64_t binary_tile_r(int64_t reduce_size) { + constexpr int64_t kMaxTileR = 8192; + return reduce_size < kMaxTileR ? reduce_size : kMaxTileR; +} +} // namespace + +#ifdef RES_CHECK +namespace { +volatile uint32_t input_ready = 0; +volatile uint32_t kernel_done[PE_NUM] = {}; +volatile uint32_t output_written = 0; +} // namespace +#endif + +int main() { + using dtype = DType; + + constexpr int64_t kTileA = 1; + constexpr int64_t kTileR = binary_tile_r(G_R); + constexpr int64_t kPowR = floor_power_of_two(G_R - 1); + static_assert(G_A > 0 && G_R > 1); + static_assert(kPowR < G_R && G_R <= 2 * kPowR); + int64_t tiling_info[5] = {G_A, G_R, kTileA, kTileR, kPowR}; + + const int64_t g_a = tiling_info[0]; + const int64_t g_r = tiling_info[1]; + + static dtype input_buf[G_A * G_R]; + static dtype output_buf[G_A * G_R]; + static float workspace_buf[K_MAX_LEVELS * G_A * K_WS_COLS]; + dtype *input = input_buf; + dtype *output = output_buf; + float *workspace = workspace_buf; + +#ifdef RES_CHECK +#ifndef CHK_DIR +#error "CHK_DIR must be set when RES_CHECK is enabled" +#endif + const uint32_t tid = get_thread_idx(); + if (tid == 0) { + readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, + static_cast(g_a) * g_r * sizeof(dtype)); + input_ready = 1; + } else { + while (input_ready == 0) { + } + } +#endif + + rms_norm_binary(input, tiling_info, output, workspace, EPS); + +#ifdef RES_CHECK + kernel_done[tid] = 1; + if (tid == 0) { + for (int pe = 0; pe < PE_NUM; ++pe) { + while (kernel_done[pe] == 0) { + } + } + writeBinaryFile(CHK_DIR "/output.bin", (uint8_t *)output, + static_cast(g_a) * g_r * sizeof(dtype)); + output_written = 1; + } else { + while (output_written == 0) { + } + } +#endif +} diff --git a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py similarity index 96% rename from benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py rename to benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py index 75d6ba46..6b3aac00 100755 --- a/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py +++ b/benchmark/one-level-arch/test/solution/normalization/rms_norm_binary/src/rms_norm_binary_data_compare.py @@ -12,9 +12,9 @@ SCRIPT_DIR = Path(__file__).resolve().parent DEFAULT_CMP_DIR = ( - SCRIPT_DIR.parents[4] + SCRIPT_DIR.parents[5] / "compare" - / "kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half" + / "solution_normalization_rms_norm_binary_rms_norm_binary_DType__half_gA16_gR16384_PE4" ) From 4230e5f0a97613ee35825bce8909ed040899c492 Mon Sep 17 00:00:00 2001 From: wangyuascend-spec Date: Mon, 7 Sep 2026 14:38:42 +0800 Subject: [PATCH 2/2] refactor(normalization): keep implementations only in solution --- benchmark/one-level-arch/compile_all.sh | 2 - .../kernels/multi_thread/README.md | 13 +- .../normalization/rms_norm/rms_norm.hpp | 71 -- .../rms_norm_binary/rms_norm_binary.hpp | 29 - .../group_norm_grad/group_norm_grad.hpp | 614 ------------------ .../group_norm_grad_1d/group_norm_grad_1d.hpp | 532 --------------- .../normalization/rms_norm/rms_norm.hpp | 167 ----- .../normalization/rms_norm_binary/SKILL.md | 189 ------ .../rms_norm_binary/rms_norm_binary.hpp | 377 ----------- .../normalization/rms_norm/Makefile | 10 - .../normalization/rms_norm/compile.all | 4 - .../normalization/rms_norm/src/rms_norm.cpp | 38 -- .../normalization/rms_norm_binary/Makefile | 10 - .../normalization/rms_norm_binary/compile.all | 4 - .../rms_norm_binary/src/rms_norm_binary.cpp | 37 -- .../test/kernel/multi_thread/res_check_all.py | 9 - 16 files changed, 1 insertion(+), 2105 deletions(-) delete mode 100644 benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm/rms_norm.hpp delete mode 100644 benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm_binary/rms_norm_binary.hpp delete mode 100644 benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad/group_norm_grad.hpp delete mode 100644 benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp delete mode 100644 benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm/rms_norm.hpp delete mode 100644 benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md delete mode 100644 benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/Makefile delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/src/rms_norm.cpp delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/Makefile delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/compile.all delete mode 100644 benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/src/rms_norm_binary.cpp diff --git a/benchmark/one-level-arch/compile_all.sh b/benchmark/one-level-arch/compile_all.sh index b0988b95..9a2316fb 100755 --- a/benchmark/one-level-arch/compile_all.sh +++ b/benchmark/one-level-arch/compile_all.sh @@ -70,8 +70,6 @@ compile_operator "$REPO_ROOT/test/kernel/multi_thread/concat" "multi_thread/conc compile_operator "$REPO_ROOT/test/kernel/multi_thread/conv2d" "multi_thread/conv2d" compile_operator "$REPO_ROOT/test/kernel/multi_thread/element_wise/gelu" "multi_thread/element_wise/gelu" compile_operator "$REPO_ROOT/test/kernel/multi_thread/gather" "multi_thread/gather" -compile_operator "$REPO_ROOT/test/kernel/multi_thread/normalization/rms_norm" "multi_thread/normalization/rms_norm" -compile_operator "$REPO_ROOT/test/kernel/multi_thread/normalization/rms_norm_binary" "multi_thread/normalization/rms_norm_binary" compile_operator "$REPO_ROOT/test/kernel/multi_thread/reduction/cumsum_row" "multi_thread/reduction/cumsum_row" compile_operator "$REPO_ROOT/test/kernel/multi_thread/reduction/reducemax_row" "multi_thread/reduction/reducemax_row" compile_operator "$REPO_ROOT/test/kernel/multi_thread/reduction/reduceprod_row" "multi_thread/reduction/reduceprod_row" diff --git a/benchmark/one-level-arch/kernels/multi_thread/README.md b/benchmark/one-level-arch/kernels/multi_thread/README.md index 0078f2fc..2b9edd93 100644 --- a/benchmark/one-level-arch/kernels/multi_thread/README.md +++ b/benchmark/one-level-arch/kernels/multi_thread/README.md @@ -28,8 +28,6 @@ divisible by the PE count. This is checked at compile time. | Shared Matmul | `matmul/matmul_shared.hpp` | | Shared-B-reuse Matmul | `matmul/matmul_shared_reuseB.hpp` | | Low-precision Matmul | `matmul/matmul_shared_lowp.hpp` | -| RMSNorm | `normalization/rms_norm/rms_norm.hpp` | -| Binary-accumulation RMSNorm | `normalization/rms_norm_binary/rms_norm_binary.hpp` | | Row Cumsum | `reduction/cumsum_rowvec.hpp` | | Row Max/Prod/Sum | `reduction/reducemax_rowvec.hpp`, `reduction/reduceprod_rowvec.hpp`, `reduction/reducesum_rowvec.hpp` | | 2D Transpose | `transpose/transpose.hpp` | @@ -43,7 +41,6 @@ Kernel and test paths mirror the single-PE tree. For example: | `kernels/single_thread/gather/gather.hpp` | `kernels/multi_thread/gather/gather.hpp` | | `test/kernel/gather/` | `test/kernel/multi_thread/gather/` | | `test/kernel/element_wise/gelu/` | `test/kernel/multi_thread/element_wise/gelu/` | -| `test/solution/normalization/rms_norm_binary/` | `test/kernel/multi_thread/normalization/rms_norm_binary/` | Each operator directory has its own `Makefile`, `compile.all`, and `src/` instead of sharing a mixed test source. One model failure therefore does not @@ -187,7 +184,6 @@ Run one or more named cases: ```bash python3 benchmark/one-level-arch/test/kernel/multi_thread/res_check_all.py \ - broadcast fa matmul_shared rms_norm ``` Run the complete representative numerical portfolio: @@ -242,20 +238,13 @@ coverage for the other precision variants. | `matmul_shared` | PASS | 0 | | `matmul_reuseB` | PASS | 0 | | `matmul_lowp` FP8 | PASS | 0 | -| `rms_norm` | PASS | 0.000976562 | -| `rms_norm_binary` | PASS | 0.000976562 | | `cumsum_row` | PASS | 0 | | `reducemax_row` | FAIL | 0.999606; row-result physical stride mismatch | | `reduceprod_row` | FAIL | 1.06037; row-result physical stride mismatch | | `reducesum_row` | PASS | 5.72205e-06 | | `transpose` | PASS | 0 | | `tadd` | PASS | 0 | -| **Total** | **14 PASS / 4 FAIL / 0 TIMEOUT** | **18 representative cases** | - -Two RMSNorm cases initially failed because the old Newton iteration used -`TRECIP(x)` as the inverse-square-root seed. With the current compiler's -`TRSQRT` support, replacing that sequence with `TRSQRT` reduced the maximum -absolute error to 0.000976562 and made both cases pass. +| **Total** | **12 PASS / 4 FAIL / 0 TIMEOUT** | **16 representative cases** | The remaining failures are kernel/API-model issues exposed by numerical checking, rather than binary I/O or four-PE synchronization failures: diff --git a/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm/rms_norm.hpp b/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm/rms_norm.hpp deleted file mode 100644 index 4f162590..00000000 --- a/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm/rms_norm.hpp +++ /dev/null @@ -1,71 +0,0 @@ -#pragma once - -#include "single_thread/normalization/rms_norm/rms_norm.hpp" - -// Compile-time four-PE RMSNorm. The A dimension is divided equally across -// PEs; each PE processes one contiguous row range and writes its own output. -template -void rms_norm(dtype *x, dtype *out, float eps = 1e-6f) { - static_assert(gA > 0 && gR > 0 && tA > 0 && tR > 0); - static_assert(peA > 0 && gA % peA == 0, "gA must be divisible by peA"); - static_assert(peA >= tA, "peA must cover one tile_a"); - static_assert(tA <= 128, "RMSNorm row-state carrier supports tA <= 128"); - static_assert(tR == gR, "static rms_norm is a single R-tile; use rms_norm_binary for R-split"); - constexpr int Mb = peA / tA; - constexpr int rmd_A = peA % tA; - constexpr float inv_r = 1.0f / static_cast(gR); - - using gm_t = global_tensor>; - using tile_h = Tile; - using tile_f = Tile; - using tile_v = - Tile; - using it_t = global_iterator; - - const uint32_t tid = get_thread_idx(); - const uint32_t gm_offset = tid * static_cast(peA * gR); - - it_t gI(x + gm_offset); - it_t gO(out + gm_offset); - - for (int ia = 0; ia < Mb; ++ia) { - tile_h src_h, dst_h; - tile_f src, squared, dst; - tile_v sqrsum, mean, denom, rms; - auto gi = gI(ia, 0); - auto go = gO(ia, 0); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(squared, src, src); - TROWSUM(sqrsum, squared); - TMULS(mean, sqrsum, inv_r); - TADDS(denom, mean, eps); - rms_detail::rsqrt_newton(rms, denom); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - if constexpr (rmd_A) { - using tile_h_r = Tile; - using tile_f_r = Tile; - using tile_v_r = Tile; - tile_h_r src_h, dst_h; - tile_f_r src, squared, dst; - tile_v_r sqrsum, mean, denom, rms; - auto gi = gI(Mb, 0); - auto go = gO(Mb, 0); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(squared, src, src); - TROWSUM(sqrsum, squared); - TMULS(mean, sqrsum, inv_r); - TADDS(denom, mean, eps); - rms_detail::rsqrt_newton(rms, denom); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } -} diff --git a/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm_binary/rms_norm_binary.hpp b/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm_binary/rms_norm_binary.hpp deleted file mode 100644 index 3260526e..00000000 --- a/benchmark/one-level-arch/kernels/multi_thread/normalization/rms_norm_binary/rms_norm_binary.hpp +++ /dev/null @@ -1,29 +0,0 @@ -#pragma once - -#include "multi_thread/utils/spmd_partition.hpp" -#include "single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp" - -#include - -namespace supernpu::multi_thread { - -// Rows are independent. Workspace is also divided by row so no cache level is -// shared between PEs. -template -void rms_norm_binary(DType *input, DType *output, float *workspace, - float eps = 1e-6f) { - using Partition = ContiguousPartition; - constexpr int kRowsPerPe = static_cast(Partition::kItemsPerPe); - constexpr std::size_t kWorkspaceValuesPerRow = - rms_bin::kMaxLevels * rms_bin::kWsCols; - const std::size_t row_offset = Partition::item_offset(); - - ::rms_norm_binary( - input + row_offset * Columns, output + row_offset * Columns, - workspace + row_offset * kWorkspaceValuesPerRow, eps); -} - -} // namespace supernpu::multi_thread diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad/group_norm_grad.hpp b/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad/group_norm_grad.hpp deleted file mode 100644 index eb4283b6..00000000 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad/group_norm_grad.hpp +++ /dev/null @@ -1,614 +0,0 @@ -// ============================================================================= -// group_norm_grad.hpp — GroupNorm backward, HxW > 1 (one-level PTO) -// ============================================================================= -// -// Matches PyTorch GroupNormBackwardKernelImplInternal -// (aten/src/ATen/native/cuda/group_norm_kernel.cu): -// 1) Spatial reduce ds/db = Σ_hw (dY*X), Σ_hw dY per (n,c) -// 2) Fused c2/c3 from ds/db over channels → workspace -// 3) dX = (rstd*gamma)*dY + c2*X + c3 -// 4) dgamma / dbeta from ds/db over N -// -// Layout: X/dY/dX [N,C,HxW]; mean/rstd [N,G] fp32; gamma/dgamma/dbeta [C] -// workspace float[2*N*C+2*N*G] -// All pointers required (dy,x,mean,rstd,gamma,dx,dgamma,dbeta,workspace). -// tiling[5] = {N, C, G, HxW, tile_hw} -// tile_hw <= 0 → min(HxW, tCap); spatial R-split when HxW > tile_hw. -// dX still requires HxW <= tCap (one spatial tile; RF-limited). -// -// tCap: logical tile >= 512B (TileOP IsValidActiveSize / TSize=1..7). -// fp16 → Cols>=256; fp32 → Cols>=128. tile_v is always fp32 → Cols=128. -// Cols=1024 like rms_norm overflows Tile RF here. -// -// Torch CUDA launch 总览 (NVIDIA, warp=32): -// Step1 ComputeInternalGradientsCUDAKernel -// grid=N*C, block=(HxW<512)?32:512 -// Step2 ComputeBackwardFusedParamsCUDAKernel -// grid=dim3(N,G), block=(D<512)?32:512 -// Step3 dX gpu_kernel (+ optional c1) -// block=128, vt=4(fp16)/2(fp32), grid=ceil(numel/(128*vt)) -// Step4 GammaBetaBackwardCUDAKernel1/2 -// N<=128: grid=ceil(C/256), block=256 -// N>128: grid=ceil(C/32), block=dim3(32,16) -// ============================================================================= -#ifndef SUPERNPU_GROUP_NORM_GRAD_PTO_HPP -#define SUPERNPU_GROUP_NORM_GRAD_PTO_HPP - -#include - -#include - -namespace gn_grad { - -inline int64_t workspace_elems(int64_t N, int64_t C, int64_t G) { - return 2 * N * C + 2 * N * G; -} - -// --------------------------------------------------------------------------- -// Step 1: spatial reduce for one (n, c) → ds[nc], db[nc] (HxW R-split) -// -// Torch: ComputeInternalGradientsCUDAKernel -// grid = N * C // 每个 (n,c) 一个 block；本函数 = 其中一个 block -// block = (HxW < 512) ? 32 : 512 -// 线程: threadIdx.x 沿 hw 做 grid-stride + warp/block reduce -// → ds[n,c]=Σ dY*X, db[n,c]=Σ dY -// --------------------------------------------------------------------------- -template -inline void spatial_reduce_nc(dtype *dy, dtype *x, float *ds, float *db, - int64_t N, int64_t C, int64_t HxW, - int64_t tile_hw, int64_t n, int64_t c) { - const int64_t nc = n * C + c; - const int64_t base = nc * HxW; - - gm_f gds(ds + nc, static_cast(N * C), 1); - gm_f gdb(db + nc, static_cast(N * C), 1); - - tile_v ds_acc(1); - tile_v db_acc(1); - tile_v cur(1); - TEXPANDS(ds_acc, 0.0f); - TEXPANDS(db_acc, 0.0f); - - // Torch: for (hw = threadIdx.x; hw < HxW; hw += blockDim.x) + BlockReduce - // PTO: tile 覆盖一段 HxW，TROWSUM 代替 block 内线程归约 - for (int64_t hw0 = 0; hw0 < HxW; hw0 += tile_hw) { - const size_t vh = static_cast( - (hw0 + tile_hw <= HxW) ? tile_hw : (HxW - hw0)); - const int64_t offset = base + hw0; - - gm_h gdy(dy + offset, static_cast(N * C), static_cast(HxW)); - gm_h gx(x + offset, static_cast(N * C), static_cast(HxW)); - - tile_h h0(1, vh); - tile_f x_f(1, vh); - tile_f dy_f(1, vh); - tile_f prod(1, vh); - - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TMUL(prod, dy_f, x_f); - TROWSUM(cur, prod); - TADD(ds_acc, ds_acc, cur); - TROWSUM(cur, dy_f); - TADD(db_acc, db_acc, cur); - } - - TSTORE(gds, ds_acc); - TSTORE(gdb, db_acc); -} - -// --------------------------------------------------------------------------- -// Step 2: fused c2/c3 for one (n, g) → c2[ng], c3[ng] -// -// Torch: ComputeBackwardFusedParamsCUDAKernel -// grid = dim3(N, G) // blockIdx.x=n, blockIdx.y=g；本函数 = 其中一个 -// block = (D < 512) ? 32 : 512 -// 线程: threadIdx.x 沿 group 内通道 i∈[0,D) stride，再 block reduce -// → c2,c3 每 (n,g) 各一个标量 -// --------------------------------------------------------------------------- -template -inline void fused_params_group(dtype *gamma, float *mean, float *rstd, - float *ds, float *db, float *c2_buf, - float *c3_buf, int64_t N, int64_t C, int64_t G, - int64_t D, int64_t n, int64_t g, float s) { - const int64_t ng = n * G + g; - const int64_t c0 = g * D; - const size_t active_d = static_cast(D); - - gm_f gmean(mean + ng, static_cast(N * G), 1); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - gm_f gds(ds + n * C + c0, 1, static_cast(C)); - gm_f gdb(db + n * C + c0, 1, static_cast(C)); - gm_f gc2(c2_buf + ng, static_cast(N * G), 1); - gm_f gc3(c3_buf + ng, static_cast(N * G), 1); - - tile_f ds_f(1, active_d); - tile_f db_f(1, active_d); - tile_f gamma_f(1, active_d); - tile_f t0(1, active_d); - tile_h h0(1, active_d); - tile_v mean_t(1); - tile_v rstd_t(1); - tile_v sum1(1); - tile_v sum2(1); - tile_v c2(1); - tile_v c3(1); - - TLOAD(ds_f, gds); - TLOAD(db_f, gdb); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - - { - gm_h gg(gamma + c0, 1, static_cast(C)); - TLOAD(h0, gg); - TCVT(gamma_f, h0); - } - - // Torch: threads 各算 ds*gamma / db*gamma 再 reduce → sum1/sum2 - TMUL(t0, ds_f, gamma_f); - TROWSUM(sum1, t0); - TMUL(t0, db_f, gamma_f); - TROWSUM(sum2, t0); - - // c2/c3 由 block 内 thread 0（归约后）写出；此处标量 tile 完成同样公式 - TMUL(c2, sum2, mean_t); - TSUB(c2, c2, sum1); - TMUL(c3, rstd_t, rstd_t); - TMUL(c3, c3, rstd_t); - TMUL(c2, c2, c3); - TMULS(c2, c2, s); - - TMUL(c3, c2, mean_t); - TMULS(c3, c3, -1.0f); - TMUL(sum1, sum2, rstd_t); - TMULS(sum1, sum1, s); - TSUB(c3, c3, sum1); - - TSTORE(gc2, c2); - TSTORE(gc3, c3); -} - -// --------------------------------------------------------------------------- -// Step 3: dX for one (n, c) using stored c2/c3 and rstd*gamma -// -// Torch: gpu_kernel 元素级 (可选先算 c1) -// block = 128 -// vt = 4 (fp16/bf16) / 2 (fp32+) -// grid = ceil(numel / (128 * vt)) // numel = N*C*HxW -// 线程: 线性下标覆盖全部元素；c2/c3 按 (n,g) 广播 -// 本函数一次处理一个 (n,c) 的整段 HxW（Tile 覆盖空间维） -// --------------------------------------------------------------------------- -template -inline void dx_nc(dtype *dy, dtype *x, dtype *gamma, float *rstd, float *c2_buf, - float *c3_buf, dtype *dx, int64_t N, int64_t C, int64_t G, - int64_t D, int64_t HxW, int64_t n, int64_t c) { - const int64_t g = c / D; - const int64_t ng = n * G + g; - const int64_t offset = (n * C + c) * HxW; - const size_t active_hw = static_cast(HxW); - - gm_h gdy(dy + offset, static_cast(N * C), static_cast(HxW)); - gm_h gx(x + offset, static_cast(N * C), static_cast(HxW)); - gm_h gdx(dx + offset, static_cast(N * C), static_cast(HxW)); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - gm_f gc2(c2_buf + ng, static_cast(N * G), 1); - gm_f gc3(c3_buf + ng, static_cast(N * G), 1); - - tile_h h0(1, active_hw); - tile_f x_f(1, active_hw); - tile_f dy_f(1, active_hw); - tile_f dx_f(1, active_hw); - tile_f tmp(1, active_hw); - tile_v rstd_t(1); - tile_v c1(1); - tile_v c2(1); - tile_v c3(1); - - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(rstd_t, grstd); - TLOAD(c2, gc2); - TLOAD(c3, gc3); - - // Torch 可选 c1 预计算同为 gpu_kernel block=128；此处 c1 = rstd*gamma[c] - // TCVT must keep matching logical shapes (PTO 0.58). tile_h Cols=tCap - // (fp16→256) while tile_v Cols=128, so convert via tile_f then TROWSUM. - { - gm_h gg(gamma + c, 1, 1); - tile_h hg(1, 1); - tile_f gf(1, 1); - tile_v gv(1); - TLOAD(hg, gg); - TCVT(gf, hg); - TROWSUM(gv, gf); - TMUL(c1, gv, rstd_t); - } - - TROWEXPANDMUL(dx_f, dy_f, c1); - TROWEXPANDMUL(tmp, x_f, c2); - TADD(dx_f, dx_f, tmp); - TROWEXPANDADD(dx_f, dx_f, c3); - - TCVT(h0, dx_f); - TSTORE(gdx, h0); -} - -// --------------------------------------------------------------------------- -// Step 4a: dbeta — dbeta[c] = Σ_n db[n,c] -// -// Torch: GammaBetaBackwardCUDAKernel1/2（与 dgamma 同一次 launch） -// N<=128: grid=ceil(C/256), block=256；每线程一个 c，循环 n -// N>128: grid=ceil(C/32), block=dim3(32,16) -// 本函数按 group 一次写 D 个通道（对 N 串行累加） -// --------------------------------------------------------------------------- -template -inline void dbeta_group(float *db, dtype *dbeta, int64_t N, int64_t C, - int64_t D, int64_t g) { - const int64_t c0 = g * D; - const size_t active_d = static_cast(D); - - tile_f acc(1, active_d); - tile_f cur(1, active_d); - tile_h h0(1, active_d); - TEXPANDS(acc, 0.0f); - - for (int64_t n = 0; n < N; ++n) { - gm_f gdb(db + n * C + c0, 1, static_cast(C)); - TLOAD(cur, gdb); - TADD(acc, acc, cur); - } - - gm_h gout(dbeta + c0, 1, static_cast(C)); - TCVT(h0, acc); - TSTORE(gout, h0); -} - -// --------------------------------------------------------------------------- -// Step 4b: dgamma — dgamma[c] = Σ_n (ds - db*mean)*rstd -// -// Torch: 与 dbeta 同 Kernel1/2 launch（见上） -// --------------------------------------------------------------------------- -template -inline void dgamma_group(float *ds, float *db, float *mean, float *rstd, - dtype *dgamma, int64_t N, int64_t C, int64_t G, - int64_t D, int64_t g) { - const int64_t c0 = g * D; - const size_t active_d = static_cast(D); - - tile_f acc(1, active_d); - tile_f ds_f(1, active_d); - tile_f db_f(1, active_d); - tile_f t0(1, active_d); - tile_h h0(1, active_d); - tile_v mean_t(1); - tile_v rstd_t(1); - TEXPANDS(acc, 0.0f); - - for (int64_t n = 0; n < N; ++n) { - const int64_t ng = n * G + g; - gm_f gds(ds + n * C + c0, 1, static_cast(C)); - gm_f gdb(db + n * C + c0, 1, static_cast(C)); - gm_f gmean(mean + ng, static_cast(N * G), 1); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - - TLOAD(ds_f, gds); - TLOAD(db_f, gdb); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - - TROWEXPANDMUL(t0, db_f, mean_t); - TSUB(t0, ds_f, t0); - TROWEXPANDMUL(t0, t0, rstd_t); - TADD(acc, acc, t0); - } - - gm_h gout(dgamma + c0, 1, static_cast(C)); - TCVT(h0, acc); - TSTORE(gout, h0); -} - -} // namespace gn_grad - -// tiling: [N, C, G, HxW, tile_hw] -// workspace: float[2*N*C + 2*N*G] -// -// 入口循环 ↔ Torch 各 kernel 的 grid 遍历： -// for n,c spatial_reduce ↔ grid = N*C -// for n,g fused_params ↔ grid = dim3(N,G) -// for n,c dx_nc ↔ numel 上 gpu_kernel 线性网格 -// for g dbeta/dgamma ↔ 按通道写回（Kernel1/2） -template -void group_norm_grad(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *gamma, const int64_t *tiling, dtype *dx, - dtype *dgamma, dtype *dbeta, float *workspace) { - // Capacity in elements: every Tile buffer >= 512B (dtype strip + float strip). - constexpr int64_t tCapDtype = - (512 + static_cast(sizeof(dtype)) - 1) / - static_cast(sizeof(dtype)); - constexpr int64_t tCap = tCapDtype > 128 ? tCapDtype : 128; - constexpr int64_t tV = 128; // float scalar/broadcast strip: 128*4B = 512B - - const int64_t N = tiling[0]; - const int64_t C = tiling[1]; - const int64_t G = tiling[2]; - const int64_t HxW = tiling[3]; - const int64_t tile_hw = - tiling[4] > 0 ? tiling[4] : (HxW < tCap ? HxW : tCap); - const int64_t D = C / G; - - float *ds = workspace; - float *db = workspace + N * C; - float *c2_buf = workspace + 2 * N * C; - float *c3_buf = c2_buf + N * G; - - using gm_h = global_tensor>; - using gm_f = global_tensor>; - using tile_h = - Tile; - using tile_f = - Tile; - using tile_v = - Tile; - - const float s = 1.0f / static_cast(D * HxW); - - for (int64_t n = 0; n < N; ++n) { - for (int64_t c = 0; c < C; ++c) { - gn_grad::spatial_reduce_nc(dy, x, ds, db, N, C, HxW, - tile_hw, n, c); - } - } - - for (int64_t n = 0; n < N; ++n) { - for (int64_t g = 0; g < G; ++g) { - gn_grad::fused_params_group(gamma, mean, rstd, ds, db, - c2_buf, c3_buf, N, C, G, D, n, - g, s); - } - } - for (int64_t n = 0; n < N; ++n) { - for (int64_t c = 0; c < C; ++c) { - gn_grad::dx_nc( - dy, x, gamma, rstd, c2_buf, c3_buf, dx, N, C, G, D, HxW, n, c); - } - } - - for (int64_t g = 0; g < G; ++g) { - gn_grad::dbeta_group(db, dbeta, N, - C, D, g); - } - for (int64_t g = 0; g < G; ++g) { - gn_grad::dgamma_group( - ds, db, mean, rstd, dgamma, N, C, G, D, g); - } -} - -// Compile-time N,C,G,HxW,tile_hw. Physical Cols = tCap (>=512B); Valid = tiling. -template -void group_norm_grad(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *gamma, dtype *dx, dtype *dgamma, dtype *dbeta, - float *workspace) { - static_assert(N > 0 && C > 0 && G > 0 && HxW > 0 && tile_hw > 0); - static_assert(C % G == 0); - constexpr int D = C / G; - constexpr int tCapDtype = - (512 + static_cast(sizeof(dtype)) - 1) / - static_cast(sizeof(dtype)); - constexpr int tCap = tCapDtype > 128 ? tCapDtype : 128; - constexpr int tV = 128; - static_assert(tile_hw <= tCap && HxW <= tCap && D <= tCap); - constexpr int n_hw = HxW / tile_hw; - constexpr int rmd_hw = HxW % tile_hw; - constexpr float s = 1.0f / static_cast(D * HxW); - - float *ds = workspace; - float *db = workspace + N * C; - float *c2_buf = workspace + 2 * N * C; - float *c3_buf = c2_buf + N * G; - - using gm_h = global_tensor>; - using gm_c = global_tensor>; - using gm_f1 = global_tensor>; - using gm_fC = global_tensor>; - using gm_fG = global_tensor>; - using tile_h_hw = - Tile; - using tile_f_hw = - Tile; - using tile_h_x = - Tile; - using tile_f_x = - Tile; - using tile_h_d = - Tile; - using tile_f_d = - Tile; - using tile_h_1 = - Tile; - using tile_f_1 = - Tile; - using tile_v = - Tile; - - for (int n = 0; n < N; ++n) { - for (int c = 0; c < C; ++c) { - const int nc = n * C + c; - const int64_t base = static_cast(nc) * HxW; - gm_f1 gds(ds + nc); - gm_f1 gdb(db + nc); - tile_v ds_acc, db_acc, cur; - TEXPANDS(ds_acc, 0.0f); - TEXPANDS(db_acc, 0.0f); - for (int k = 0; k < n_hw; ++k) { - gm_h gx(x + base + k * tile_hw); - gm_h gdy(dy + base + k * tile_hw); - tile_h_hw h0; - tile_f_hw x_f, dy_f, prod; - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TMUL(prod, dy_f, x_f); - TROWSUM(cur, prod); - TADD(ds_acc, ds_acc, cur); - TROWSUM(cur, dy_f); - TADD(db_acc, db_acc, cur); - } - if constexpr (rmd_hw) { - using tile_h_r = Tile; - using tile_f_r = Tile; - gm_h gx(x + base + n_hw * tile_hw); - gm_h gdy(dy + base + n_hw * tile_hw); - tile_h_r h0; - tile_f_r x_f, dy_f, prod; - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TMUL(prod, dy_f, x_f); - TROWSUM(cur, prod); - TADD(ds_acc, ds_acc, cur); - TROWSUM(cur, dy_f); - TADD(db_acc, db_acc, cur); - } - TSTORE(gds, ds_acc); - TSTORE(gdb, db_acc); - } - } - - for (int n = 0; n < N; ++n) { - for (int g = 0; g < G; ++g) { - const int ng = n * G + g; - const int c0 = g * D; - gm_fG gmean(mean + ng); - gm_fG grstd(rstd + ng); - gm_fC gds(ds + n * C + c0); - gm_fC gdb(db + n * C + c0); - gm_fG gc2(c2_buf + ng); - gm_fG gc3(c3_buf + ng); - gm_c gg(gamma + c0); - tile_f_d ds_f, db_f, gamma_f, t0; - tile_h_d h0; - tile_v mean_t, rstd_t, sum1, sum2, c2, c3; - TLOAD(ds_f, gds); - TLOAD(db_f, gdb); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - TLOAD(h0, gg); - TCVT(gamma_f, h0); - TMUL(t0, ds_f, gamma_f); - TROWSUM(sum1, t0); - TMUL(t0, db_f, gamma_f); - TROWSUM(sum2, t0); - TMUL(c2, sum2, mean_t); - TSUB(c2, c2, sum1); - TMUL(c3, rstd_t, rstd_t); - TMUL(c3, c3, rstd_t); - TMUL(c2, c2, c3); - TMULS(c2, c2, s); - TMUL(c3, c2, mean_t); - TMULS(c3, c3, -1.0f); - TMUL(sum1, sum2, rstd_t); - TMULS(sum1, sum1, s); - TSUB(c3, c3, sum1); - TSTORE(gc2, c2); - TSTORE(gc3, c3); - } - } - - for (int n = 0; n < N; ++n) { - for (int c = 0; c < C; ++c) { - const int g = c / D; - const int ng = n * G + g; - const int64_t offset = (static_cast(n) * C + c) * HxW; - gm_h gdy(dy + offset); - gm_h gx(x + offset); - gm_h gdx(dx + offset); - gm_fG grstd(rstd + ng); - gm_fG gc2(c2_buf + ng); - gm_fG gc3(c3_buf + ng); - tile_h_x h0; - tile_f_x x_f, dy_f, dx_f, tmp; - tile_v rstd_t, c1, c2, c3; - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(rstd_t, grstd); - TLOAD(c2, gc2); - TLOAD(c3, gc3); - { - gm_c gg(gamma + c); - tile_h_1 hg; - tile_f_1 gf; - tile_v gv; - TLOAD(hg, gg); - TCVT(gf, hg); - TROWSUM(gv, gf); - TMUL(c1, gv, rstd_t); - } - TROWEXPANDMUL(dx_f, dy_f, c1); - TROWEXPANDMUL(tmp, x_f, c2); - TADD(dx_f, dx_f, tmp); - TROWEXPANDADD(dx_f, dx_f, c3); - TCVT(h0, dx_f); - TSTORE(gdx, h0); - } - } - - for (int g = 0; g < G; ++g) { - const int c0 = g * D; - tile_f_d acc, cur; - tile_h_d h0; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - gm_fC gdb(db + n * C + c0); - TLOAD(cur, gdb); - TADD(acc, acc, cur); - } - gm_c gout(dbeta + c0); - TCVT(h0, acc); - TSTORE(gout, h0); - } - for (int g = 0; g < G; ++g) { - const int c0 = g * D; - tile_f_d acc, ds_f, db_f, t0; - tile_h_d h0; - tile_v mean_t, rstd_t; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - const int ng = n * G + g; - gm_fC gds(ds + n * C + c0); - gm_fC gdb(db + n * C + c0); - gm_fG gmean(mean + ng); - gm_fG grstd(rstd + ng); - TLOAD(ds_f, gds); - TLOAD(db_f, gdb); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - TROWEXPANDMUL(t0, db_f, mean_t); - TSUB(t0, ds_f, t0); - TROWEXPANDMUL(t0, t0, rstd_t); - TADD(acc, acc, t0); - } - gm_c gout(dgamma + c0); - TCVT(h0, acc); - TSTORE(gout, h0); - } -} - -#endif // SUPERNPU_GROUP_NORM_GRAD_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp b/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp deleted file mode 100644 index ff4e675c..00000000 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/group_norm_grad_1d/group_norm_grad_1d.hpp +++ /dev/null @@ -1,532 +0,0 @@ -// ============================================================================= -// group_norm_grad_1d.hpp — GroupNorm backward, HxW == 1 (one-level PTO) -// ============================================================================= -// -// Matches PyTorch GroupNorm1dBackward -// (aten/src/ATen/native/cuda/group_norm_kernel.cu): -// Stage A1 reduce → c2/c3 (per n,g) -// Stage A2 dX = (rstd*gamma)*dY + c2*X + c3 -// Stage B dgamma / dbeta -// -// Layout: X/dY/dX [N,C]; mean/rstd [N,G] fp32; gamma/dgamma/dbeta [C] -// All pointers required (dy,x,mean,rstd,gamma,dx,dgamma,dbeta). -// -// tiling[4] = {N, C, G, tile_d} -// tile_d <= 0 → min(D, tD); channel R-split on dgamma/dbeta when D > tile_d. -// Stage A requires D <= tD (one tile). -// -// Tile capacity: logical tile >= 512B (TileOP IsValidActiveSize / TSize=1..7). -// fp16 → Cols>=256; fp32 → Cols>=128. tile_v is always fp32 → Cols=128. -// Reduce and dX are separate passes so large tiles do not stay live across both. -// -// Torch CUDA launch 总览 (NVIDIA, warp=32; HxW==1 特化): -// A1 Compute1dBackwardFusedParamsCUDAKernel -// grid=dim3(N,G), block=(D<512)?32:512 -// A2 dX gpu_kernel -// block=128, vt=4(fp16)/2(fp32), grid=ceil(N*C/(128*vt)) -// B GammaBeta1dBackwardCUDAKernel1/2 -// N<=128: grid=ceil(C/256), block=256 -// N>128: grid=ceil(C/32), block=dim3(32,16) -// ============================================================================= -#ifndef SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP -#define SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP - -#include - -#include - -namespace gn_grad_1d { - -// --------------------------------------------------------------------------- -// Stage A1: channel reduce → c2/c3 for one (n, g) -// scratch[2] = {c2, c3} -// -// Torch: Compute1dBackwardFusedParamsCUDAKernel -// grid = dim3(N, G) // blockIdx.x=n, blockIdx.y=g；本函数 = 其中一个 -// block = (D < 512) ? 32 : 512 -// 线程: threadIdx.x 沿 i∈[0,D) stride，读 dY/X/gamma 累加后 block reduce -// → sum1=Σ dY*X*gamma, sum2=Σ dY*gamma → c2,c3 -// --------------------------------------------------------------------------- -template -inline void fused_params_group(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *gamma, float *scratch, int64_t N, - int64_t C, int64_t G, int64_t D, int64_t n, - int64_t g, float s) { - const int64_t ng = n * G + g; - const int64_t c0 = g * D; - const int64_t offset = n * C + c0; - const size_t active_d = static_cast(D); - - gm_h gdy(dy + offset, static_cast(N), static_cast(C)); - gm_h gx(x + offset, static_cast(N), static_cast(C)); - gm_f gmean(mean + ng, static_cast(N * G), 1); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - gm_f gc2(scratch + 0, 1, 1); - gm_f gc3(scratch + 1, 1, 1); - - tile_h h0(1, active_d); - tile_h h1(1, active_d); - tile_f x_f(1, active_d); - tile_f dy_f(1, active_d); - tile_f t0(1, active_d); - tile_f t1(1, active_d); - tile_v mean_t(1); - tile_v rstd_t(1); - tile_v sum1(1); - tile_v sum2(1); - tile_v c2(1); - tile_v c3(1); - - // Torch: 各 thread 读本组一段通道；PTO 一次 Tile 覆盖整组 D - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - - { - gm_h gg(gamma + c0, 1, static_cast(C)); - TLOAD(h1, gg); - TCVT(t0, h1); // gamma - } - - // sum2 = Σ dy*gamma ; sum1 = Σ dy*gamma*x ↔ thread 局部累加 + BlockReduce - TMUL(t1, dy_f, t0); - TROWSUM(sum2, t1); - TMUL(t1, t1, x_f); - TROWSUM(sum1, t1); - - // c2 = (sum2*mean - sum1) * rstd^3 * s （归约后标量，通常 thread0 写） - TMUL(c2, sum2, mean_t); - TSUB(c2, c2, sum1); - TMUL(c3, rstd_t, rstd_t); - TMUL(c3, c3, rstd_t); - TMUL(c2, c2, c3); - TMULS(c2, c2, s); - - // c3 = -c2*mean - sum2*rstd*s - TMUL(c3, c2, mean_t); - TMULS(c3, c3, -1.0f); - TMUL(sum1, sum2, rstd_t); - TMULS(sum1, sum1, s); - TSUB(c3, c3, sum1); - - TSTORE(gc2, c2); - TSTORE(gc3, c3); -} - -// --------------------------------------------------------------------------- -// Stage A2: dX for one (n, g) from spilled c2/c3 -// -// Torch: gpu_kernel 元素级 -// block = 128 -// vt = 4 (fp16/bf16) / 2 (fp32+) -// grid = ceil(N*C / (128*vt)) -// 线程: 线性下标覆盖 [N,C]；c2/c3 按 (n,g) 广播到组内通道 -// 本函数一次写完一组 D 个通道（HxW=1） -// --------------------------------------------------------------------------- -template -inline void dx_group(dtype *dy, dtype *x, float *rstd, dtype *gamma, - float *scratch, dtype *dx, int64_t N, int64_t C, - int64_t G, int64_t D, int64_t n, int64_t g) { - const int64_t ng = n * G + g; - const int64_t c0 = g * D; - const int64_t offset = n * C + c0; - const size_t active_d = static_cast(D); - - gm_h gdy(dy + offset, static_cast(N), static_cast(C)); - gm_h gx(x + offset, static_cast(N), static_cast(C)); - gm_h gdx(dx + offset, static_cast(N), static_cast(C)); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - gm_f gc2(scratch + 0, 1, 1); - gm_f gc3(scratch + 1, 1, 1); - - tile_h h0(1, active_d); - tile_h h1(1, active_d); - tile_f x_f(1, active_d); - tile_f dy_f(1, active_d); - tile_f t0(1, active_d); - tile_f t1(1, active_d); - tile_v rstd_t(1); - tile_v c2(1); - tile_v c3(1); - - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(rstd_t, grstd); - TLOAD(c2, gc2); - TLOAD(c3, gc3); - - { - gm_h gg(gamma + c0, 1, static_cast(C)); - TLOAD(h1, gg); - TCVT(t0, h1); // gamma - } - - // dX = (rstd*gamma)*dY + c2*X + c3 - TROWEXPANDMUL(t1, t0, rstd_t); - TMUL(t1, t1, dy_f); - TROWEXPANDMUL(t0, x_f, c2); - TADD(t1, t1, t0); - TROWEXPANDADD(t1, t1, c3); - - TCVT(h0, t1); - TSTORE(gdx, h0); -} - -// --------------------------------------------------------------------------- -// Stage B: dbeta — dbeta[c] = Σ_n dY[n,c] -// -// Torch: GammaBeta1dBackwardCUDAKernel1/2（与 dgamma 同 launch） -// N<=128: grid=ceil(C/256), block=256；每线程一个 c，循环 n -// N>128: grid=ceil(C/32), block=dim3(32,16) -// --------------------------------------------------------------------------- -template -inline void dbeta_group(dtype *dy, dtype *dbeta, int64_t N, int64_t C, - int64_t D, int64_t tile_d, int64_t g) { - const int64_t c0 = g * D; - - for (int64_t d0 = 0; d0 < D; d0 += tile_d) { - const size_t vd = static_cast( - (d0 + tile_d <= D) ? tile_d : (D - d0)); - - tile_h h0(1, vd); - tile_f dy_f(1, vd); - tile_f acc(1, vd); - TEXPANDS(acc, 0.0f); - - // Torch Kernel1: 单线程 for(n) 累加；此处 Tile 一次累加一组通道 - for (int64_t n = 0; n < N; ++n) { - gm_h gdy(dy + n * C + c0 + d0, static_cast(N), - static_cast(C)); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TADD(acc, acc, dy_f); - } - - gm_h gdb(dbeta + c0 + d0, 1, static_cast(C)); - TCVT(h0, acc); - TSTORE(gdb, h0); - } -} - -// --------------------------------------------------------------------------- -// Stage B: dgamma — dgamma[c] = Σ_n dY*(X-mean)*rstd -// -// Torch: 与 dbeta 同 Kernel1/2 launch（见上） -// --------------------------------------------------------------------------- -template -inline void dgamma_group(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *dgamma, int64_t N, int64_t C, int64_t G, - int64_t D, int64_t tile_d, int64_t g) { - const int64_t c0 = g * D; - - for (int64_t d0 = 0; d0 < D; d0 += tile_d) { - const size_t vd = static_cast( - (d0 + tile_d <= D) ? tile_d : (D - d0)); - - tile_h h0(1, vd); - tile_f dy_f(1, vd); - tile_f x_f(1, vd); - tile_f t0(1, vd); - tile_f acc(1, vd); - tile_v mean_t(1); - tile_v rstd_t(1); - TEXPANDS(acc, 0.0f); - - for (int64_t n = 0; n < N; ++n) { - const int64_t ng = n * G + g; - const int64_t offset = n * C + c0 + d0; - - gm_h gdy(dy + offset, static_cast(N), static_cast(C)); - gm_h gx(x + offset, static_cast(N), static_cast(C)); - gm_f gmean(mean + ng, static_cast(N * G), 1); - gm_f grstd(rstd + ng, static_cast(N * G), 1); - - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - - TROWEXPANDMUL(t0, x_f, rstd_t); - TMUL(t0, t0, dy_f); - TROWEXPANDMUL(x_f, dy_f, mean_t); - TROWEXPANDMUL(x_f, x_f, rstd_t); - TSUB(t0, t0, x_f); - TADD(acc, acc, t0); - } - - gm_h gdg(dgamma + c0 + d0, 1, static_cast(C)); - TCVT(h0, acc); - TSTORE(gdg, h0); - } -} - -} // namespace gn_grad_1d - -// tiling: [N, C, G, tile_d] -// -// 入口循环 ↔ Torch grid： -// for n,g fused_params + dx_group ↔ grid=dim3(N,G) 再接 numel 上 gpu_kernel -// for g dbeta/dgamma ↔ Kernel1/2 按通道写回 -template -void group_norm_grad_1d(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *gamma, const int64_t *tiling, dtype *dx, - dtype *dgamma, dtype *dbeta) { - // Capacity in elements: every Tile buffer >= 512B (dtype strip + float strip). - constexpr int64_t tDDtype = - (512 + static_cast(sizeof(dtype)) - 1) / - static_cast(sizeof(dtype)); - constexpr int64_t tD = tDDtype > 128 ? tDDtype : 128; - constexpr int64_t tV = 128; // float scalar/broadcast strip: 128*4B = 512B - - const int64_t N = tiling[0]; - const int64_t C = tiling[1]; - const int64_t G = tiling[2]; - const int64_t D = C / G; - const int64_t tile_d = tiling[3] > 0 ? tiling[3] : (D < tD ? D : tD); - - using gm_h = global_tensor>; - using gm_f = global_tensor>; - using tile_h = - Tile; - using tile_f = - Tile; - using tile_v = - Tile; - - const float s = 1.0f / static_cast(D); - float scratch[2]; // c2, c3 for one (n,g) - - for (int64_t n = 0; n < N; ++n) { - for (int64_t g = 0; g < G; ++g) { - gn_grad_1d::fused_params_group(dy, x, mean, rstd, gamma, - scratch, N, C, G, D, n, g, - s); - gn_grad_1d::dx_group( - dy, x, rstd, gamma, scratch, dx, N, C, G, D, n, g); - } - } - - for (int64_t g = 0; g < G; ++g) { - gn_grad_1d::dbeta_group( - dy, dbeta, N, C, D, tile_d, g); - } - for (int64_t g = 0; g < G; ++g) { - gn_grad_1d::dgamma_group( - dy, x, mean, rstd, dgamma, N, C, G, D, tile_d, g); - } -} - -template -void group_norm_grad_1d(dtype *dy, dtype *x, float *mean, float *rstd, - dtype *gamma, dtype *dx, dtype *dgamma, - dtype *dbeta) { - static_assert(N > 0 && C > 0 && G > 0); - static_assert(C % G == 0); - constexpr int D = C / G; - constexpr int tDDtype = - (512 + static_cast(sizeof(dtype)) - 1) / - static_cast(sizeof(dtype)); - constexpr int tD = tDDtype > 128 ? tDDtype : 128; - constexpr int tV = 128; - constexpr int kTileD = tile_d > 0 ? tile_d : (D < tD ? D : tD); - static_assert(kTileD > 0 && kTileD <= tD && D <= tD); - constexpr int n_d = D / kTileD; - constexpr int rmd_d = D % kTileD; - constexpr float s = 1.0f / static_cast(D); - - using gm_h = global_tensor>; - using gm_c = global_tensor>; - using gm_f1 = global_tensor>; - using gm_fG = global_tensor>; - using tile_h_d = - Tile; - using tile_f_d = - Tile; - using tile_h_td = - Tile; - using tile_f_td = - Tile; - using tile_v = - Tile; - - float scratch[2]; - - for (int n = 0; n < N; ++n) { - for (int g = 0; g < G; ++g) { - const int ng = n * G + g; - const int c0 = g * D; - const int offset = n * C + c0; - gm_h gdy(dy + offset); - gm_h gx(x + offset); - gm_h gdx(dx + offset); - gm_fG gmean(mean + ng); - gm_fG grstd(rstd + ng); - gm_f1 gc2(scratch + 0); - gm_f1 gc3(scratch + 1); - gm_c gg(gamma + c0); - - tile_h_d h0, h1; - tile_f_d x_f, dy_f, t0, t1; - tile_v mean_t, rstd_t, sum1, sum2, c2, c3; - - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - TLOAD(h1, gg); - TCVT(t0, h1); - TMUL(t1, dy_f, t0); - TROWSUM(sum2, t1); - TMUL(t1, t1, x_f); - TROWSUM(sum1, t1); - TMUL(c2, sum2, mean_t); - TSUB(c2, c2, sum1); - TMUL(c3, rstd_t, rstd_t); - TMUL(c3, c3, rstd_t); - TMUL(c2, c2, c3); - TMULS(c2, c2, s); - TMUL(c3, c2, mean_t); - TMULS(c3, c3, -1.0f); - TMUL(sum1, sum2, rstd_t); - TMULS(sum1, sum1, s); - TSUB(c3, c3, sum1); - TSTORE(gc2, c2); - TSTORE(gc3, c3); - - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(rstd_t, grstd); - TLOAD(c2, gc2); - TLOAD(c3, gc3); - TLOAD(h1, gg); - TCVT(t0, h1); - TROWEXPANDMUL(t1, t0, rstd_t); - TMUL(t1, t1, dy_f); - TROWEXPANDMUL(t0, x_f, c2); - TADD(t1, t1, t0); - TROWEXPANDADD(t1, t1, c3); - TCVT(h0, t1); - TSTORE(gdx, h0); - } - } - - for (int g = 0; g < G; ++g) { - const int c0 = g * D; - for (int d0 = 0; d0 < n_d; ++d0) { - tile_h_td h0; - tile_f_td dy_f, acc; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - gm_h gdy(dy + n * C + c0 + d0 * kTileD); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TADD(acc, acc, dy_f); - } - gm_c gdb(dbeta + c0 + d0 * kTileD); - TCVT(h0, acc); - TSTORE(gdb, h0); - } - if constexpr (rmd_d) { - using tile_h_r = - Tile; - using tile_f_r = - Tile; - tile_h_r h0; - tile_f_r dy_f, acc; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - gm_h gdy(dy + n * C + c0 + n_d * kTileD); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TADD(acc, acc, dy_f); - } - gm_c gdb(dbeta + c0 + n_d * kTileD); - TCVT(h0, acc); - TSTORE(gdb, h0); - } - } - - for (int g = 0; g < G; ++g) { - const int c0 = g * D; - for (int d0 = 0; d0 < n_d; ++d0) { - tile_h_td h0; - tile_f_td dy_f, x_f, t0, acc; - tile_v mean_t, rstd_t; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - const int ng = n * G + g; - const int offset = n * C + c0 + d0 * kTileD; - gm_h gdy(dy + offset); - gm_h gx(x + offset); - gm_fG gmean(mean + ng); - gm_fG grstd(rstd + ng); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - TROWEXPANDMUL(t0, x_f, rstd_t); - TMUL(t0, t0, dy_f); - TROWEXPANDMUL(x_f, dy_f, mean_t); - TROWEXPANDMUL(x_f, x_f, rstd_t); - TSUB(t0, t0, x_f); - TADD(acc, acc, t0); - } - gm_c gdg(dgamma + c0 + d0 * kTileD); - TCVT(h0, acc); - TSTORE(gdg, h0); - } - if constexpr (rmd_d) { - using tile_h_r = - Tile; - using tile_f_r = - Tile; - tile_h_r h0; - tile_f_r dy_f, x_f, t0, acc; - tile_v mean_t, rstd_t; - TEXPANDS(acc, 0.0f); - for (int n = 0; n < N; ++n) { - const int ng = n * G + g; - const int offset = n * C + c0 + n_d * kTileD; - gm_h gdy(dy + offset); - gm_h gx(x + offset); - gm_fG gmean(mean + ng); - gm_fG grstd(rstd + ng); - TLOAD(h0, gdy); - TCVT(dy_f, h0); - TLOAD(h0, gx); - TCVT(x_f, h0); - TLOAD(mean_t, gmean); - TLOAD(rstd_t, grstd); - TROWEXPANDMUL(t0, x_f, rstd_t); - TMUL(t0, t0, dy_f); - TROWEXPANDMUL(x_f, dy_f, mean_t); - TROWEXPANDMUL(x_f, x_f, rstd_t); - TSUB(t0, t0, x_f); - TADD(acc, acc, t0); - } - gm_c gdg(dgamma + c0 + n_d * kTileD); - TCVT(h0, acc); - TSTORE(gdg, h0); - } - } -} - -#endif // SUPERNPU_GROUP_NORM_GRAD_1D_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm/rms_norm.hpp b/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm/rms_norm.hpp deleted file mode 100644 index 8a27ebab..00000000 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm/rms_norm.hpp +++ /dev/null @@ -1,167 +0,0 @@ -// ============================================================================= -// rms_norm.hpp — RMSNorm (one-level PTO) -// ============================================================================= -// -// Shape dims: A (outer / row), R (reduce / col). -// -// out[a] = x[a] * rsqrt(mean(x[a]^2) + eps) -// -// Entry: -// rms_norm(x, tiling, out, eps); -// tiling[4] = {g_a, g_r, tile_a, tile_r} (int64_t) -// tile_r <= 0 means use g_r (full-row tile). -// -// Pipeline (fp16 in/out, fp32 compute): -// TLOAD → TCVT → TMUL(x,x) → TROWSUM → TMULS(1/g_r) → TADDS(eps) -// → Newton rsqrt → TROWEXPANDMUL → TCVT → TSTORE -// -// Dynamic ValidRow/ValidCol: Tile Valid = -1, ctor passes runtime values. -// Full A tiles in the main loop; trailing rows handled separately. -// ============================================================================= -#ifndef SUPERNPU_RMS_NORM_PTO_HPP -#define SUPERNPU_RMS_NORM_PTO_HPP - -#include - -#include - -namespace rms_detail { - -template -inline void rsqrt_newton(TileVec &out, TileVec &a) { - TRSQRT(out, a); -} - -template -inline void rms_norm_tile(dtype *x, dtype *out, int64_t gA, int64_t gR, - int64_t a_off, int64_t active_a, int64_t active_r, - float inv_r, float eps) { - const int64_t offset = a_off * gR; - gm_t gi(x + offset, static_cast(gA), static_cast(gR)); - gm_t go(out + offset, static_cast(gA), static_cast(gR)); - - tile_h src_h(static_cast(active_a), - static_cast(active_r)); - tile_h dst_h(static_cast(active_a), - static_cast(active_r)); - tile_f src(static_cast(active_a), - static_cast(active_r)); - tile_f squared(static_cast(active_a), - static_cast(active_r)); - tile_f dst(static_cast(active_a), - static_cast(active_r)); - tile_v sqrsum(static_cast(active_a)); - tile_v mean(static_cast(active_a)); - tile_v denom(static_cast(active_a)); - tile_v rms(static_cast(active_a)); - - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(squared, src, src); - TROWSUM(sqrsum, squared); - TMULS(mean, sqrsum, inv_r); - TADDS(denom, mean, eps); - rsqrt_newton(rms, denom); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); -} - -} // namespace rms_detail - -// tiling: [g_a, g_r, tile_a, tile_r] -template -void rms_norm(dtype *x, const int64_t *tiling, dtype *out, float eps = 1e-6f) { - // Physical capacity (Rows×Cols); Valid comes from tiling (tile_a,tile_r). - // Size must cover ValidRow×ValidCol; SoftCore should not require Rows≥ValidRow. - constexpr int64_t tA = 1; - constexpr int64_t tR = 1024; - - const int64_t gA = tiling[0]; - const int64_t gR = tiling[1]; - const int64_t tile_a = tiling[2] > 0 ? tiling[2] : tA; - const int64_t tile_r = tiling[3] > 0 ? tiling[3] : gR; - - using gm_t = global_tensor>; - using tile_h = Tile; - using tile_f = Tile; - // Row expansion requires the broadcast source to have physical Cols=1. - // Keep a 512-byte carrier by placing the padding in Rows instead. - using tile_v = - Tile; - - const float inv_r = 1.0f / static_cast(gR); - - // Full A tiles; peel the last iteration for the trailing block. - int64_t ia = 0; - for (; ia + tile_a < gA; ia += tile_a) { - rms_detail::rms_norm_tile( - x, out, gA, gR, ia, tile_a, tile_r, inv_r, eps); - } - // Tail (or sole) block: ValidRow = remaining rows along A. - rms_detail::rms_norm_tile( - x, out, gA, gR, ia, gA - ia, tile_r, inv_r, eps); -} - -// Compile-time shape / tiling (gelu/gather style). tR must cover the full row. -template -void rms_norm(dtype *x, dtype *out, float eps = 1e-6f) { - static_assert(gA > 0 && gR > 0 && tA > 0 && tR > 0); - static_assert(tR == gR, "static rms_norm is a single R-tile; use rms_norm_binary for R-split"); - constexpr int Mb = gA / tA; - constexpr int rmd_A = gA % tA; - constexpr float inv_r = 1.0f / static_cast(gR); - - using gm_t = global_tensor>; - using tile_h = Tile; - using tile_f = Tile; - static_assert(tA <= 128, "RMSNorm row-state carrier supports tA <= 128"); - using tile_v = - Tile; - using it_t = global_iterator; - - it_t gI(x); - it_t gO(out); - - for (int ia = 0; ia < Mb; ++ia) { - tile_h src_h, dst_h; - tile_f src, squared, dst; - tile_v sqrsum, mean, denom, rms; - auto gi = gI(ia, 0); - auto go = gO(ia, 0); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(squared, src, src); - TROWSUM(sqrsum, squared); - TMULS(mean, sqrsum, inv_r); - TADDS(denom, mean, eps); - rms_detail::rsqrt_newton(rms, denom); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - if constexpr (rmd_A) { - using tile_h_r = Tile; - using tile_f_r = Tile; - using tile_v_r = - Tile; - tile_h_r src_h, dst_h; - tile_f_r src, squared, dst; - tile_v_r sqrsum, mean, denom, rms; - auto gi = gI(Mb, 0); - auto go = gO(Mb, 0); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(squared, src, src); - TROWSUM(sqrsum, squared); - TMULS(mean, sqrsum, inv_r); - TADDS(denom, mean, eps); - rms_detail::rsqrt_newton(rms, denom); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } -} - -#endif // SUPERNPU_RMS_NORM_PTO_HPP diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md b/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md deleted file mode 100644 index df25db21..00000000 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/SKILL.md +++ /dev/null @@ -1,189 +0,0 @@ ---- -name: rms-norm-binary -description: >- - Build, run, and debug the one-level rms_norm_binary kernel (R-split RMSNorm) - with SuperNPUBench + run_op.py + gfrun/gfsim precision checks. Use when - editing rms_norm_binary.hpp, rms_norm_binary tests, workspace/GetCacheId - reduce, TADD cross-tile sum, or verifying [1,8192] fp16 binary RMSNorm. - Shape dims are A (outer) and R (reduce): g_a/g_r, tile_a/tile_r, tA/tR. ---- - -# rms_norm_binary — kernel & verification - -Default root: `/home/wangyu/Code/SuperScalar`（下文 `$ROOT`）。 - -## Shape naming - -| Old | New | Meaning | -|-----|-----|---------| -| M / `g_m` / `tM` / `tile_m` | **A** / `g_a` / `tA` / `tile_a` | outer / row | -| N / `g_n` / `tN` / `tile_n` | **R** / `g_r` / `tR` / `tile_r` | reduce / col | -| `Nb` | `Rb` | `# R-tiles = ceil(g_r / tile_r)` | - -`tiling[4] = {g_a, g_r, tile_a, tile_r}`. - -## What it is - -One-level PTO RMSNorm when `g_r > tile_r`: - -```text -out[a] = x[a] * rsqrt(mean(x[a]^2) + eps) -``` - -R is split into `Rb = ceil(g_r / tile_r)` tiles. Each tile does local -`TROWSUM(x^2)`, then tiles are reduced to a full-row sum. - -Current default test shape: **`[1, 8192]`**, `tile_r=1024` → **`Rb=8`**, fp16. - -## Key paths - -| Role | Path | -|------|------| -| Kernel | `$ROOT/SuperNPUBench/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp` | -| Reference (single-tile) | `.../kernels/single_thread/normalization/rms_norm/rms_norm.hpp` | -| Testcase | `.../test/solution/normalization/rms_norm_binary/` | -| Host entry | `.../rms_norm_binary/src/rms_norm_binary.cpp` | -| Gen golden | `.../rms_norm_binary/src/gen_rms_norm_binary_data.py` | -| Compare | `.../rms_norm_binary/src/rms_norm_binary_data_compare.py` | -| Runner | `$ROOT/run_op.py` preset `rms_norm_binary` | -| Toolchain | `$ROOT/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin` → `COMPILER_DIR` | -| Sims | `$ROOT/SuperScalarModel/bin/gfrun`, `gfsim` | -| Related skill | `.../kernels/single_thread/reduction/binary-accumulation-cache-id/SKILL.md` | - -ELF after build: - -```text -.../output/kernel/normalization/rms_norm_binary/elf/ - kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half.elf -``` - -Compare dir (precision): - -```text -.../compare/kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half/ - input.bin golden.bin output.bin tiling_info.bin -``` - -## Kernel pipeline (current) - -File: `rms_norm_binary.hpp`. **No `rms_norm_dyn_ops.hpp`.** TEPL style like -`rms_norm.hpp`. - -```text -Pass1: - TLOAD(sum ← zeros) - per R-tile: TLOAD → TCVT → TMUL(x,x) → TROWSUM → TADD(sum, sum, cur) - -Pass1.5: - TMULS(mean, sum, 1/g_r) → TADDS(eps) → Newton rsqrt → rms - -Pass2 (per R-tile): - TLOAD → TCVT → TROWEXPANDMUL(x, rms) → TCVT → TSTORE -``` - -Important implementation notes: - -1. **Cross-tile sum is streaming** (`sum += cur`), not GetCacheId carry-merge. -2. Zero-init `sum` outside the R loop; uniform `TADD` inside (no first-tile branch). -3. `tile_v`: `Cols=32`, **static `Valid=1,1`** (`tile_a==1`) so TEPL `B.DIM` - immediates are legal. -4. `workspace` argument is kept in the API but **currently unused**. -5. Do **not** put early-return parameter checks in the kernel (caller owns tiling). - -## How to verify - -```bash -export COMPILER_DIR=$ROOT/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin -cd $ROOT -python3 run_op.py rms_norm_binary -``` - -What `run_op.py` does: - -1. `gen_rms_norm_binary_data.py` → write `input.bin` / `golden.bin` / tiling -2. `make TESTCASE=rms_norm_binary DType=__half res_check=on` → ELF with I/O -3. `gfrun` functional sim → writes `output.bin` -4. `rms_norm_binary_data_compare.py` → atol/rtol vs golden -5. Rebuild **without** `res_check` (res_check ELF often crashes `gfsim`) -6. `gfsim` performance / cycle sim -7. After compile: write `.diss` via `llvm-objdump -dl` (disable: `--no-diss`) - -Useful flags: - -```bash -python3 run_op.py rms_norm_binary --func-only # gfrun + precision only -python3 run_op.py rms_norm_binary --perf-only # gfsim only -python3 run_op.py rms_norm_binary --compile-only -python3 run_op.py rms_norm_binary --no-check-precision -python3 run_op.py rms_norm_binary --skip-compile -python3 run_op.py rms_norm_binary --no-diss -``` - -Manual make (same case): - -```bash -cd $ROOT/SuperNPUBench/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary -make TESTCASE=rms_norm_binary DType=__half COMPILER_DIR=$COMPILER_DIR -# or: bash compile.all -``` - -## Expected results (as of current kernel) - -| Step | Typical result | -|------|----------------| -| Compile | OK | -| gfrun | **PASS** (`Success to Reach the End`) | -| Precision | **PASS** (`max_abs` often `0.0` on `[1,8192]`) | -| gfsim | **FAIL** | - -### gfsim failure (known) - -Symptom: - -```text -TMOV ... INVALID -FATAL: gfsim received signal 11 -# or: Bank store offset out of range! -``` - -Cause: compiler-inserted `TMOV` (tile rename) for small `tile_v` kept live across -R-split loops; timing sim corrupts tile metadata. **Not a golden mismatch** — -gfrun + compare already pass. - -Baseline `rms_norm` (no cross-tile accumulate / second R loop) usually **PASS**es gfsim. - -## Testcase layout - -`rms_norm_binary.cpp` defaults: - -```cpp -G_A=1, G_R=8192, TILE_A=1, TILE_R=1024 -workspace_buf[K_MAX_LEVELS * G_A * K_WS_COLS] // kept for ABI; unused by kernel -``` - -Precision scripts default shape `--g-r 8192`, `--tile-r 1024`. - -`run_op.py` preset name is exactly **`rms_norm_binary`** (no size suffix). - -## Agent checklist when changing the kernel - -1. Keep compute TEPL-only; do not reintroduce `rms_norm_dyn_ops.hpp` unless asked. -2. Prefer `Valid=1,1` on `tile_v` when using TEPL ops with NTTP `B.DIM`. -3. Avoid taking addresses of `tile_v` / large pointer arrays of tiles (Liveouts / - illegal spill). -4. Do not mix `TROWSUM` u-reg lineage with `TLOAD` of small reduce tiles in the - same hot path without verifying Match Instruction / gfsim. -5. After edits: `python3 run_op.py rms_norm_binary` (or `--func-only` if only - checking correctness). -6. If implementing true GetCacheId carry + workspace reload, also read - `binary-accumulation-cache-id/SKILL.md` and expect toolchain/sim constraints - above. - -## Anti-patterns - -- Naming the run_op preset `rms_norm_binary_1x8192` / `..._1x32768` — canonical - name is `rms_norm_binary`. -- Treating gfsim FAIL as a precision bug when gfrun+compare already PASS. -- Putting workspace spill between `rsqrt` and `TROWEXPANDMUL` (clobbers `rms`). -- Using TEPL `TADD` with `Valid=-1` (`Match Instruction Error`). -- Reverting shape names to M/N — use **A/R** consistently. diff --git a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp b/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp deleted file mode 100644 index d89a378e..00000000 --- a/benchmark/one-level-arch/kernels/single_thread/normalization/rms_norm_binary/rms_norm_binary.hpp +++ /dev/null @@ -1,377 +0,0 @@ -// ============================================================================= -// rms_norm_binary.hpp — RMSNorm for g_r > tile_r (R-split) -// ============================================================================= -// -// tiling[5] = {g_a, g_r, tile_a, tile_r, pow_r} -// -// 每块 RowSum 后立刻 UpdateCache（workspace = cacheBuffer），对齐 AscendC： -// DataCopy(aReg, src); -// for (j = 0; j < cid; ++j) { -// DataCopy(bReg, cache + j * stride); -// Add(aReg, aReg, bReg); -// } -// DataCopy(cache + cid * stride, aReg); -// cid = GetCacheId(idx) = ctz(idx+1) -// sum = cache[GetCacheId(r-1)] （r 为 2^k） -// -// workspace: [0, kMaxLevels) cache 档 -// ============================================================================= -#ifndef SUPERNPU_RMS_NORM_BINARY_PTO_HPP -#define SUPERNPU_RMS_NORM_BINARY_PTO_HPP - -#include - -#include - -namespace rms_bin { - -constexpr int kWsCols = 128; -constexpr int kMaxLevels = 6; - -inline int64_t GetCacheId(int64_t idx) { - return static_cast( - __builtin_ctzll(static_cast(idx + 1))); -} - -template -inline void rsqrt_newton(TileVec &out, TileVec &a) { - TRSQRT(out, a); -} - -} // namespace rms_bin - -template -void rms_norm_binary(dtype *x, const int64_t *tiling, dtype *out, - float *workspace, float eps = 1e-6f) { - constexpr int64_t tA = 1; - constexpr int64_t tR = 1024; - - const int64_t gA = tiling[0]; - const int64_t gR = tiling[1]; - const int64_t tile_r = tiling[3] > 0 ? tiling[3] : tR; - const int64_t powR = tiling[4]; - - const int64_t remR = gR - powR; - const int64_t headR = powR - remR; - const int64_t n_rem_full = remR / tile_r; - const int64_t rem_tail = remR - n_rem_full * tile_r; - const int64_t n_head_full = headR / tile_r; - const int64_t head_tail = headR - n_head_full * tile_r; - const int64_t n_full = gR / tile_r; - const int64_t tail_r = gR - n_full * tile_r; - const float inv_r = 1.0f / static_cast(gR); - - using gm_t = global_tensor>; - using gm_f = global_tensor>; - using tile_h = Tile; - using tile_f = Tile; - using tile_v = Tile; - - for (int64_t ia = 0; ia < gA; ++ia) { - constexpr size_t active_a = 1; - const size_t full_r = static_cast(tile_r); - - tile_v cur, buf, sum, mean, denom, rms, zero; - TEXPANDS(zero, 0.0f); - - float *cache = workspace + ia * rms_bin::kWsCols; - const int64_t stride = gA * rms_bin::kWsCols; - - for (int64_t lv = 0; lv < rms_bin::kMaxLevels; ++lv) { - gm_f go(cache + lv * stride, 1, rms_bin::kWsCols); - TSTORE(go, zero); - } - - int64_t r = 0; - - // UpdateCache（AscendC 同构） -#define RMS_BIN_UPDATE_CACHE() \ - do { \ - const uint16_t cid = \ - static_cast(rms_bin::GetCacheId(r)); \ - for (uint16_t j = 0; j < cid; ++j) { \ - gm_f gj(cache + static_cast(j) * stride, 1, \ - rms_bin::kWsCols); \ - TLOAD(buf, gj); \ - TADD(cur, cur, buf); \ - } \ - gm_f gc(cache + static_cast(cid) * stride, 1, \ - rms_bin::kWsCols); \ - TSTORE(gc, cur); \ - ++r; \ - } while (0) - - for (int64_t tr = 0; tr < n_rem_full; ++tr) { - const int64_t offset = ia * gR + tr * tile_r; - gm_t gi0(x + offset, static_cast(gA), static_cast(gR)); - gm_t gi1(x + offset + powR, static_cast(gA), - static_cast(gR)); - tile_h src0_h(active_a, full_r); - tile_h src1_h(active_a, full_r); - tile_f src0(active_a, full_r); - tile_f src1(active_a, full_r); - tile_f sq0(active_a, full_r); - tile_f sq1(active_a, full_r); - - TLOAD(src0_h, gi0); - TLOAD(src1_h, gi1); - TCVT(src0, src0_h); - TCVT(src1, src1_h); - TMUL(sq0, src0, src0); - TMUL(sq1, src1, src1); - TADD(sq0, sq0, sq1); - TROWSUM(cur, sq0); - RMS_BIN_UPDATE_CACHE(); - } - - if (rem_tail > 0) { - const int64_t offset = ia * gR + n_rem_full * tile_r; - const size_t ar = static_cast(rem_tail); - gm_t gi0(x + offset, static_cast(gA), static_cast(gR)); - gm_t gi1(x + offset + powR, static_cast(gA), - static_cast(gR)); - tile_h src0_h(active_a, ar); - tile_h src1_h(active_a, ar); - tile_f src0(active_a, ar); - tile_f src1(active_a, ar); - tile_f sq0(active_a, ar); - tile_f sq1(active_a, ar); - - TLOAD(src0_h, gi0); - TLOAD(src1_h, gi1); - TCVT(src0, src0_h); - TCVT(src1, src1_h); - TMUL(sq0, src0, src0); - TMUL(sq1, src1, src1); - TADD(sq0, sq0, sq1); - TROWSUM(cur, sq0); - RMS_BIN_UPDATE_CACHE(); - } - - for (int64_t tr = 0; tr < n_head_full; ++tr) { - const int64_t offset = ia * gR + remR + tr * tile_r; - gm_t gi(x + offset, static_cast(gA), static_cast(gR)); - tile_h src_h(active_a, full_r); - tile_f src(active_a, full_r); - tile_f sq(active_a, full_r); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(sq, src, src); - TROWSUM(cur, sq); - RMS_BIN_UPDATE_CACHE(); - } - if (head_tail > 0) { - const int64_t offset = ia * gR + remR + n_head_full * tile_r; - const size_t ar = static_cast(head_tail); - gm_t gi(x + offset, static_cast(gA), static_cast(gR)); - tile_h src_h(active_a, ar); - tile_f src(active_a, ar); - tile_f sq(active_a, ar); - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(sq, src, src); - TROWSUM(cur, sq); - RMS_BIN_UPDATE_CACHE(); - } -#undef RMS_BIN_UPDATE_CACHE - - { - const int64_t rid = r > 0 ? rms_bin::GetCacheId(r - 1) : 0; - gm_f gr(cache + rid * stride, 1, rms_bin::kWsCols); - TLOAD(sum, gr); - } - - TMULS(mean, sum, inv_r); - TADDS(denom, mean, eps); - rms_bin::rsqrt_newton(rms, denom); - - for (int64_t tr = 0; tr < n_full; ++tr) { - const int64_t offset = ia * gR + tr * tile_r; - gm_t gi(x + offset, static_cast(gA), static_cast(gR)); - gm_t go(out + offset, static_cast(gA), static_cast(gR)); - tile_h src_h(active_a, full_r); - tile_h dst_h(active_a, full_r); - tile_f src(active_a, full_r); - tile_f dst(active_a, full_r); - TLOAD(src_h, gi); - TCVT(src, src_h); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - if (tail_r > 0) { - const int64_t offset = ia * gR + n_full * tile_r; - const size_t ar = static_cast(tail_r); - gm_t gi(x + offset, static_cast(gA), static_cast(gR)); - gm_t go(out + offset, static_cast(gA), static_cast(gR)); - tile_h src_h(active_a, ar); - tile_h dst_h(active_a, ar); - tile_f src(active_a, ar); - tile_f dst(active_a, ar); - TLOAD(src_h, gi); - TCVT(src, src_h); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - } -} - -// Compile-time shape / tiling. Same algorithm as the dynamic entry. -template -void rms_norm_binary(dtype *x, dtype *out, float *workspace, - float eps = 1e-6f) { - static_assert(gA > 0 && gR > 0 && tA == 1 && tR > 0 && powR > 0); - static_assert(powR < gR && gR <= 2 * powR); - constexpr int remR = gR - powR; - constexpr int headR = powR - remR; - constexpr int n_rem_full = remR / tR; - constexpr int rem_tail = remR % tR; - constexpr int n_head_full = headR / tR; - constexpr int head_tail = headR % tR; - constexpr int n_full = gR / tR; - constexpr int tail_r = gR % tR; - constexpr float inv_r = 1.0f / static_cast(gR); - - using gm_t = global_tensor>; - using gm_f = global_tensor>; - using tile_h = Tile; - using tile_f = Tile; - using tile_v = Tile; - - for (int64_t ia = 0; ia < gA; ++ia) { - tile_v cur, buf, sum, mean, denom, rms, zero; - TEXPANDS(zero, 0.0f); - - float *cache = workspace + ia * rms_bin::kWsCols; - const int64_t stride = static_cast(gA) * rms_bin::kWsCols; - - for (int64_t lv = 0; lv < rms_bin::kMaxLevels; ++lv) { - gm_f go(cache + lv * stride); - TSTORE(go, zero); - } - - int64_t r = 0; - -#define RMS_BIN_UPDATE_CACHE_S() \ - do { \ - const uint16_t cid = \ - static_cast(rms_bin::GetCacheId(r)); \ - for (uint16_t j = 0; j < cid; ++j) { \ - gm_f gj(cache + static_cast(j) * stride); \ - TLOAD(buf, gj); \ - TADD(cur, cur, buf); \ - } \ - gm_f gc(cache + static_cast(cid) * stride); \ - TSTORE(gc, cur); \ - ++r; \ - } while (0) - - for (int tr = 0; tr < n_rem_full; ++tr) { - const int64_t offset = ia * gR + tr * tR; - gm_t gi0(x + offset); - gm_t gi1(x + offset + powR); - tile_h src0_h, src1_h; - tile_f src0, src1, sq0, sq1; - TLOAD(src0_h, gi0); - TLOAD(src1_h, gi1); - TCVT(src0, src0_h); - TCVT(src1, src1_h); - TMUL(sq0, src0, src0); - TMUL(sq1, src1, src1); - TADD(sq0, sq0, sq1); - TROWSUM(cur, sq0); - RMS_BIN_UPDATE_CACHE_S(); - } - if constexpr (rem_tail) { - using tile_h_r = - Tile; - using tile_f_r = - Tile; - const int64_t offset = ia * gR + n_rem_full * tR; - gm_t gi0(x + offset); - gm_t gi1(x + offset + powR); - tile_h_r src0_h, src1_h; - tile_f_r src0, src1, sq0, sq1; - TLOAD(src0_h, gi0); - TLOAD(src1_h, gi1); - TCVT(src0, src0_h); - TCVT(src1, src1_h); - TMUL(sq0, src0, src0); - TMUL(sq1, src1, src1); - TADD(sq0, sq0, sq1); - TROWSUM(cur, sq0); - RMS_BIN_UPDATE_CACHE_S(); - } - for (int tr = 0; tr < n_head_full; ++tr) { - const int64_t offset = ia * gR + remR + tr * tR; - gm_t gi(x + offset); - tile_h src_h; - tile_f src, sq; - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(sq, src, src); - TROWSUM(cur, sq); - RMS_BIN_UPDATE_CACHE_S(); - } - if constexpr (head_tail) { - using tile_h_r = - Tile; - using tile_f_r = - Tile; - const int64_t offset = ia * gR + remR + n_head_full * tR; - gm_t gi(x + offset); - tile_h_r src_h; - tile_f_r src, sq; - TLOAD(src_h, gi); - TCVT(src, src_h); - TMUL(sq, src, src); - TROWSUM(cur, sq); - RMS_BIN_UPDATE_CACHE_S(); - } -#undef RMS_BIN_UPDATE_CACHE_S - - { - const int64_t rid = r > 0 ? rms_bin::GetCacheId(r - 1) : 0; - gm_f gr(cache + rid * stride); - TLOAD(sum, gr); - } - - TMULS(mean, sum, inv_r); - TADDS(denom, mean, eps); - rms_bin::rsqrt_newton(rms, denom); - - for (int tr = 0; tr < n_full; ++tr) { - const int64_t offset = ia * gR + tr * tR; - gm_t gi(x + offset); - gm_t go(out + offset); - tile_h src_h, dst_h; - tile_f src, dst; - TLOAD(src_h, gi); - TCVT(src, src_h); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - if constexpr (tail_r) { - using tile_h_r = - Tile; - using tile_f_r = - Tile; - const int64_t offset = ia * gR + n_full * tR; - gm_t gi(x + offset); - gm_t go(out + offset); - tile_h_r src_h, dst_h; - tile_f_r src, dst; - TLOAD(src_h, gi); - TCVT(src, src_h); - TROWEXPANDMUL(dst, src, rms); - TCVT(dst_h, dst); - TSTORE(go, dst_h); - } - } -} - -#endif // SUPERNPU_RMS_NORM_BINARY_PTO_HPP diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/Makefile b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/Makefile deleted file mode 100644 index 8671b49e..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/Makefile +++ /dev/null @@ -1,10 +0,0 @@ -TESTCASE ?= rms_norm - -ifeq ($(TESTCASE),rms_norm) -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/rms_norm.cpp -TARGET = $(ELF_HEAD)_rms_norm_PE4.elf -else -$(error Unsupported TESTCASE=$(TESTCASE); expected rms_norm) -endif - -include ../../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/compile.all b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/compile.all deleted file mode 100644 index 88679c80..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/compile.all +++ /dev/null @@ -1,4 +0,0 @@ -#!/bin/bash -set -euo pipefail -: "${COMPILER_DIR:?Set COMPILER_DIR to the Linx compiler bin directory}" -make TESTCASE=rms_norm COMPILER_DIR="$COMPILER_DIR" diss diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/src/rms_norm.cpp b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/src/rms_norm.cpp deleted file mode 100644 index 702c29bb..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm/src/rms_norm.cpp +++ /dev/null @@ -1,38 +0,0 @@ -#include "benchmark.h" -#include "fileop.h" -#include "multi_thread_res_check.h" -#include "multi_thread/normalization/rms_norm/rms_norm.hpp" - -namespace { -constexpr int kPeCount = 4; -constexpr int kRowsPerPe = 1; -constexpr int kColumns = 8192; -alignas(4096) __half input[kPeCount * kRowsPerPe * kColumns]; -alignas(4096) __half output[kPeCount * kRowsPerPe * kColumns]; -#ifdef RES_CHECK -MultiThreadResCheckSync res_check_sync{}; -#endif -} // namespace - -int main() { - const std::uint32_t tid = get_thread_idx(); -#ifdef RES_CHECK - if (tid == 0) { - readBinaryFile(CHK_DIR "/input.bin", reinterpret_cast(input), - sizeof(input)); - } - res_check_publish_inputs(res_check_sync, tid); -#endif - BENCHSTART; - rms_norm<__half, kRowsPerPe, kPeCount * kRowsPerPe, - kColumns, 1, kColumns>(input, output); - BENCHEND; -#ifdef RES_CHECK - res_check_wait_for_all(res_check_sync, tid); - if (tid == 0) { - writeBinaryFile(CHK_DIR "/output.bin", - reinterpret_cast(output), sizeof(output)); - } -#endif - return 0; -} diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/Makefile b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/Makefile deleted file mode 100644 index b127a664..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/Makefile +++ /dev/null @@ -1,10 +0,0 @@ -TESTCASE ?= rms_norm_binary - -ifeq ($(TESTCASE),rms_norm_binary) -SRC_FILE += $(TEST_ROOT)/$(CASE_SRC_DIR)/rms_norm_binary.cpp -TARGET = $(ELF_HEAD)_rms_norm_binary_PE4.elf -else -$(error Unsupported TESTCASE=$(TESTCASE); expected rms_norm_binary) -endif - -include ../../../../common/Makefile.common diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/compile.all b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/compile.all deleted file mode 100644 index 94b4f5a0..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/compile.all +++ /dev/null @@ -1,4 +0,0 @@ -#!/bin/bash -set -euo pipefail -: "${COMPILER_DIR:?Set COMPILER_DIR to the Linx compiler bin directory}" -make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" diss diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/src/rms_norm_binary.cpp b/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/src/rms_norm_binary.cpp deleted file mode 100644 index 86cd790c..00000000 --- a/benchmark/one-level-arch/test/kernel/multi_thread/normalization/rms_norm_binary/src/rms_norm_binary.cpp +++ /dev/null @@ -1,37 +0,0 @@ -#include "benchmark.h" -#include "fileop.h" -#include "multi_thread_res_check.h" -#include "multi_thread/normalization/rms_norm_binary/rms_norm_binary.hpp" - -namespace { -alignas(4096) __half input[4 * 8192]; -alignas(4096) __half output[4 * 8192]; -alignas(4096) float - workspace[4 * rms_bin::kMaxLevels * rms_bin::kWsCols]; -#ifdef RES_CHECK -MultiThreadResCheckSync res_check_sync{}; -#endif -} // namespace - -int main() { - const std::uint32_t tid = get_thread_idx(); -#ifdef RES_CHECK - if (tid == 0) { - readBinaryFile(CHK_DIR "/input.bin", reinterpret_cast(input), - sizeof(input)); - } - res_check_publish_inputs(res_check_sync, tid); -#endif - BENCHSTART; - supernpu::multi_thread::rms_norm_binary< - __half, 4, 8192, 1, 1024, 4096>(input, output, workspace); - BENCHEND; -#ifdef RES_CHECK - res_check_wait_for_all(res_check_sync, tid); - if (tid == 0) { - writeBinaryFile(CHK_DIR "/output.bin", - reinterpret_cast(output), sizeof(output)); - } -#endif - return 0; -} diff --git a/benchmark/one-level-arch/test/kernel/multi_thread/res_check_all.py b/benchmark/one-level-arch/test/kernel/multi_thread/res_check_all.py index 78b08ab5..610332ea 100644 --- a/benchmark/one-level-arch/test/kernel/multi_thread/res_check_all.py +++ b/benchmark/one-level-arch/test/kernel/multi_thread/res_check_all.py @@ -95,13 +95,6 @@ def prep_gather(case_dir: Path) -> np.ndarray: return table[indexes].reshape(-1) -def prep_rms(case_dir: Path) -> np.ndarray: - rng = np.random.default_rng(2) - x = rng.uniform(-1.0, 1.0, (4, 8192)).astype(np.float16) - write(case_dir, "input.bin", x) - xf = x.astype(np.float32) - return (xf / np.sqrt(np.mean(xf * xf, axis=1, keepdims=True) + 1e-6)).astype(np.float16).reshape(-1) - def prep_rows(case_dir: Path, operation: str) -> np.ndarray: rng = np.random.default_rng(3) @@ -180,8 +173,6 @@ def prep_fa(case_dir: Path) -> np.ndarray: Case("matmul_shared", "matmul/elf/kernel_multi_thread_matmul_matmul_shared_B1_M256_N256_K256_tM128_tN256_tK128.elf", prep_matmul, output_name="res.bin", atol=1e-3, rtol=1e-3), Case("matmul_reuseB", "matmul/elf/kernel_multi_thread_matmul_matmul_reuseB_B1_M256_N256_K256_tM128_tN256_tK128.elf", prep_matmul, output_name="res.bin", atol=1e-3, rtol=1e-3), Case("matmul_lowp_fp8", "matmul/elf/kernel_multi_thread_matmul_matmul_lowp_FP8_B1_M256_N256_K512_tM128_tN256_tK512.elf", prep_matmul_lowp, output_name="res.bin"), - Case("rms_norm", "normalization/rms_norm/elf/kernel_multi_thread_normalization_rms_norm_rms_norm_PE4.elf", prep_rms, output_dtype=np.float16, atol=3e-2, rtol=3e-2), - Case("rms_norm_binary", "normalization/rms_norm_binary/elf/kernel_multi_thread_normalization_rms_norm_binary_rms_norm_binary_PE4.elf", prep_rms, output_dtype=np.float16, atol=3e-2, rtol=3e-2), Case("cumsum_row", "reduction/cumsum_row/elf/kernel_multi_thread_reduction_cumsum_row_cumsum_row_PE4.elf", lambda p: prep_rows(p, "cumsum"), atol=2e-3, rtol=2e-3), Case("reducemax_row", "reduction/reducemax_row/elf/kernel_multi_thread_reduction_reducemax_row_reducemax_row_PE4.elf", lambda p: prep_rows(p, "max")), Case("reduceprod_row", "reduction/reduceprod_row/elf/kernel_multi_thread_reduction_reduceprod_row_reduceprod_row_PE4.elf", lambda p: prep_rows(p, "prod"), atol=2e-3, rtol=2e-3),