Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
11 changes: 6 additions & 5 deletions artifacts/blogs/flash-attention-4/MANIFEST.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -4,11 +4,12 @@ code_present: true
total_blocks: 2
generated_by: scripts/extract_blog_code.py
files:
- local_path: code/01-software-exp-cody-waite-horner.cu
heading_path: '## Key Code > ### Software exp (Cody-Waite + Horner)'
- local_path: code/01-software-exp-published-range-reduction-and-rounded-polynomia.cu
heading_path: '## Illustrative Code > ### Software exp (published range reduction
and rounded polynomial)'
fence_lang: cuda
sha256: 6fba9c537831bba7b175ee3a050e3adb5ddf99193dee550b7122a4663e3f9164
sha256: 89c085df18aa5ae9b8693d1085bf43881b8bc3644068761e0312c9ee7628e71c
- local_path: code/02-2-cta-cooperative-backward.cu
heading_path: '## Key Code > ### 2-CTA cooperative backward'
heading_path: '## Illustrative Code > ### 2-CTA cooperative backward'
fence_lang: cuda
sha256: e089feff40c10ebf42fa1ff878f775c21f4f044fc4b6f2f510b43c2666c7e219
sha256: 5747bb5ef24ec8a82d9fd80d433d0dc231d0de042af83eeec9414b03a1a53de9

This file was deleted.

Original file line number Diff line number Diff line change
@@ -0,0 +1,15 @@
// Extracted from sources/blogs/flash-attention-4.md by scripts/extract_blog_code.py
// Heading: ## Illustrative Code > ### Software exp (published range reduction and rounded polynomial)
// Original fence language: cuda
// See artifacts/blogs/flash-attention-4/code/PROVENANCE.yaml for origin + license metadata.

// KernelWiki scalar illustration derived from the FA4 blog equations.
// This is not verbatim upstream FA4 code and omits selection and clamping.
#include <cmath>

__host__ __device__ inline float fa4_blog_exp2_reference(float x) {
const int n = static_cast<int>(floorf(x));
const float f = x - static_cast<float>(n); // f in [0, 1)
const float p = 1.0f + f * (0.6951f + f * (0.2276f + f * 0.0771f));
return ldexpf(p, n);
}
Original file line number Diff line number Diff line change
@@ -1,10 +1,14 @@
// Extracted from sources/blogs/flash-attention-4.md by scripts/extract_blog_code.py
// Heading: ## Key Code > ### 2-CTA cooperative backward
// Heading: ## Illustrative Code > ### 2-CTA cooperative backward
// Original fence language: cuda
// See artifacts/blogs/flash-attention-4/code/PROVENANCE.yaml for origin + license metadata.

// 2-CTA cooperative backward: paired CTAs in a cluster share a single TMEM
// accumulator half, halving SMEM traffic for dK/dV accumulation.
asm volatile(
"tcgen05.mma.cta_group::2.kind::f16 [%0], %1, %2, %3, 1;"
: : "r"(tmem_acc_shared), "l"(desc_a), "l"(desc_b), "r"(0));
// KernelWiki schematic derived from the FA4 paper/blog dimensions.
// This is not upstream inline PTX or a complete kernel.
struct Fa4TwoCtaBackwardShape {
static constexpr int cta_group = 2;
static constexpr int mma_m = 256;
static constexpr int mma_n = 128;
static constexpr int mma_k = 128;
static constexpr int backward_gemm_count = 5;
};
15 changes: 8 additions & 7 deletions artifacts/blogs/flash-attention-4/code/PROVENANCE.yaml
Original file line number Diff line number Diff line change
@@ -1,21 +1,22 @@
origin_url: https://tridao.me/blog/2026/flash4/
upstream_repo: blog
upstream_sha: c9a560f44002da92d82680036e3482d0ac4939a3
upstream_sha: none
license: inherits-from-source-blog
retrieved_at: 2026-04-27
retrieved_at: 2026-08-08
asset_mode: extracted
generated_by: scripts/extract_blog_code.py
size_cap_truncated: false
files:
- local_path: 01-software-exp-cody-waite-horner.cu
- local_path: 01-software-exp-published-range-reduction-and-rounded-polynomia.cu
role: extracted-block
mode: extracted
upstream_path: inline-in-blog-markdown
heading_path: '## Key Code > ### Software exp (Cody-Waite + Horner)'
sha256: 6fba9c537831bba7b175ee3a050e3adb5ddf99193dee550b7122a4663e3f9164
heading_path: '## Illustrative Code > ### Software exp (published range reduction
and rounded polynomial)'
sha256: 89c085df18aa5ae9b8693d1085bf43881b8bc3644068761e0312c9ee7628e71c
- local_path: 02-2-cta-cooperative-backward.cu
role: extracted-block
mode: extracted
upstream_path: inline-in-blog-markdown
heading_path: '## Key Code > ### 2-CTA cooperative backward'
sha256: e089feff40c10ebf42fa1ff878f775c21f4f044fc4b6f2f510b43c2666c7e219
heading_path: '## Illustrative Code > ### 2-CTA cooperative backward'
sha256: 5747bb5ef24ec8a82d9fd80d433d0dc231d0de042af83eeec9414b03a1a53de9
19 changes: 10 additions & 9 deletions artifacts/blogs/flashmla/MANIFEST.yaml
Original file line number Diff line number Diff line change
@@ -1,14 +1,15 @@
slug: flashmla
origin_url: https://github.com/deepseek-ai/FlashMLA
origin_url: https://github.com/deepseek-ai/FlashMLA/tree/71c737929f2567bd0a094ae140f8f60f390b1232
code_present: true
total_blocks: 2
generated_by: scripts/extract_blog_code.py
files:
- local_path: code/01-mla-decode-inner-loop.cu
heading_path: '## Key Code > ### MLA decode inner loop'
fence_lang: cuda
sha256: 4d3950b24b3da7a36e1db9d60a0492ed963d1b49ad6a1b97adeab54abaf4548f
- local_path: code/02-sparse-mla-kv-retrieval-kernel-v3-2.cu
heading_path: '## Key Code > ### Sparse-MLA KV-retrieval kernel (V3.2)'
fence_lang: cuda
sha256: 9b0499c0ddb77a9cebd370e84929d7149ebdc03163768369c95a12ae0b5c4f7f
- local_path: code/01-v3-fp8-sparse-decode-byte-check.py
heading_path: '## Exact V3 FP8 Sparse-Decode Layout > ### V3 FP8 sparse-decode byte
check'
fence_lang: python
sha256: bfc0a6df6b2e04ce4f894a9df04926d4120b4e370354d68e80a55c1645565c7e
- local_path: code/02-decode-page-index-round-trip.py
heading_path: '## Sparse Index Contracts > ### Decode page-index round trip'
fence_lang: python
sha256: a0bbc6f0248adc67b8113e3faa66c1be8924aafce665ebf89b43f0e2f0e4a405
31 changes: 0 additions & 31 deletions artifacts/blogs/flashmla/code/01-mla-decode-inner-loop.cu

This file was deleted.

Original file line number Diff line number Diff line change
@@ -0,0 +1,16 @@
# Extracted from sources/blogs/flashmla.md by scripts/extract_blog_code.py
# Heading: ## Exact V3 FP8 Sparse-Decode Layout > ### V3 FP8 sparse-decode byte check
# Original fence language: python
# See artifacts/blogs/flashmla/code/PROVENANCE.yaml for origin + license metadata.

# KernelWiki-derived contract check; not upstream FlashMLA code.
NOPE_FP8_VALUES = 512
GROUPS = 4
FP32_BYTES = 4
ROPE_BF16_VALUES = 64
BF16_BYTES = 2

V3_FP8_SPARSE_BYTES = (
NOPE_FP8_VALUES + GROUPS * FP32_BYTES + ROPE_BF16_VALUES * BF16_BYTES
)
assert V3_FP8_SPARSE_BYTES == 656
15 changes: 15 additions & 0 deletions artifacts/blogs/flashmla/code/02-decode-page-index-round-trip.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,15 @@
# Extracted from sources/blogs/flashmla.md by scripts/extract_blog_code.py
# Heading: ## Sparse Index Contracts > ### Decode page-index round trip
# Original fence language: python
# See artifacts/blogs/flashmla/code/PROVENANCE.yaml for origin + license metadata.

# KernelWiki-derived contract check; not upstream FlashMLA code.
def encode_page_index(physical_page: int, offset: int, page_size: int) -> int:
assert physical_page >= 0 and 0 <= offset < page_size
return physical_page * page_size + offset

def decode_page_index(encoded: int, page_size: int) -> tuple[int, int]:
assert encoded >= 0 and page_size > 0
return divmod(encoded, page_size)

assert decode_page_index(encode_page_index(7, 13, 64), 64) == (7, 13)

This file was deleted.

17 changes: 9 additions & 8 deletions artifacts/blogs/flashmla/code/PROVENANCE.yaml
Original file line number Diff line number Diff line change
@@ -1,21 +1,22 @@
origin_url: https://github.com/deepseek-ai/FlashMLA
origin_url: https://github.com/deepseek-ai/FlashMLA/tree/71c737929f2567bd0a094ae140f8f60f390b1232
upstream_repo: blog
upstream_sha: 71c737929f2567bd0a094ae140f8f60f390b1232
upstream_sha: none
license: inherits-from-source-blog
retrieved_at: 2026-04-27
asset_mode: extracted
generated_by: scripts/extract_blog_code.py
size_cap_truncated: false
files:
- local_path: 01-mla-decode-inner-loop.cu
- local_path: 01-v3-fp8-sparse-decode-byte-check.py
role: extracted-block
mode: extracted
upstream_path: inline-in-blog-markdown
heading_path: '## Key Code > ### MLA decode inner loop'
sha256: 4d3950b24b3da7a36e1db9d60a0492ed963d1b49ad6a1b97adeab54abaf4548f
- local_path: 02-sparse-mla-kv-retrieval-kernel-v3-2.cu
heading_path: '## Exact V3 FP8 Sparse-Decode Layout > ### V3 FP8 sparse-decode byte
check'
sha256: bfc0a6df6b2e04ce4f894a9df04926d4120b4e370354d68e80a55c1645565c7e
- local_path: 02-decode-page-index-round-trip.py
role: extracted-block
mode: extracted
upstream_path: inline-in-blog-markdown
heading_path: '## Key Code > ### Sparse-MLA KV-retrieval kernel (V3.2)'
sha256: 9b0499c0ddb77a9cebd370e84929d7149ebdc03163768369c95a12ae0b5c4f7f
heading_path: '## Sparse Index Contracts > ### Decode page-index round trip'
sha256: a0bbc6f0248adc67b8113e3faa66c1be8924aafce665ebf89b43f0e2f0e4a405
16 changes: 3 additions & 13 deletions artifacts/blogs/gated-delta-net/MANIFEST.yaml
Original file line number Diff line number Diff line change
@@ -1,14 +1,4 @@
slug: gated-delta-net
origin_url: https://github.com/NVlabs/GatedDeltaNet
code_present: true
total_blocks: 2
generated_by: scripts/extract_blog_code.py
files:
- local_path: code/01-chunk-parallel-prefill-reference-pytorch.py
heading_path: '## Key Code > ### Chunk-parallel prefill reference (PyTorch)'
fence_lang: python
sha256: 709cd74d72c22e764ffb77d85bbecd0b1c47b52f2cc0c48bab279954d0f24a22
- local_path: code/02-triton-decode-step-kernel-streaming.py
heading_path: '## Key Code > ### Triton decode-step kernel (streaming)'
fence_lang: python
sha256: 09a1bc93abdf8165a3d7809e485ef17a3d1a980ae579e36ba81b47665ba35b22
origin_url: https://github.com/NVlabs/GatedDeltaNet/tree/b53d6d3a161267432a79c1c04af69fa52bddc921
code_present: false
generated_by: KernelWiki verifier remediation

This file was deleted.

This file was deleted.

21 changes: 0 additions & 21 deletions artifacts/blogs/gated-delta-net/code/PROVENANCE.yaml

This file was deleted.

4 changes: 4 additions & 0 deletions artifacts/kernels/flash-attention-4/full/PROVENANCE.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@ retrieved_at: '2026-04-17'
asset_mode: verbatim
size_cap_truncated: false
generated_by: Phase 3 Round 2 task11 anchor bundle
description: >-
Adjacent NVIDIA CUTLASS SM100 FMHA backward MLA reference. This bundle is
not the Dao-AILab FlashAttention-4 implementation; it is retained only as a
separately pinned comparison implementation.
source_pr_id: pr-cutlass-2466
files:
- local_path: sm100_fmha_bwd_mla_kernel_tma_warpspecialized.hpp
Expand Down
Loading