nikil-ravi/ThunderKittens
Languages
Tile primitives for speedy kernels
Harness | Input / Output Cost | ||||||
|---|---|---|---|---|---|---|---|
1 |  Mini-SWE-agent | 16 / 30 | $0.07 | $0.2/$1.2 | 3m55s |
Key Takeaways
- This 30-task result is directional rather than a smoke test.
- GPT-5.6 Luna with Mini-SWE-agent records 235.10 seconds of latency per test.
Cost Analysis
Average Token Use / Test
Cost is the clearest tradeoff in this comparison. GPT-5.6 Luna leads at 53.33% for $0.07 per test. No other model in this comparison is cheaper.
Latency Analysis
Average Response Time / Test
GPT-5.6 Luna is both the most accurate and fastest model in this comparison at 53.33% and 3m 55s.
Tasks with failures
| Models | ||||||||||||||
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| GPT-5.6 Luna |
Task detail
bd9dcceIssue statement
Update the ring-attention and Ulysses-attention benchmark reference paths to use the FlashAttention 4 CUTE interface on both Hopper (H100) and Blackwell (B200). The ring benchmark must preserve non-causal attention and write into its preallocated output tensor. The Ulysses benchmark must use a single causal forward path that writes both output and log-sum-exp data into the provided tensors, regardless of GPU type. Because Ulysses partitions attention heads across ranks, its log-sum-exp buffer must be sized for the local head count rather than the global head count. Update the benchmark installation guidance to match the FlashAttention 4 package.
View Hidden Tests
diff --git a/valsmith_tests/test_flash_attention4_benchmarks.py b/valsmith_tests/test_flash_attention4_benchmarks.pynew file mode 100644index 00000000..9f903b2a--- /dev/null+++ b/valsmith_tests/test_flash_attention4_benchmarks.py@@ -0,0 +1,65 @@+import ast+from pathlib import Path+++ROOT = Path(__file__).resolve().parents[1]+++def parsed(relative_path):+ return ast.parse((ROOT / relative_path).read_text())+++def function(tree, name):+ return next(node for node in ast.walk(tree) if isinstance(node, ast.FunctionDef) and node.name == name)+++def fa4_imported(tree):+ return any(+ isinstance(node, ast.ImportFrom)+ and node.module == "flash_attn.cute.interface"+ and any(alias.name == "_flash_attn_fwd" for alias in node.names)+ for node in ast.walk(tree)+ )+++def calls_named(node, name):+ return [call for call in ast.walk(node) if isinstance(call, ast.Call) and isinstance(call.func, ast.Name) and call.func.id == name]+++def test_ring_benchmark_uses_flash_attention4_output_api():+ tree = parsed("kernels/parallel/ring_attn/benchmark.py")+ assert fa4_imported(tree)+ calls = calls_named(function(tree, "flash_attn_fwd"), "_flash_attn_fwd")+ assert len(calls) == 1+ keywords = {keyword.arg: keyword.value for keyword in calls[0].keywords}+ assert isinstance(keywords.get("out"), ast.Name) and keywords["out"].id == "O_flash"+ assert isinstance(keywords.get("causal"), ast.Constant) and keywords["causal"].value is False+++def test_ulysses_uses_one_flash_attention4_path_for_both_gpu_types():+ tree = parsed("kernels/parallel/ulysses_attn/benchmark.py")+ assert fa4_imported(tree)+ definitions = [node for node in ast.walk(tree) if isinstance(node, ast.FunctionDef) and node.name == "flash_attn_fwd_raw"]+ assert len(definitions) == 1+ calls = calls_named(definitions[0], "_flash_attn_fwd")+ assert len(calls) == 1+ keywords = {keyword.arg: keyword.value for keyword in calls[0].keywords}+ assert isinstance(keywords.get("out"), ast.Name) and keywords["out"].id == "O"+ assert isinstance(keywords.get("lse"), ast.Name) and keywords["lse"].id == "L"+ assert isinstance(keywords.get("causal"), ast.Constant) and keywords["causal"].value is True+++def test_ulysses_lse_buffer_matches_local_head_count():+ tree = parsed("kernels/parallel/ulysses_attn/benchmark.py")+ run = function(tree, "run")+ assignment = next(+ node for node in ast.walk(run)+ if isinstance(node, ast.Assign)+ and any(isinstance(target, ast.Name) and target.id == "L_tk" for target in node.targets)+ )+ call = assignment.value+ assert isinstance(call, ast.Call)+ head_dimension = call.args[1]+ assert isinstance(head_dimension, ast.BinOp) and isinstance(head_dimension.op, ast.FloorDiv)+ assert isinstance(head_dimension.left, ast.Name) and head_dimension.left.id == "H"+ assert isinstance(head_dimension.right, ast.Name) and head_dimension.right.id == "local_world_size"