Repository Analysis

Dao-AILab/flash-attention

Fast and memory-efficient exact attention

13.8 Low AI signal View on GitHub

Analysis Overview

This report presents the forensic synthetic code analysis of Dao-AILab/flash-attention, a Python project with 24,447 GitHub stars. SynthScan v2.0 examined 144,172 lines of code across 429 source files, recording 1238 pattern matches distributed across 18 syntactic categories. The overall adjusted score of 13.8 places this repository in the Low AI signal band.

The scanner applied 160+ deterministic lexical heuristics, multi-line block detectors, abstract syntax tree depth profilers, and a cross-file Jaccard similarity matrix to construct a statistically normalised synthetic code estimate. All matches are individually weighted by severity coefficient and contextual multiplier before summation, and the resulting headline score is temporally discounted to account for the repository's development history relative to the commercial emergence of large language model coding tooling (November 2022 onward).

13.8
Adjusted Score
13.8
Raw Score
100%
Time Factor
2026-07-14
Last Push
24.4K
Stars
Python
Language
144.2K
Lines of Code
429
Files
1.2K
Pattern Hits
2026-07-14
Scan Date
0.16
HC Hit Rate

What These Metrics Mean

Adjusted Score
Primary synthetic code indicator. Raw score normalised per 1,000 lines of code and multiplied by the temporal discount factor. This is the definitive comparative metric — use it to rank repositories by AI authorship density.
Raw Score
The unmodified sum of all severity-weighted, context-multiplied pattern match scores before temporal discounting. Reflects the absolute signal strength independent of when the repository was last active.
Time Factor
The temporal discount multiplier (0–100%) applied to the raw score. Repositories last updated before ChatGPT's launch (Nov 2022) receive a 5% factor. Full signal is only assigned to repositories active in the post-adoption era (Jan 2024+).
Pattern Hits
Total count of individual pattern matches across all files and categories. A high hit count with a low score may indicate a very large codebase with isolated AI snippets; a low count with a high score indicates dense, concentrated AI signatures.
HC Hit Rate
High+Critical pattern hits per file, averaged across the repository. This orthogonal signal catches repositories where a few files are densely packed with high-severity AI tells — a strong indicator even when the normalised score appears moderate due to codebase size.
Lines of Code / Files
Total lines and files analysed. The scanner examines 94 file extensions. These denominators are used to normalise the score, enabling fair comparison between repositories of vastly different sizes.

Score History

This chart maps the temporal evolution of the adjusted synthetic code score across successive scan runs. An upward trajectory indicates ongoing incorporation of AI-generated code or expanding LLM-assisted scaffolding; a stable or declining trajectory may reflect active human refactoring, code removal, or the adoption of stricter authorship policies. The dashed secondary line (right axis) independently tracks total raw pattern hit count, which can diverge from the normalised score when codebase size changes significantly between scans.

Severity Breakdown

Classifies detected patterns by their diagnostic confidence and structural impact. CRITICAL patterns (coefficient 10) represent definitive synthetic signatures — hallucinated imports, explicit LLM attribution metadata — virtually never produced by human authors. HIGH (5) indicates strong structural tells such as cross-file repetition or cross-linguistic idioms. MEDIUM (2) covers recognisable conversational padding and AI-specific vocabulary. LOW (1) captures subtle indicators like tautological comments and generic boilerplate that require density to carry independent signal.

CRITICAL 10HIGH 57MEDIUM 170LOW 1001

Directory Score Breakdown

This horizontal bar chart decomposes the repository's raw synthetic code score by top-level directory, allowing you to pinpoint precisely which modules or components carry the highest AI authorship density. Directories with disproportionately high scores relative to their size warrant targeted manual review: concentrated AI signatures often trace back to mass-generated configuration layers, auto-ported test suites, LLM-scaffolded boilerplate classes, or entire subsystems authored under heavy copilot assistance. Use this view to prioritise your human code-review effort.

Pattern Findings

The scanner identified 1238 distinct pattern matches across 18 syntactic categories. Each entry below represents a discrete location in the source code where the engine recorded a statistically significant AI authorship indicator. Expand any category row to inspect the individual file paths, line numbers, code snippets, and the lexical context (CODE, COMMENT, or STRING) in which each match was detected.

Reading the findings table: The Severity column indicates the diagnostic confidence level (CRITICAL / HIGH / MEDIUM / LOW). The Context column identifies whether the match occurred inside executable code, an inline comment, or a string literal — comment-context matches receive a ×1.5 weight because LLMs systematically over-annotate. The ⚡ bolt icon marks clustered matches: three or more patterns within a 10-line window, each receiving an additional ×1.5 density multiplier as dense clusters constitute far stronger evidence of synthetic authorship than isolated hits.

Decorative Section Separators152 hits · 513 pts
SeverityFileLineSnippetContext
MEDIUMtools/sass_diff.py24# ── Parsing ──────────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/sass_diff.py92# ── Diffing ──────────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/sass_diff.py111# ── Display ──────────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/sass_diff.py219# ── Main ─────────────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/ci/run_fa4_ci.py28# ── GPU helpers ───────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/ci/run_fa4_ci.py61# ── Runtime DSL pin (decouples cutlass-dsl from the baked image) ─────────────────COMMENT
MEDIUMtools/ci/run_fa4_ci.py86# ── Step plan ─────────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/ci/run_fa4_ci.py129# ── Step runner ───────────────────────────────────────────────────────────────COMMENT
MEDIUMtools/ci/run_fa4_ci.py203# ── CLI ───────────────────────────────────────────────────────────────────────COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py494# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py496# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py56# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py58# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py95# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py97# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py244# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py246# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py421# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py423# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py619# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py623# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py755# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod_varlen.py757# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py70# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py72# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py178# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py180# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py401# =============================================================================COMMENT
MEDIUMtests/cute/test_score_mod_varlen.py403# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py7# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py10# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py485# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py487# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py197# =============================================================================COMMENT
MEDIUMtests/cute/score_mod_definitions.py201# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod.py857# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod.py861# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod.py1533# =============================================================================COMMENT
MEDIUMtests/cute/test_mask_mod.py1535# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py14# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py17# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py19# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py21# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py270# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py272# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py333# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py337# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py176# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py180# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py388# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py390# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py492# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py494# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py596# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py598# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py748# =============================================================================COMMENT
MEDIUMtests/cute/mask_mod_definitions.py750# =============================================================================COMMENT
MEDIUMtests/cute/test_flash_attn.py2830# ---------------------------------------------------------------------------COMMENT
MEDIUMtests/cute/test_flash_attn.py2832# ---------------------------------------------------------------------------COMMENT
MEDIUMtests/cute/test_flash_attn.py2892# ---------------------------------------------------------------------------COMMENT
92 more matches not shown…
Hyper-Verbose Identifiers353 hits · 339 pts
SeverityFileLineSnippetContext
LOWsetup.py92def get_cuda_bare_metal_version(cuda_dir):CODE
LOWsetup.py215def validate_and_update_archs(archs):CODE
LOWcsrc/layer_norm/setup.py16def get_cuda_bare_metal_version(cuda_dir):CODE
LOWcsrc/layer_norm/setup.py25def check_cuda_torch_binary_vs_bare_metal(cuda_dir):CODE
LOWcsrc/fused_dense_lib/setup.py10def get_cuda_bare_metal_version(cuda_dir):CODE
LOWcsrc/flash_attn/src/generate_kernels.py43def get_fwd_split_align_template() -> str:CODE
LOWhopper/test_attn_kvcache.py155def test_flash_attn_kvcache_nosplit(nheads_kv, gqa_ratio, num_requests, query_seqlen, context_seqlen, headdim, causal, gCODE
LOWhopper/test_attn_kvcache.py292def test_flash_attn_kvcache_output(nheads_kv, gqa_ratio, num_requests, query_seqlen, context_seqlen, headdim, causal, usCODE
LOWhopper/test_torch_compile_and_export.py61def test_compile_and_package_model():CODE
LOWhopper/test_flash_attn_triton_amd.py334def test_flash_attn_varlen_output(CODE
LOWhopper/test_flash_attn_triton_amd.py1042def test_flash_attn_race_condition(seqlen_q, seqlen_k, d, causal, dtype):CODE
LOWhopper/test_util.py9def generate_random_padding_mask(max_seqlen, batch_size, device, mode="random", zero_lengths=False):CODE
LOWhopper/setup.py338def get_cuda_bare_metal_version(cuda_dir):STRING
LOWhopper/test_flash_attn_bwd_determinism.py391def test_flash_attn_varlen_output(CODE
LOWhopper/test_flash_attn.py404def test_flash_attn_varlen_output(CODE
LOWhopper/test_flash_attn.py1133def test_flash_attn_race_condition(seqlen_q, seqlen_k, d, causal, dtype):CODE
LOWhopper/flash_attn_interface.py313def _flash_attn_backward_fake(CODE
LOWhopper/flash_attn_interface.py747def flash_attn_qkvpacked_func(CODE
LOWtraining/src/metrics/num_tokens.py39 def _forward_reduce_state_update(self, *args: Any, **kwargs: Any) -> Any:CODE
LOWtraining/src/callbacks/speed_monitor.py35 def on_validation_epoch_start(self, trainer: "pl.Trainer", pl_module: "pl.LightningModule") -> None:CODE
LOWtraining/src/optim/param_grouping.py15def group_parameters_for_optimizer(model, optimizer_cfg, bias_weight_decay=False,CODE
LOWtraining/src/utils/gpu_affinity.py57def set_single_unique_affinity(gpu_id, nproc_per_node):CODE
LOWtraining/src/utils/gpu_affinity.py80def set_socket_unique_affinity(gpu_id, nproc_per_node, mode):CODE
LOWtraining/src/utils/checkpoint.py32def blockdiag_to_dense_mlp_bert(state_dict):CODE
LOWtraining/src/utils/checkpoint.py41def interpolate_pos_embedding(state_dict, out_seqlen, pos_embedding_name='model.pos_encoder.pe', interleave=False):CODE
LOWtraining/src/utils/ddp_zero1.py24def get_zero_optimizer_state_dict_local(optimizer, global_rank):CODE
LOWtests/test_flash_attn_triton_amd.py44def attn_bias_from_alibi_slopes(CODE
LOWtests/test_flash_attn_triton_amd.py73def generate_random_padding_mask(max_seqlen, batch_size, device, mode="random"):CODE
LOWtests/test_flash_attn_triton_amd.py397def attention_blocksparse_ref(qkv, blockmask, attn_mask, dropout_p, dropout_mask):CODE
LOWtests/test_flash_attn_triton_amd.py601def test_flash_attn_qkvpacked(seqlen, d, dropout_p, causal, local, alibi, deterministic, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py748def test_flash_attn_varlen_qkvpacked(CODE
LOWtests/test_flash_attn_triton_amd.py1191def test_flash_attn_varlen_output(CODE
LOWtests/test_flash_attn_triton_amd.py1619def test_flash_attn_varlen_causal(CODE
LOWtests/test_flash_attn_triton_amd.py2230def test_flash_attn_race_condition(seqlen_q, seqlen_k, d, dropout_p, causal, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py2279def test_flash_attn_bwd_overflow(seqlen, d, causal, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py2336def test_flash_attn_bwd_transpose(seqlen, d, causal, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py2389def test_flash_attn_bwd_varlen_overflow(d, causal, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py2448def test_flash_attn_deterministic(seqlen_q, seqlen_k, swap_sq_sk, d, causal, local, dtype):CODE
LOWtests/test_flash_attn_triton_amd.py2507def test_flash_attn_varlen_deterministic(seqlen_q, seqlen_k, swap_sq_sk, d, causal, local, dtype):CODE
LOWtests/test_rotary.py229def test_rotary_emb_varlen_func(inplace, interleaved, rotary_fraction, seqlen_offsets_type, dtype):CODE
LOWtests/test_flash_attn_ck.py42def get_bwd_unsupported_reason(d, deterministic):CODE
LOWtests/test_flash_attn_ck.py49def ck_randval_to_dropout_mask(randval, p):CODE
LOWtests/test_flash_attn_ck.py56def pad_rearrange_dropout_mask_hts_to_bhss(S_dmask, cu_seqlens_q, seqlen_q_rounded, seqlen_k_rounded):CODE
LOWtests/test_flash_attn_ck.py88def test_flash_attn_qkvpacked(seqlen, d, dropout_p, causal, local, alibi, deterministic, dtype):CODE
LOWtests/test_flash_attn_ck.py186def test_flash_attn_varlen_qkvpacked(seqlen, d, dropout_p, causal, local, alibi, deterministic, dtype):CODE
LOWtests/test_flash_attn_ck.py537def test_flash_attn_varlen_output(CODE
LOWtests/test_flash_attn_ck.py895def test_flash_attn_varlen_causal(CODE
LOWtests/test_flash_attn_ck.py1327def test_flash_attn_race_condition(seqlen_q, seqlen_k, d, dropout_p, causal, dtype):CODE
LOWtests/test_flash_attn_ck.py1374def test_flash_attn_bwd_overflow(seqlen, d, causal, dtype):CODE
LOWtests/test_flash_attn_ck.py1433def test_flash_attn_bwd_transpose(seqlen, d, causal, dtype):CODE
LOWtests/test_flash_attn_ck.py1486def test_flash_attn_bwd_varlen_overflow(d, causal, dtype):CODE
LOWtests/test_flash_attn_ck.py1522def test_flash_attn_bwd_varlen_seqq_zero(d, causal, nheads_kv, deterministic, dtype):CODE
LOWtests/test_flash_attn_ck.py1583def test_flash_attn_deterministic(seqlen_q, seqlen_k, swap_sq_sk, d, causal, local, dtype):CODE
LOWtests/test_flash_attn_ck.py1634def test_flash_attn_varlen_deterministic(seqlen_q, seqlen_k, swap_sq_sk, d, causal, local, dtype):CODE
LOWtests/test_util.py8def generate_random_padding_mask(max_seqlen, batch_size, device, mode="random", zero_lengths=False):CODE
LOWtests/test_flash_attn.py29def attn_bias_from_alibi_slopes(CODE
LOWtests/test_flash_attn.py58def generate_random_padding_mask(max_seqlen, batch_size, device, mode="random"):CODE
LOWtests/test_flash_attn.py382def attention_blocksparse_ref(qkv, blockmask, attn_mask, dropout_p, dropout_mask):CODE
LOWtests/test_flash_attn.py586def test_flash_attn_qkvpacked(seqlen, d, dropout_p, causal, local, alibi, deterministic, dtype):CODE
LOWtests/test_flash_attn.py733def test_flash_attn_varlen_qkvpacked(CODE
293 more matches not shown…
Cross-File Repetition48 hits · 240 pts
SeverityFileLineSnippetContext
HIGHREADME.md0if k and v are not none, k_cache and v_cache will be updated *inplace* with the new values from k and v. this is useful STRING
HIGHhopper/flash_attn_interface.py0if k and v are not none, k_cache and v_cache will be updated *inplace* with the new values from k and v. this is useful STRING
HIGHflash_attn/flash_attn_interface.py0if k and v are not none, k_cache and v_cache will be updated *inplace* with the new values from k and v. this is useful STRING
HIGHhopper/test_attn_kvcache.py0arguments: q: (batch_size, seqlen_q, nheads, head_dim) k: (batch_size, seqlen_k, nheads_k, head_dim) v: (batch_size, seqSTRING
HIGHtests/test_flash_attn_triton_amd.py0arguments: q: (batch_size, seqlen_q, nheads, head_dim) k: (batch_size, seqlen_k, nheads_k, head_dim) v: (batch_size, seqSTRING
HIGHtests/test_util.py0arguments: q: (batch_size, seqlen_q, nheads, head_dim) k: (batch_size, seqlen_k, nheads_k, head_dim) v: (batch_size, seqSTRING
HIGHtests/test_flash_attn.py0arguments: q: (batch_size, seqlen_q, nheads, head_dim) k: (batch_size, seqlen_k, nheads_k, head_dim) v: (batch_size, seqSTRING
HIGHhopper/test_kvcache.py0use pytorch benchmark on the forward pass of an arbitrary function.STRING
HIGHbenchmarks/benchmark_gemm.py0use pytorch benchmark on the forward pass of an arbitrary function.STRING
HIGHflash_attn/cute/benchmark.py0use pytorch benchmark on the forward pass of an arbitrary function.STRING
HIGHflash_attn/utils/benchmark.py0use pytorch benchmark on the forward pass of an arbitrary function.STRING
HIGHhopper/benchmark_flash_attention_fp8.py0arguments: qkv: (batch_size, seqlen, 3, nheads, head_dim) dropout_p: float output: output: (batch_size, seqlen, nheads, STRING
HIGHbenchmarks/benchmark_causal.py0arguments: qkv: (batch_size, seqlen, 3, nheads, head_dim) dropout_p: float output: output: (batch_size, seqlen, nheads, STRING
HIGHbenchmarks/benchmark_flash_attention.py0arguments: qkv: (batch_size, seqlen, 3, nheads, head_dim) dropout_p: float output: output: (batch_size, seqlen, nheads, STRING
HIGHhopper/test_util.py0arguments: q: (batch_size, seqlen_q, nheads, d) k: (batch_size, seqlen_k, nheads_k, d) v: (batch_size, seqlen_k, nheads_STRING
HIGHflash_attn/utils/testing.py0arguments: q: (batch_size, seqlen_q, nheads, d) k: (batch_size, seqlen_k, nheads_k, d) v: (batch_size, seqlen_k, nheads_STRING
HIGHtests/test_flash_attn_triton_amd.py0arguments: q: (batch_size, seqlen_q, nheads, d) k: (batch_size, seqlen_k, nheads_k, d) v: (batch_size, seqlen_k, nheads_STRING
HIGHtests/test_util.py0arguments: q: (batch_size, seqlen_q, nheads, d) k: (batch_size, seqlen_k, nheads_k, d) v: (batch_size, seqlen_k, nheads_STRING
HIGHtests/test_flash_attn.py0arguments: q: (batch_size, seqlen_q, nheads, d) k: (batch_size, seqlen_k, nheads_k, d) v: (batch_size, seqlen_k, nheads_STRING
HIGHtests/test_flash_attn_triton_amd.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/test_flash_attn_ck.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/test_flash_attn.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/test_flash_attn_triton_amd.py0we previously had a bug where we were using the wrong strides of dout, which shows up when dout is not contiguous.STRING
HIGHtests/test_flash_attn_ck.py0we previously had a bug where we were using the wrong strides of dout, which shows up when dout is not contiguous.STRING
HIGHtests/test_flash_attn.py0we previously had a bug where we were using the wrong strides of dout, which shows up when dout is not contiguous.STRING
HIGHtests/test_flash_attn_triton_amd.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/test_flash_attn_ck.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/test_flash_attn.py0we previously had a bug where not masking elements beyond seqlen_k caused nan in dq, in the case where seqlen % 128 != 0STRING
HIGHtests/models/test_llama.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_bigcode.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_opt.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_falcon.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_gptj.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_btlm.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_baichuan.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_gpt_neox.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_gpt.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_bert.py0check that our implementation of bert (without any optimizations enabled) matches the hf implementation: the output of oSTRING
HIGHtests/models/test_llama.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHtests/models/test_falcon.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHtests/models/test_baichuan.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHtests/models/test_opt.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHtests/models/test_gpt_generation_parallel.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHtests/models/test_gpt.py0check that our implementation of gpt2 generation matches the hf implementation: the scores in fp16 should be around the STRING
HIGHflash_attn/cute/flash_bwd.py0check if the kernel can be implemented with the given parameters. :param dtype: data type :type dtype: cutlass.numeric :STRING
HIGHflash_attn/cute/flash_bwd_postprocess.py0check if the kernel can be implemented with the given parameters. :param dtype: data type :type dtype: cutlass.numeric :STRING
HIGHflash_attn/cute/flash_fwd.py0check if the kernel can be implemented with the given parameters. :param dtype: data type :type dtype: cutlass.numeric :STRING
HIGHflash_attn/cute/flash_bwd_preprocess.py0check if the kernel can be implemented with the given parameters. :param dtype: data type :type dtype: cutlass.numeric :STRING
Unused Imports199 hits · 197 pts
SeverityFileLineSnippetContext
LOWsetup.py25CODE
LOWcsrc/layer_norm/setup.py2CODE
LOWcsrc/layer_norm/setup.py3CODE
LOWcsrc/layer_norm/setup.py8CODE
LOWcsrc/layer_norm/setup.py9CODE
LOWcsrc/fused_dense_lib/setup.py5CODE
LOWtools/ci/run_fa4_ci.py7CODE
LOWtools/ci/assert_dsl_floor.py13CODE
LOWhopper/benchmark_mla_decode.py11CODE
LOWhopper/test_attn_kvcache.py4CODE
LOWhopper/test_attn_kvcache.py6CODE
LOWhopper/test_attn_kvcache.py8CODE
LOWhopper/test_kvcache.py9CODE
LOWhopper/benchmark_flash_attention_fp8.py3CODE
LOWhopper/benchmark_flash_attention_fp8.py7CODE
LOWhopper/benchmark_flash_attention_fp8.py10CODE
LOWhopper/benchmark_flash_attention_fp8.py12CODE
LOWhopper/benchmark_flash_attention_fp8.py12CODE
LOWhopper/benchmark_flash_attention_fp8.py13CODE
LOWhopper/benchmark_flash_attention_fp8.py13CODE
LOWhopper/benchmark_flash_attention_fp8.py15CODE
LOWhopper/benchmark_flash_attention_fp8.py16CODE
LOWhopper/test_flash_attn_triton_amd.py7CODE
LOWhopper/benchmark_attn.py1CODE
LOWhopper/benchmark_attn.py2CODE
LOWhopper/benchmark_attn.py7CODE
LOWhopper/benchmark_attn.py8CODE
LOWhopper/benchmark_attn.py21CODE
LOWhopper/benchmark_attn.py24CODE
LOWhopper/benchmark_attn.py24CODE
LOWhopper/benchmark_attn.py24CODE
LOWhopper/benchmark_attn.py24CODE
LOWhopper/benchmark_attn.py24CODE
LOWhopper/setup.py25CODE
LOWhopper/test_flash_attn_bwd_determinism.py2CODE
LOWhopper/test_flash_attn_bwd_determinism.py7CODE
LOWhopper/test_flash_attn_bwd_determinism.py8CODE
LOWhopper/test_flash_attn_bwd_determinism.py10CODE
LOWhopper/test_flash_attn_bwd_determinism.py16CODE
LOWhopper/test_flash_attn_bwd_determinism.py16CODE
LOWhopper/test_flash_attn_bwd_determinism.py23CODE
LOWhopper/test_flash_attn_bwd_determinism.py24CODE
LOWhopper/test_flash_attn_bwd_determinism.py24CODE
LOWhopper/test_flash_attn.py7CODE
LOWhopper/flash_attn_interface.py3CODE
LOWhopper/flash_attn_interface.py7CODE
LOWhopper/benchmark_split_kv.py5CODE
LOWhopper/flash_attn_3/flash_attn_config.py1CODE
LOWhopper/flash_attn_3/flash_attn_interface.py2CODE
LOWtraining/tests/datamodules/test_language_modeling_hf.py6CODE
LOWtraining/src/eval.py1CODE
LOWtraining/src/eval.py8CODE
LOWtraining/src/metrics/perplexity.py9CODE
LOWtraining/src/metrics/accuracy.py1CODE
LOWtraining/src/metrics/accuracy.py4CODE
LOWtraining/src/tasks/seq.py1CODE
LOWtraining/src/tasks/seq.py4CODE
LOWtraining/src/distributed/ddp_comm_hooks.py3CODE
LOWtraining/src/distributed/ddp_comm_hooks.py3CODE
LOWtraining/src/callbacks/flop_count.py2CODE
139 more matches not shown…
Hallucination Indicators10 hits · 130 pts
SeverityFileLineSnippetContext
CRITICALhopper/test_flash_attn_triton_amd.py1140 assert torch.ops.flash_attn_3.fwd.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn_triton_amd.py1153 assert torch.ops.flash_attn_3.bwd.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn_triton_amd.py1161 assert torch.ops.flash_attn_3.fwd_combine.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn_triton_amd.py1165 assert torch.ops.flash_attn_3.get_scheduler_metadata.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn.py1230 assert torch.ops.flash_attn_3.fwd.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn.py1243 assert torch.ops.flash_attn_3.bwd.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn.py1251 assert torch.ops.flash_attn_3.fwd_combine.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALhopper/test_flash_attn.py1255 assert torch.ops.flash_attn_3.get_scheduler_metadata.default._schema.is_backward_compatible_with(parse_schema(CODE
CRITICALtests/models/test_btlm.py214 assert model.transformer.embeddings.word_embeddings.weight.mean().abs() < 1e-4CODE
CRITICALtests/models/test_btlm.py216 model.transformer.embeddings.word_embeddings.weight.std()CODE
Deep Nesting127 hits · 118 pts
SeverityFileLineSnippetContext
LOWsetup.py101CODE
LOWtools/sass_diff.py128CODE
LOWhopper/benchmark_flash_attention_fp8.py34CODE
LOWhopper/test_flash_attn_triton_amd.py628CODE
LOWhopper/benchmark_attn.py76CODE
LOWhopper/setup.py138CODE
LOWhopper/test_flash_attn_bwd_determinism.py110CODE
LOWhopper/test_flash_attn_bwd_determinism.py391CODE
LOWhopper/test_flash_attn.py715CODE
LOWhopper/benchmark_split_kv.py35CODE
LOWtraining/src/train.py32CODE
LOWtraining/src/callbacks/norm_monitor.py33CODE
LOWtraining/src/optim/param_grouping.py15CODE
LOWtraining/src/utils/gpu_affinity.py80CODE
LOWtraining/src/utils/gpu_affinity.py127CODE
LOWtraining/src/utils/ema.py228CODE
LOWtraining/src/utils/distributed.py70CODE
LOWtraining/src/models/modules/seq_common.py15CODE
LOWtests/cute/benchmark_mask_mod.py154CODE
LOWtests/cute/benchmark_mask_mod.py448CODE
LOWtests/cute/test_mask_mod_varlen.py249CODE
LOWtests/cute/test_mask_mod_varlen.py903CODE
LOWtests/cute/test_flash_attn_race_condition.py391CODE
LOWtests/cute/test_score_mod_varlen.py602CODE
LOWtests/cute/test_score_mod_varlen.py950CODE
LOWtests/cute/test_mask_mod.py2193CODE
LOWtests/cute/test_mask_mod.py2239CODE
LOWtests/cute/test_flash_attn.py159CODE
LOWtests/cute/test_flash_attn.py621CODE
LOWtests/cute/test_flash_attn.py1128CODE
LOWtests/cute/test_flash_attn.py2366CODE
LOWtests/cute/test_flash_attn.py2755CODE
LOWtests/cute/test_block_sparsity.py43CODE
LOWtests/cute/test_block_sparsity.py484CODE
LOWbenchmarks/tune_ex2_emu.py33CODE
LOWbenchmarks/tune_ex2_emu.py225CODE
LOWbenchmarks/benchmark_attn.py363CODE
LOWbenchmarks/bench_sm90.py334CODE
LOWbenchmarks/bench_sm90.py367CODE
LOWbenchmarks/bench_sm90.py397CODE
LOWbenchmarks/bench_sm90.py452CODE
LOWbenchmarks/bench_sm90.py489CODE
LOWflash_attn/flash_attn_triton.py66CODE
LOWflash_attn/flash_attn_triton.py365CODE
LOWflash_attn/flash_attn_interface.py31CODE
LOWflash_attn/cute/sm90_config_search.py174CODE
LOWflash_attn/cute/sm90_config_search.py315CODE
LOWflash_attn/cute/mask.py76CODE
LOWflash_attn/cute/mask.py177CODE
LOWflash_attn/cute/mask.py497CODE
LOWflash_attn/cute/mask.py615CODE
LOWflash_attn/cute/mask.py777CODE
LOWflash_attn/cute/mask.py1444CODE
LOWflash_attn/cute/mask.py1586CODE
LOWflash_attn/cute/mask.py1661CODE
LOWflash_attn/cute/sm100_hd256_2cta_fmha_forward.py589CODE
LOWflash_attn/cute/sm100_hd256_2cta_fmha_forward.py1560CODE
LOWflash_attn/cute/flash_bwd_mla_dq_dqv_sm100.py437CODE
LOWflash_attn/cute/ampere_helpers.py35CODE
LOWflash_attn/cute/flash_bwd_mla_sm100.py1499CODE
67 more matches not shown…
Over-Commented Block108 hits · 104 pts
SeverityFileLineSnippetContext
LOWcsrc/flash_attn_ck/flash_common.hpp1/******************************************************************************COMMENT
LOWcsrc/layer_norm/static_switch.h1// Inspired by https://github.com/NVIDIA/DALI/blob/main/include/dali/core/static_switch.hCOMMENT
LOWcsrc/fused_dense_lib/fused_dense.cpp1// Adapted from https://github.com/NVIDIA/apex/blob/master/csrc/fused_dense.cppCOMMENT
LOWcsrc/flash_attn/flash_api.cpp1/******************************************************************************COMMENT
LOWcsrc/flash_attn/src/flash_fwd_kernel.h1/******************************************************************************COMMENT
LOWcsrc/flash_attn/src/flash_fwd_kernel.h221 // for (int i = 0; i < size(tScQ); ++i) {COMMENT
LOWcsrc/flash_attn/src/utils.h1/******************************************************************************COMMENT
LOWcsrc/flash_attn/src/utils.h321 cute::clear(D(_, m, _));COMMENT
LOWcsrc/flash_attn/src/utils.h341 // if (Is_even_MN || get<0>(identity_MN(0, m, 0)) < max_MN) {COMMENT
LOWcsrc/flash_attn/src/flash_fwd_launch_template.h1/******************************************************************************COMMENT
LOWcsrc/flash_attn/src/flash_fwd_launch_template.h261 } else {COMMENT
LOWcsrc/flash_attn/src/flash_bwd_launch_template.h1/******************************************************************************COMMENT
LOWcsrc/flash_attn/src/flash_bwd_launch_template.h21#define KERNEL_PARAM_MODIFIER __grid_constant__COMMENT
LOWcsrc/flash_attn/src/flash_bwd_launch_template.h181 // run_flash_bwd<Flash_bwd_kernel_traits<Headdim, 128, 128, 8, 4, 4, 4, true, false, T>, Is_dropout>(params,COMMENT
LOWcsrc/flash_attn/src/flash_bwd_launch_template.h241 // printf("max_smem_per_block = %d\n", max_smem_per_block);COMMENT
LOWcsrc/flash_attn/src/dropout.h41 uint2 rowcol = make_uint2(block_row_start, block_col_start);COMMENT
LOWcsrc/flash_attn/src/flash_bwd_kernel.h1/***************************************************************************************************COMMENT
LOWcsrc/flash_attn/src/flash_bwd_kernel.h321 // If not local, we're guaranteed that m_block_min <= m_block:COMMENT
LOWcsrc/flash_attn/src/static_switch.h1// Inspired byCOMMENT
LOWtools/ci/build_sif.sh1#!/usr/bin/env bashCOMMENT
LOWhopper/mainloop_fwd_sm90_tma_gmma_ws.hpp1/******************************************************************************COMMENT
LOWhopper/utils.h1/******************************************************************************COMMENT
LOWhopper/mainloop_bwd_sm90_tma_gmma_ws.hpp1/******************************************************************************COMMENT
LOWhopper/benchmark_mla_decode.py121 print(f"Arithmetic intensity: {flops / mem_io:.1f}")COMMENT
LOWhopper/mainloop_bwd_sm80.hpp1/******************************************************************************COMMENT
LOWhopper/mainloop_bwd_sm80.hpp561 #pragma unrollCOMMENT
LOWhopper/mainloop_bwd_sm80.hpp621 // Instead of passing in tQcQ, we pass in t0QcQ and subtract the offset from the limitCOMMENT
LOWhopper/flash_fwd_launch_template.h1/******************************************************************************COMMENT
LOWhopper/benchmark_flash_attention_fp8.py321 # )()COMMENT
LOWhopper/test_flash_attn_triton_amd.py61@pytest.mark.parametrize("deterministic", [False])COMMENT
LOWhopper/test_flash_attn_triton_amd.py221 and dtype != torch.float8_e4m3fn COMMENT
LOWhopper/test_flash_attn_triton_amd.py241 # causal,COMMENT
LOWhopper/test_flash_attn_triton_amd.py501 ):COMMENT
LOWhopper/test_flash_attn_triton_amd.py581# @pytest.mark.parametrize("mha_type", ["mha"])COMMENT
LOWhopper/test_flash_attn_triton_amd.py901 out = output_pad_fn(out)COMMENT
LOWhopper/tile_scheduler.hpp641 // Total number of blocks for the next 31 batchesCOMMENT
LOWhopper/tile_scheduler.hpp741 int split_idx = bidh - bidh_actual * num_splits;COMMENT
LOWhopper/benchmark_attn.py41def time_fwd(func, *args, repeats=30, verbose=True, desc="", **kwargs):COMMENT
LOWhopper/benchmark_attn.py241# bs_seqlen_vals = [(32, 512), (16, 1024)]COMMENT
LOWhopper/benchmark_attn.py401 # print(time_f)COMMENT
LOWhopper/flash_bwd_launch_template.h1/******************************************************************************COMMENT
LOWhopper/test_flash_attn_bwd_determinism.py61# @pytest.mark.parametrize("mha_type", ["mqa"])COMMENT
LOWhopper/test_flash_attn_bwd_determinism.py341# @pytest.mark.parametrize("dtype", [torch.float8_e4m3fn])COMMENT
LOWhopper/test_flash_attn_bwd_determinism.py401 # batch_size = 40COMMENT
LOWhopper/static_switch.h1// Inspired byCOMMENT
LOWhopper/mainloop_fwd_sm80.hpp1/******************************************************************************COMMENT
LOWhopper/test_flash_attn.py121# @pytest.mark.parametrize("has_qv", [True])COMMENT
LOWhopper/test_flash_attn.py301 # k,COMMENT
LOWhopper/test_flash_attn.py361@pytest.mark.parametrize("softcap", [0.0] + ([15.0] if not DISABLE_SOFTCAP else []))COMMENT
LOWhopper/test_flash_attn.py601 # None,COMMENT
LOWhopper/test_flash_attn.py621 dv.masked_fill_(k_zero_masking, 0.0)COMMENT
LOWhopper/flash_fwd_kernel_sm90.h1/******************************************************************************COMMENT
LOWhopper/flash_api.cpp1221 #ifndef FLASHATTENTION_DISABLE_HDIM256COMMENT
LOWhopper/benchmark_split_kv.py121 causal=causal,COMMENT
LOWhopper/flash_api_stable.cpp1/******************************************************************************COMMENT
LOWhopper/flash_api_stable.cpp541 #endifCOMMENT
LOWhopper/flash_api_stable.cpp1241 if (out_type == torch::headeronly::ScalarType::BFloat16) {COMMENT
LOWhopper/flash_api_stable.cpp1281 if (params.d_rounded == 64) { return run_mha_bwd_<Arch, cutlass::half_t, 64, Has_softcap>(params, stream); }COMMENT
LOWhopper/flash_api_stable.cpp1301 #endifCOMMENT
LOWtraining/configs/experiment/owt/gpt2xl-flash.yaml1# @package _global_COMMENT
48 more matches not shown…
AI Structural Patterns105 hits · 100 pts
SeverityFileLineSnippetContext
LOWhopper/test_attn_kvcache.py45CODE
LOWhopper/test_torch_compile_and_export.py36CODE
LOWhopper/test_util.py226CODE
LOWhopper/flash_attn_interface.py60CODE
LOWhopper/flash_attn_interface.py154CODE
LOWhopper/flash_attn_interface.py259CODE
LOWhopper/flash_attn_interface.py313CODE
LOWhopper/flash_attn_interface.py747CODE
LOWhopper/flash_attn_interface.py809CODE
LOWhopper/flash_attn_interface.py890CODE
LOWhopper/flash_attn_interface.py942CODE
LOWhopper/flash_attn_interface.py1106CODE
LOWhopper/flash_attn_interface.py455CODE
LOWhopper/flash_attn_interface.py555CODE
LOWhopper/flash_attn_interface.py645CODE
LOWtraining/src/datamodules/language_modeling_hf.py42CODE
LOWtraining/src/datamodules/imagenet.py63CODE
LOWtraining/src/utils/gpu_affinity.py42CODE
LOWtraining/src/utils/checkpoint.py66CODE
LOWtests/test_flash_attn_triton_amd.py232CODE
LOWtests/test_flash_attn_triton_amd.py322CODE
LOWtests/test_flash_attn_triton_amd.py355CODE
LOWtests/test_util.py185CODE
LOWtests/test_flash_attn.py217CODE
LOWtests/test_flash_attn.py307CODE
LOWtests/test_flash_attn.py340CODE
LOWtests/cute/test_flash_attn_varlen.py51CODE
LOWtests/cute/score_mod_definitions.py338CODE
LOWtests/cute/test_mask_mod.py1537CODE
LOWbenchmarks/clc_bench.py758CODE
LOWbenchmarks/bench_sm90.py107CODE
LOWflash_attn/flash_attn_interface.py154CODE
LOWflash_attn/flash_attn_interface.py207CODE
LOWflash_attn/flash_attn_interface.py1019CODE
LOWflash_attn/flash_attn_interface.py1078CODE
LOWflash_attn/flash_attn_interface.py1156CODE
LOWflash_attn/flash_attn_interface.py1233CODE
LOWflash_attn/flash_attn_interface.py1299CODE
LOWflash_attn/flash_attn_interface.py1391CODE
LOWflash_attn/flash_attn_interface.py1485CODE
LOWflash_attn/losses/cross_entropy.py10CODE
LOWflash_attn/cute/mask.py615CODE
LOWflash_attn/cute/mask.py1078CODE
LOWflash_attn/cute/mask.py1583CODE
LOWflash_attn/cute/sm100_hd256_2cta_fmha_forward.py36CODE
LOWflash_attn/cute/sm100_hd256_2cta_fmha_forward.py168CODE
LOWflash_attn/cute/seqlen_info.py84CODE
LOWflash_attn/cute/flash_bwd_mla_sm100.py44CODE
LOWflash_attn/cute/flash_bwd.py32CODE
LOWflash_attn/cute/flash_bwd.py374CODE
LOWflash_attn/cute/flash_fwd_mla_sm100.py49CODE
LOWflash_attn/cute/flash_fwd_mla_sm100.py350CODE
LOWflash_attn/cute/compute_block_sparsity.py334CODE
LOWflash_attn/cute/flash_bwd_sm90.py49CODE
LOWflash_attn/cute/flash_bwd_sm90.py344CODE
LOWflash_attn/cute/flash_bwd_sm90.py630CODE
LOWflash_attn/cute/flash_bwd_sm100.py52CODE
LOWflash_attn/cute/flash_bwd_sm100.py444CODE
LOWflash_attn/cute/interface.py299CODE
LOWflash_attn/cute/interface.py1217CODE
45 more matches not shown…
Structural Annotation Overuse34 hits · 58 pts
SeverityFileLineSnippetContext
LOWhopper/flash_fwd_combine_kernel.h229 // Step 1: load LSE_partial from gmem -> smemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h274 // Step 2: Load O_partial from gmem -> smem for split = 0, 1, ..., kStages - 2.COMMENT
LOWhopper/flash_fwd_combine_kernel.h335 // Step 3: load and transpose LSE_partial from smem -> rmemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h345 // Step 4: compute the final LSE along the split dimensionCOMMENT
LOWhopper/flash_fwd_combine_kernel.h394 // Step 5: store final LSE back to gmemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h417 // Step 6: read O_partial from gmem -> smem -> rmem and accumulate the final OCOMMENT
LOWhopper/flash_fwd_combine_kernel.h460 // Step 7: Write the final O to gmemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h174 // Step 1: load dQaccum from gmem to smemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h200 // Step 2: Load dQaccum from smem to register, then convert fp32 -> fp16/bf16COMMENT
LOWhopper/flash_bwd_postprocess_kernel.h218 // Step 3: Copy dQ from register to smemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h229 // Step 4: Copy dQ from smem to register to prepare for coalesced write to gmemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h247 // Step 5: Copy dQ from register to gmemCOMMENT
LOWhopper/epilogue_fwd.hpp251 // Step 1: Write O from rmem -> smemCOMMENT
LOWhopper/epilogue_fwd.hpp281 // Step 2: Write LSE from rmem -> gmemCOMMENT
LOWhopper/epilogue_fwd.hpp310 // Step 3: Write O from smem -> gmemCOMMENT
LOWAI/DEBUG_2CTA.md5### Step 1: Build a minimal reproCOMMENT
LOWAI/DEBUG_2CTA.md12### Step 2: Add printf to locate the hangCOMMENT
LOWAI/DEBUG_2CTA.md42### Step 3: Identify the deadlock chainCOMMENT
LOWAI/DEBUG_2CTA.md54### Step 4: Vary the problem size systematicallyCOMMENT
LOWAI/DEBUG_2CTA.md67### Step 5: Check barrier byte counts (tx_count)COMMENT
LOWAI/DEBUG_2CTA.md77### Step 6: Check phase / parity trackingCOMMENT
LOWAI/DEBUG_2CTA.md81### Step 7: Beware compiler-as-bug-sourceCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py492 # Step 1: load dQaccum from gmem to smemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py501 # Step 2: load dQ from smem to rmemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py534 # Step 3: Copy dQ from register to smemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py568 # Step 4: Copy dQ from smem to register to prepare for coalesced write to gmemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py577 # Step 5: Copy dQ from register to gmemCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py405 # Step 1: Load LSE_partial from gmem to shared memoryCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py442 # Step 2: Load O_partial for pipeline stagesCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py495 # Step 3: Load and transpose LSE from smem to registersCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py513 # Step 4: Compute final LSE along split dimensionCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py573 # Step 5: Store final LSE to gmemCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py595 # Step 6: Read O_partial and accumulate final OCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py642 # Step 7: Write final O to gmemCOMMENT
Verbosity Indicators32 hits · 55 pts
SeverityFileLineSnippetContext
LOWhopper/flash_fwd_combine_kernel.h229 // Step 1: load LSE_partial from gmem -> smemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h274 // Step 2: Load O_partial from gmem -> smem for split = 0, 1, ..., kStages - 2.COMMENT
LOWhopper/flash_fwd_combine_kernel.h335 // Step 3: load and transpose LSE_partial from smem -> rmemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h345 // Step 4: compute the final LSE along the split dimensionCOMMENT
LOWhopper/flash_fwd_combine_kernel.h394 // Step 5: store final LSE back to gmemCOMMENT
LOWhopper/flash_fwd_combine_kernel.h417 // Step 6: read O_partial from gmem -> smem -> rmem and accumulate the final OCOMMENT
LOWhopper/flash_fwd_combine_kernel.h460 // Step 7: Write the final O to gmemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h174 // Step 1: load dQaccum from gmem to smemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h200 // Step 2: Load dQaccum from smem to register, then convert fp32 -> fp16/bf16COMMENT
LOWhopper/flash_bwd_postprocess_kernel.h218 // Step 3: Copy dQ from register to smemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h229 // Step 4: Copy dQ from smem to register to prepare for coalesced write to gmemCOMMENT
LOWhopper/flash_bwd_postprocess_kernel.h247 // Step 5: Copy dQ from register to gmemCOMMENT
LOWhopper/epilogue_fwd.hpp251 // Step 1: Write O from rmem -> smemCOMMENT
LOWhopper/epilogue_fwd.hpp281 // Step 2: Write LSE from rmem -> gmemCOMMENT
LOWhopper/epilogue_fwd.hpp310 // Step 3: Write O from smem -> gmemCOMMENT
LOWflash_attn/cute/flash_bwd.py256 # Do we need to check if we overshot kBlockM when we load Q?COMMENT
LOWflash_attn/cute/flash_bwd.py258 # Do we need to check if we overshot kBlockN when we load K?COMMENT
LOWflash_attn/cute/flash_bwd.py265 # Do we need to check if we overshot kBlockN when we load V?COMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py492 # Step 1: load dQaccum from gmem to smemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py501 # Step 2: load dQ from smem to rmemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py534 # Step 3: Copy dQ from register to smemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py568 # Step 4: Copy dQ from smem to register to prepare for coalesced write to gmemCOMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py577 # Step 5: Copy dQ from register to gmemCOMMENT
LOWflash_attn/cute/flash_fwd.py496 # Do we need to check if we overshoot kBlockN when we load K?COMMENT
LOWflash_attn/cute/flash_fwd.py542 # Do we need to check if we overshoot kBlockN when we load V?COMMENT
LOWflash_attn/cute/flash_fwd_combine.py405 # Step 1: Load LSE_partial from gmem to shared memoryCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py442 # Step 2: Load O_partial for pipeline stagesCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py495 # Step 3: Load and transpose LSE from smem to registersCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py513 # Step 4: Compute final LSE along split dimensionCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py573 # Step 5: Store final LSE to gmemCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py595 # Step 6: Read O_partial and accumulate final OCOMMENT
LOWflash_attn/cute/flash_fwd_combine.py642 # Step 7: Write final O to gmemCOMMENT
Cross-Language Confusion9 hits · 42 pts
SeverityFileLineSnippetContext
HIGHtools/ci/run_fa4_ci.py165 '"$SP"/quack "$SP"/quack_kernels* 2>/dev/null || true'CODE
HIGHhopper/setup.py308 blocks.append(cuda_compile_rule) # type: ignore[possibly-undefined]STRING
HIGHhopper/setup.py309 blocks.append(cuda_compile_rule_sm80) # type: ignore[possibly-undefined]STRING
HIGHhopper/setup.py310 blocks.append(cuda_compile_rule_sm80_sm90) # type: ignore[possibly-undefined]STRING
HIGHhopper/setup.py311 blocks.append(cuda_compile_rule_sm100) # type: ignore[possibly-undefined]STRING
HIGHtests/cute/test_flash_attn_combine.py183 # Only compare valid positions (beyond seqused, output is undefined)COMMENT
HIGHAI/parse_clc_log.py246let selectedSm = null;CODE
HIGHAI/parse_clc_log.py309 if (id === query || id.includes(query)) {{CODE
HIGHAI/parse_clc_log.py326 selectedSm = null;CODE
Modern AI Meta-Vocabulary11 hits · 23 pts
SeverityFileLineSnippetContext
MEDIUMREADME.md297 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMhopper/flash_attn_interface.py964 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMhopper/flash_attn_interface.py1117 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMtraining/src/datamodules/language_modeling_hf.py118 # However, it's useful for zero-shot transfer from Openwebtext,COMMENT
MEDIUMflash_attn/flash_attn_interface.py1024 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1084 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1163 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1240 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1309 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1402 window_size=(-1, -1), # -1 means infinite context windowCODE
MEDIUMflash_attn/flash_attn_interface.py1499 window_size=(-1, -1), # -1 means infinite context windowCODE
AI Slop Vocabulary11 hits · 19 pts
SeverityFileLineSnippetContext
LOWhopper/generate_kernels.py134 # so we should just pass in packgqa=False to avoid the `_packgqa` in the filename.STRING
LOWtests/cute/score_mod_definitions.py478 # Don't read from aux_tensors at all - just add the global index as biasCOMMENT
MEDIUMtests/cute/test_mask_mod.py6# (identity, document, block_diagonal, etc.) with comprehensive seqlen coverageCOMMENT
LOWflash_attn/flash_attn_triton.py145 # [2022-10-30] TD: Triton bug - in the case of EVEN_M=True and EVEN_N=False, if we just callCOMMENT
LOWflash_attn/flash_attn_triton.py347 # if we just call tl.store(dv_ptrs), there's a race conditionCOMMENT
LOWflash_attn/flash_attn_triton.py442 # if we just call tl.load(k_ptrs), we get the wrong output!COMMENT
LOWflash_attn/flash_attn_triton.py521 # [2022-11-01] TD: Triton bug, there's a race condition if we just use m_mask and not d_mask.COMMENT
LOWflash_attn/cute/flash_bwd_postprocess.py187 # We can't just use kHeadDim here. E.g. if MMA shape is 64 x 96 but split across 2 WGs,COMMENT
MEDIUMflash_attn/cute/flash_fwd_sm100.py712 # CLC buffers placed here to utilize padding before sO's 1024-byte alignment.COMMENT
LOWflash_attn/cute/flash_fwd_sm90.py1392 # 2 elements. So we just call ptx directly.COMMENT
LOWflash_attn/cute/flash_fwd_sm90.py1464 # 2 elements. So we just call ptx directly.COMMENT
Excessive Try-Catch Wrapping17 hits · 18 pts
SeverityFileLineSnippetContext
LOWsetup.py186 except Exception as e:CODE
MEDIUMsetup.py180def detect_hipify_v2():CODE
LOWtests/cute/benchmark_block_sparsity.py83 except Exception as e:CODE
LOWtests/cute/benchmark_block_sparsity.py190 except Exception as e:CODE
LOWtests/cute/benchmark_block_sparsity.py375 except Exception as e:CODE
LOWbenchmarks/tune_ex2_emu.py307 except Exception as e:CODE
LOWbenchmarks/tune_ex2_emu.py370 except Exception as e:CODE
LOWbenchmarks/bench_sm90.py126 except Exception as e:CODE
LOWbenchmarks/bench_sm90.py165 except Exception as e:CODE
LOWbenchmarks/bench_sm90.py175 except Exception as e:CODE
LOWbenchmarks/benchmark_flash_attention.py119 except Exception:CODE
LOWbenchmarks/benchmark_flash_attention.py134 except Exception:CODE
LOWbenchmarks/benchmark_flash_attention.py141 except Exception:CODE
LOWflash_attn/cute/cute_dsl_ptxas.py93 except Exception as e:CODE
LOWflash_attn/cute/benchmark_flash_attention_fp8.py330 except Exception as e:CODE
LOWflash_attn/cute/benchmark_flash_attention_fp8.py402 except Exception as e:CODE
LOWflash_attn/cute/utils.py83 except Exception:CODE
Redundant / Tautological Comments9 hits · 12 pts
SeverityFileLineSnippetContext
LOWsetup.py522 # Check if torch is using hipify v2. Until CK is updated with HIPIFY_V2 macro,COMMENT
LOWhopper/setup.py401 # Set timeout to 300 seconds to prevent the request from hanging forever.STRING
LOWtests/cute/test_utils.py204 # Set __cute_hash__ to simulate Inductor-generated codeCOMMENT
LOWtests/cute/test_block_sparsity.py111 # Check if ref skipped it entirely (all masked)COMMENT
LOWflash_attn/cute/flash_bwd.py138 # Check if block size setting is out of shared memory capacityCOMMENT
LOWflash_attn/cute/compute_block_sparsity.py382 # Check if mask_mod is marked as suitable for 5-point samplingCOMMENT
LOWflash_attn/cute/interface.py2398 # Check if configuration can be implementedCOMMENT
LOWflash_attn/cute/flash_fwd.py159 # Check if block size setting is out of shared memory capacityCOMMENT
LOWflash_attn/cute/flash_fwd.py172 # Check if twice the block size is divisible by the number of threadsCOMMENT
Self-Referential Comments4 hits · 12 pts
SeverityFileLineSnippetContext
MEDIUMhopper/generate_kernels.py3# This file is run to generate the kernel instantiations for the flash_attn kernelsCOMMENT
MEDIUMtests/cute/test_flash_attn_combine.py254 # Create a permuted batch index mapping: virtual batch -> real batchCOMMENT
MEDIUMtests/models/test_llama.py578 # Create a shared test model.COMMENT
MEDIUMflash_attn/modules/embedding.py137 # Create a mask of valid vocab ids (1 means it needs to be masked).COMMENT
Modern Structural Boilerplate6 hits · 6 pts
SeverityFileLineSnippetContext
LOWtraining/src/metrics/perplexity.py18__all__ = ['Perplexity']CODE
LOW.github/scripts/bump_beta_tag.py46def set_github_output(key: str, value: str) -> None:CODE
LOWflash_attn/cute/__init__.py15__all__ = [CODE
LOWflash_attn/cute/fa_logging.py77def set_fa_log_level(level: int | str) -> None:CODE
LOWflash_attn/models/bert.py54logger = logging.getLogger(__name__)CODE
LOWflash_attn/models/gpt.py59logger = logging.getLogger(__name__)CODE
Example Usage Blocks3 hits · 4 pts
SeverityFileLineSnippetContext
LOWtools/ci/build_sif.sh4# Usage:COMMENT
LOWtests/cute/test_mask_mod_varlen.py8# Usage:COMMENT
LOWtests/cute/test_mask_mod.py10# Usage:COMMENT