--- title: Learning FlashAttention the Hard Way | daily.dev description: A deep technical walkthrough of generating efficient CUDA kernels for FlashAttention-2 using the Emmy compiler, without handwriting them manually. Starting... canonical: https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw twitter:card: summary_large_image twitter:site: @dailydotdev og:type: website og:site_name: daily.dev og:title: Learning FlashAttention the Hard Way | daily.dev og:description: A deep technical walkthrough of generating efficient CUDA kernels for FlashAttention-2 using the Emmy compiler, without handwriting them manually. Starting... og:url: https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw og:image: https://api.daily.dev/og/posts/oS5uNtnuw.png og:image:alt: Learning FlashAttention the Hard Way og:image:width: 1200 og:image:height: 630 og:locale: en --- > ## Documentation Index > Fetch the complete documentation index at: https://daily.dev/llms.txt > Use this file to discover all available pages before exploring further. # Learning FlashAttention the Hard Way **[ITNEXT](https://daily.dev/sources/itnext)** · 55 min read · 0 upvotes · 0 comments ## Summary A deep technical walkthrough of generating efficient CUDA kernels for FlashAttention-2 using the Emmy compiler, without handwriting them manually. Starting from a naive scalar streaming kernel, the post walks through four major optimizations: switching to tensor cores (42x speedup), shared memory staging (1.20x), double-buffered pipelining with cp.async (1.05x), and TMA transport (1.04x). Each optimization is explained with actual generated CUDA code, PTX inline assembly, and measured latency on an RTX 5090. The autotuner independently arrives at the same tile geometry FlashAttention-2's author chose by hand, matching its latency to the microsecond. Optimizations from FA-3 and FA-4 that were tried and rejected are also documented. ## Full article daily.dev links to this article rather than hosting it. Read it at the original source: ## Similar posts on daily.dev - [Tuning Flash Attention for Peak Performance in NVIDIA CUDA Tile](https://daily.dev/posts/tuning-flash-attention-for-peak-performance-in-nvidia-cuda-tile-ytpdwf7wh) · NVIDIA Developer · 0 upvotes · 0 comments - [We reverse-engineered Flash Attention 4](https://daily.dev/posts/we-reverse-engineered-flash-attention-4-vuuzhgl36) · Modal · 0 upvotes · 0 comments - [Overcoming Compute and Memory Bottlenecks with FlashAttention-4 on NVIDIA Blackwell](https://daily.dev/posts/overcoming-compute-and-memory-bottlenecks-with-flashattention-4-on-nvidia-blackwell-cb1uqp4vo) · NVIDIA Developer · 0 upvotes · 0 comments - [FlashAttention 4: Faster, Memory-Efficient Attention for LLMs](https://daily.dev/posts/flashattention-4-faster-memory-efficient-attention-for-llms-xzpyphq6n) · DigitalOcean Community · 1 upvotes · 0 comments --- Tags: [#gpu](https://daily.dev/tags/gpu), [#cuda](https://daily.dev/tags/cuda) [View this post on daily.dev](https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw) ```json {"@context":"https://schema.org","@graph":[{"@type":"Organization","@id":"https://daily.dev/#organization","name":"daily.dev","url":"https://daily.dev","logo":{"@type":"ImageObject","url":"https://daily.dev/apple-touch-icon.png","width":180,"height":180},"sameAs":["https://twitter.com/dailydotdev","https://github.com/dailydotdev","https://www.linkedin.com/company/daily-dev-ltd"]},{"@type":"WebSite","@id":"https://daily.dev/#website","url":"https://daily.dev","name":"daily.dev","publisher":{"@id":"https://daily.dev/#organization"},"potentialAction":{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https://daily.dev/search?q={search_term_string}"},"query-input":"required name=search_term_string"}}]} {"@context":"https://schema.org","@type":"TechArticle","headline":"Learning FlashAttention the Hard Way","url":"https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw","mainEntityOfPage":{"@type":"WebPage","@id":"https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw"},"datePublished":"2026-07-27T21:44:31.264Z","dateModified":"2026-07-27T21:44:59.476Z","description":"A deep technical walkthrough of generating efficient CUDA kernels for FlashAttention-2 using the Emmy compiler, without handwriting them manually. Starting...","image":"https://media.daily.dev/image/upload/f_auto,q_auto/v1/posts/858a4877e8a89cb61afe8ae779274cf9?_a=AQAEuop","thumbnailUrl":"https://media.daily.dev/image/upload/f_auto,q_auto/v1/posts/858a4877e8a89cb61afe8ae779274cf9?_a=AQAEuop","isAccessibleForFree":true,"articleSection":"ITNEXT","inLanguage":"en","publisher":{"@type":"Organization","name":"daily.dev","url":"https://daily.dev","logo":{"@type":"ImageObject","url":"https://daily.dev/apple-touch-icon.png","width":180,"height":180}},"author":{"@type":"Organization","name":"ITNEXT","logo":"https://media.daily.dev/image/upload/t_logo,f_auto/v1/logos/itnext2","url":"https://daily.dev/sources/itnext"},"commentCount":0,"discussionUrl":"https://daily.dev/posts/learning-flashattention-the-hard-way-os5untnuw","interactionStatistic":[{"@type":"InteractionCounter","interactionType":{"@type":"LikeAction"},"userInteractionCount":0},{"@type":"InteractionCounter","interactionType":{"@type":"CommentAction"},"userInteractionCount":0}],"keywords":"gpu,cuda","timeRequired":"PT55M"} {"@context":"https://schema.org","@type":"BreadcrumbList","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https://daily.dev"},{"@type":"ListItem","position":2,"name":"ITNEXT","item":"https://daily.dev/sources/itnext"},{"@type":"ListItem","position":3,"name":"Learning FlashAttention the Hard Way"}]} ```