separate hd pad/unpad kernels

92494a8a · danyao12 · b16fa5f0 · 92494a8a · 92494a8a · 92494a8a
Commit 92494a8a authored Jan 24, 2025 by danyao12
20 changed files
--- a/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
+++ b/example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtna.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtna.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtna_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtna_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtne.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtne.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtne_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtne_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtz.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtz.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtz_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a16_rtz_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtna.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtna.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtna_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtna_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtne.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtne.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtne_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtne_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtz.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtz.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtz_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_a32_rtz_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtna.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtna.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtna_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtna_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtne.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtne.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtne_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtne_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtz.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtz.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtz_pddv.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a16_rtz_pddv.cpp
--- a/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a32_rtna.cpp
+++ b/example/ck_tile/01_fmha/hsaco/bwd_bf16_causal_a32_rtna.cpp