gkapur
EDIT: I added to my profile after the above comment so it is a valid point.
These techniques can actually be combined and will be combined in the long-term I think. You may have LSM-based sandboxes that exist in Docker-based sandboxes.
Landscape and Perspective on Recursive Self-Improvement from a Neolab
TMA Transport for Matmul and Flash Kernels. cuBLAS and Flash Attention kernels still use the same cp.async transport on consumer Blackwell dies (in fact, most cuBLAS GEMM kernels on consumer Blackwell, including the FP16 tensorop path, are forward-ported Ampere-era cutlass_80_* kernels). Swapping cp.async with TMA allows us to reduce the number of instructions kernels need to issue, and the TMA’s swizzle drops shared-memory bank conflicts for free.
Hybrid FP16/FP32 Accumulation. The default matmul path on the production stack uses FP16 tensor cores with FP32 accumulation. However, on consumer dies, the FP16-input/FP32-accumulate HMMA runs at exactly half the rate of FP16-input/FP16-accumulate. To work around this, I use the fast atom mma_m16n8k16_f16_f16, but keep accuracy in check by promoting the FP16 partials into the FP32 registers and thus doing global accumulation accurately in FP32. So you get the FP16 tensor-core speed with an FP32 accumulation instead of paying the FP32-accumulate tax on every single mma. We measured the error of C = A@B using FP16 inputs drawn from N(0,1), comparing each accumulation strategy against an FP64 reference over the identical FP16-rounded operands. All configurations land within one standard error of each other, so the hybrid accumulation does not degrade task quality, matching the kernel-level error analysis.
The compiler approach allowed us to introduce both of these optimizations into all GEMM and Attention kernels used in the model rather than requiring hand optimization.
We were able to demonstrate that the compiler approach can match or even outperform vendor libraries. On Gemma 4 12B, a small open-weight model, running on RTX 5090, compiler-generated kernels achieved up to 1.6× speedup over cuBLASLt, and a 1.30× geometric-mean speedup over PyTorch.
Primarily, two optimizations allowed the Emmy compiler to outperform cuBLAS on common GEMM and FA shapes on RTX 5090 and RTX 4090: TMA transport (Blackwell-only) and leveraging full FP16 tensor cores with FP32 shadow accumulation registers:
TMA Transport for Matmul and Flash Kernels. cuBLAS and Flash Attention kernels still use the same cp.async transport on consumer Blackwell dies (in fact, most cuBLAS GEMM kernels on consumer Blackwell, including the FP16 tensorop path, are forward-ported Ampere-era cutlass_80_* kernels). Swapping cp.async with TMA allows us to reduce the number of instructions kernels need to issue, and the TMA’s swizzle drops shared-memory bank conflicts for free.
Hybrid FP16/FP32 Accumulation. The default matmul path on the production stack uses FP16 tensor cores with FP32 accumulation. However, on consumer dies, the FP16-input/FP32-accumulate HMMA runs at exactly half the rate of FP16-input/FP16-accumulate. To work around this, I use the fast atom mma_m16n8k16_f16_f16, but keep accuracy in check by promoting the FP16 partials into the FP32 registers and thus doing global accumulation accurately in FP32. So you get the FP16 tensor-core speed with an FP32 accumulation instead of paying the FP32-accumulate tax on every single mma. We measured the error of C = A@B using FP16 inputs drawn from N(0,1), comparing each accumulation strategy against an FP64 reference over the identical FP16-rounded operands. All configurations land within one standard error of each other, so the hybrid accumulation does not degrade task quality, matching the kernel-level error analysis. The compiler approach allowed us to introduce both of these optimizations into all GEMM and Attention kernels used in the model.
Exploring FlashAttention-3/4 optimizations on RTX GPUs
The kernel reaches parity with FA-2 (206us on RTX5090 with batch=1, heads=8, seq_len=4096, head_dim=64), but unfortunately, FA-3/4 optimizations are either not applicable or not helpful on consumer cards. It looks like FA-2 is the ceiling.
In summary: - Faster tensor-core instructions (WGMMA) are the main lever behind FA-3, but they are not available on RTX GPUs. - TMA (tensor memory accelerator) is available on sm_120 (RTX 5XXX). It helps on paper (LSU drops), but the transport isn't the bottleneck, so the final number barely moves. - Warp specialization is also available. However, it is mainly a scheduling optimization, i.e., it helps eliminate pipeline bubbles and better utilize tensor cores, but without asynchronous tensor core instructions. The result is negative: 213 vs 206 us. - In FA-4, they also simulated exp using FMA instructions because the tensor cores on the B200 are so fast that the whole pipeline became SFU-bound (special functions unit). RTX 5090 is tensor-core bound, so no point in this optimization either. In fact, even a conventional optimization of using faster exp2f instead of expf for softmax doesn't move the number.
I have tried a handful of other optimizations that could potentially work on consumer silicon, such as a deeper pipeline and register ping-pong. No luck. Since the whole pipeline is tensor-pipe-bound, I believe the FA-2 is the ceiling, and that all meaningful levers will require sacrificing some accuracy to leverage faster, lower-precision tensor cores.
Note that this is an exploration of the regular attention that dominates the prefill- and compute-bound regimes. Decoding against a large KV cache is a different, memory-bound story where split-KV/Flash-Decoding matters more than any of the above.
Since it pivoted, I imagine the story they are telling is the usage. But you can add to that Peter Fenton has had incredible success in commercial open source (Docker-ish, Elastic, JBoss, CockroachDB, TimeScaleDB, and if you include Benchmark there is Elastic, Confluent, etc.) so a bunch of these people are probably trying to ride the Peter Fenton/Benchmark train. My guess is that's what Theory is doing.
Investors are not free but they are very hype driven. There are < 10 investors who are truly experts in open source in my opinion in the market (and maybe < 5.) It's a really wacky cadence and I would argue there are different types of OSS businesses, which makes it even more difficult to be an expert. It's something I struggle with personally investing in OSS -- I think I understand the likely motion for an OSS startup and then > 50% of the time I am wrong.
I think the bet would have to be that a US Open Weight company either: 1. Gets a lot of money from Jenson who views them as a counterbalance to the big labs in his ecosystem and a way to generate leverage (the same way he is positioning neoclouds-- it also could be synergistic with neoclouds who could offer the model serving endpoints) 2. Can fast follow the same way Mistral does (which, honestly, seems like just distilling the Chinese model, which distills the US lab but is pretty innovative on a whole lot of architecture both in training and serving land.) 3. AND figure out some (maybe not super lucrative but lucrative enough) sort of business model, as well. There are lots of possible business models, so I will be curious how this whole space evolves.
The model could also be more flexible for non-coding use-cases (they show the results for reasoning being strong) so maybe the argument is to use it for non-coding use-cases to drive relatively deterministic conclusions for non-coding agents (they have also done some determinism work on kernels, which could be useful in pulling on that thread of deterministic models that are fine-tuned for everything that is not writing code.)
That said, I'm not sure how much all the work they have done actually synergizes or if the market size (at least in the short to medium term) is big enough for a huge outcome from the company's current valuation with those bets as the enterprise agent estate is taking a while to evolve. Hence companies like Anthropic and OpenAI are throwing tons of consulting money at the problem.
That said, the fine-tuning API + open weight model at least is a semblance of a viable business that could work so I will be curious about it. I'm not sure the synergy is fully there (why is someone with an open weights model privelaged to fine-tune it better if it's just QLora or Lora) but let's see!
Tutorial: Algebraic Foundations Powering FlashAttention
Part 1 is the theoretical foundation. It walks through a modern algebraic formalism showing that FlashAttention is an associative operation, which lets you treat it as a regular reduction on the GPU and apply all the same scheduling optimizations. Some recent MLSys and CVPR (ELSA) papers lean on this framing, and I find it much more powerful than the original.
This framework is particularly useful in ML compilers, where you should implement general optimizations applicable to many operations rather than writing specialized kernels. This article shows that attention actually belongs to a large family of "secretly-associative" operations, walks through a handful of examples, and links a few concepts from abstract algebra that let you identify whether an operation is secretly associative.
Overview:
- Safe softmax, Welford's variance, and FlashAttention are the same secretly-associative operation
- The twisted monoid (transport of structure), why the max-rescale coupling doesn't break associativity
- The qk_scale = log2(e)/√D like in FA-2 derived from scratch
- Numerical analysis: overflow bounds, error limits, and why tiling never amplifies error
- Bird's 3rd Homomorphism Theorem as a test for whether any loop is secretly associativ
Do you think this would scale to larger transformer models with more parameters per layer?
How would this work with MOE models or sparse models?
Surfacing a 60% performance bug in cuBLAS
A Principled ML Compiler Stack in 5k Lines of Python
Agentic systems redraw the Pareto frontier on ARC-AGI
The TLDR/key (from my understanding) is that verifying N tokens can be faster than generating N tokens.
Feels like two types of companies raised money: - Companies trying to couple the cloud with a programming language. - More recently, companies trying to couple GPUs with a programming language/alternative to CUDA.
Will be curious how this generation goes.
I don’t see those deals on Uber Eats so it feels like the real value of Uber One is for heavy Uber Eats users.
PS. Worth going through the cancellation flow when you are up for renewal as they will probably offer you 50% off Uber One.
VCs went a bit ham on this category when Auth0 got bought. I sense that the general thought process was: Auth0 multi-billion dollar company -> Auth0 will become crappy under Okta > replace Auth0 and build a bit company.
Dbt introduced a language for managing these “metrics” at scale including the ability to use variables and more complex templates (Jinja.)
Then you do dbt run (https://docs.getdbt.com/reference/commands/run) and kapow the metric is populated in your database.
More broadly dbt did two other things: 1. It pushed the paradigm from ETL to ELT (so stick all the data in your warehouse and then transform it rather than transform it at extraction time.) 2. It created the concept of an “analytics engineer” (previously know as guy who knows SQL or business analyst.)