Speaker: Hi, I'm Abhik. 30 minutes. By the end you'll know why a GPU is not a faster CPU, and how the same one line of Python travels from your interpreter down to the metal. This whole talk is that one line, c = a + b, taken four ways down to the metal.

Speaker: Quick who-am-I. I head machine learning at Cloudastructure, so my day job is ML systems in production, and nearly all of that is Python. A lot of that work is video processing, frames and pixels at scale, which is exactly the kind of work a GPU is built for. The thing I genuinely care about is performance: not switching languages when things get slow, but knowing what Python is doing underneath and making it fast where it stands. That is exactly the question this talk answers for GPUs: how fast can you go without leaving Python?

Speaker: Deliberately the most boring operation in computing. Element-wise add. The operation never changes, the machinery underneath changes completely.

Speaker: Hold this in your head. The intent is fixed. Everything from here is about who executes it, and where.

Speaker: Correct and honest about what's happening, and the slowest thing all day. Each c[i]: grab two objects, check types, unbox, add, re-box, store.

Speaker: Nothing here is about the math. It's all overhead around the math. Remember that phrase: move the loop somewhere better.

Speaker: This is the moment array programming was invented for. You stopped writing the loop, someone wrote it once in C, fast. The cells sit packed together now, one block of memory, not four scattered objects. Where did the loop go? Still there, just once in C, not a million times in Python.

Speaker: Key mental shift: you gave up the loop to gain the speed. That trade is the whole talk. Common trap: people hear "vectorized" and picture all cores lighting up. Not here. Element-wise ufuncs run in a single thread; SIMD is parallelism inside one core, several adds per instruction. Easy to verify: for a + b, cpu time equals wall time (one core); a BLAS matmul burns 12x wall on this laptop, because only BLAS-backed linalg fans out across cores.

Speaker: I changed the import and nothing else. The API is a promise: write NumPy, run on the GPU. Every c cell is a separate thread running this generated kernel. It really is that easy, until it isn't.

Speaker: This is the pivot. The easy path works right up to the moment your operation isn't in the library. Then you need to write the kernel yourself.

Speaker: This is where 'faster CPU' finally dies. You are not looping over elements. You launch a grid of programs; each grabs its BLOCK, loads, adds, stores. And it compiles all the way down to PTX. Notice the rhyme with Step 1: pure Python disassembled to CPython bytecode, Triton disassembles to GPU assembly. Same idea, opposite end of the ladder.

Speaker: Left: the loop is someone else's problem. Right: the loop is your problem again, but now on the right hardware.

Speaker: You would not run a marathon-length commute for the whole city on bikes, and you would not take a tram three blocks alone. Same trip, different machine. The bike is the CPU: one rider, off in seconds, turns down any alley, reroutes on a whim. Unbeatable for one person. The tram is the GPU: it takes a minute to fill and it only runs on rails, but once it moves it carries hundreds in lockstep. Per passenger the bike wins every time. But when the job is the whole city moving at once, which is exactly what c = a + b is, a million elements, you want the tram. Next slide: the same story in silicon.

Speaker: Same number of transistors, opposite bet. Look at the CPU side: the cores are almost the small part. The grey slab, cache, branch prediction, out-of-order machinery, is most of the die, and none of it does arithmetic. It exists so that one thread never has to wait: a latency machine. The GPU makes the opposite trade. Control shrinks to a sliver and nearly everything becomes arithmetic lanes. And here is the fact that kills the "faster CPU" model for good: each of those lanes is slower than a CPU core, roughly 1.5 gigahertz against 5, and dumber, no branch prediction, no out-of-order. Per thread, the GPU loses. And the counts make the bet concrete: a desktop CPU gives you eight or sixteen cores, a modern GPU more than ten thousand lanes. It only wins in aggregate, when you hand it the same operation over millions of elements, which is exactly our c = a + b. That is also why Step 1 was so painful: a branchy interpreter loop is the best possible food for the CPU's machinery and the worst possible food for this one.

Speaker: So what does the whole chip look like? A sea of SMs, streaming multiprocessors, plus one L2 cache they all share. For now treat an SM as a box that receives work; we will open one up on the next slide. When you launch a kernel you hand the hardware a grid of blocks, forty-eight here, and a little scheduler on the die deals them out to SMs as slots free up: any order, no promises, blocks cannot talk to each other. Each SM here holds two resident blocks, green and orange, the same pairs you saw land in Step 4. The blocks that did not fit simply wait; the moment a block retires, the next one drops in. And this picture is the answer to a question you should be asking: why does the same kernel run on my laptop and on an H100 without a single code change? Because the contract is just a bag of independent blocks. A bigger die is more SMs, which is more blocks in flight at once, and nothing else changes. That is the whole reason the grid model exists. So the obvious next question: what happens to a block once it lands on an SM? Let's open one.

Speaker: First, the name. When people say a GPU has ten thousand cores, that number is marketing: a CUDA core is just one arithmetic lane, it has no scheduler, no program counter of its own, you cannot give it work. The thing you can give work to is the SM, the streaming multiprocessor: a complete little processor with its own scheduler, its own lanes, its own registers and scratchpad. The GPU is a grid of these, and everything you launch is scheduled onto SMs, never onto cores. This is where the block you just watched land physically lives. Your block, 128 threads, gets cut into four warps of thirty-two, automatically, you never write that. The whole block lands on exactly one SM and stays there until it is done. Inside the SM: the scheduler picks one ready warp every cycle. Warp 0 is issuing, all thirty-two lanes fire together. Warp 1 is stalled on DRAM, so the scheduler simply does not pick it. Both of those are the next slide, so just note them for now. And note the bottom two boxes. Every resident warp keeps its registers on the SM the whole time it lives there, which is exactly why swapping warps costs zero cycles. And the shared memory, 192 kilobytes of on-chip scratchpad, belongs to the block. That little green box is the star of the second half of this talk: tiling is nothing but choreographing data through it. A real SM is a bit wider than this picture, four schedulers with thirty-two lanes each, but the model is the same, and the die you just saw is this box repeated about a hundred times. It is also why the block is Triton's unit of thinking. Now I have used two words without earning them, warp and stalled. Let's slow the scheduler down and watch it work.

Speaker: Two ideas here. First, SIMT: one instruction stream drives a whole warp, thirty-two lanes executing the same op in lockstep. You do not write thirty-two threads, you write one and the hardware runs the width. The warp matters because it is the unit the hardware actually schedules: it is what you just watched the SM's scheduler pick between, and it is what your 128-thread block was cut into four of. Second, the payoff of the whole throughput bet: latency hiding. When warp 0 issues a load, the data is in DRAM, hundreds of cycles away. On a CPU that would be a disaster, which is why the CPU carries that huge cache. The GPU shrugs: the scheduler parks warp 0 and issues warp 1, then warp 2, all ready to compute. And the key trick is number three: the swap costs nothing. Every resident warp keeps its own registers on the SM the whole time, so switching warps is not a context switch, it is just picking a different ready warp next cycle. Zero cost. That is why you launch thousands of threads for a chip with far fewer lanes: the surplus is not waste, it is the fuel the scheduler burns to keep the arithmetic units busy while memory crawls. Cache versus more warps, same enemy, opposite weapon. And notice what the whole trick rests on: DRAM being hundreds of cycles away. So before we pick Triton back up, one more zoom out: where do the bytes actually live?

Speaker: The picture people carry is a GPU chip next to some memory. Make it accurate. The die is a sea of SMs, and the fast memory lives with them: registers and L1 slash shared are inside every single SM, a couple hundred kilobytes each, and one L2, around forty megabytes, is shared by all of them. That is the fast stuff, tens of terabytes a second, effectively instant. Notice how little of it there is: all the SRAM on an A100 adds up to well under a hundred megabytes. The big memory, the VRAM you see quoted in gigabytes, is off the die entirely, HBM stacks sitting next to the chip, reached over the memory bus at about 1.6 terabytes a second. That is still huge, but it is roughly twelve times slower than on-chip, and every miss pays that toll. One correction people always need: that die-to-VRAM link is the memory bus, not PCIe. PCIe is a completely different, much slower road, thirty-odd times slower again, and it only matters when you copy between the CPU and the card. So there are two gaps, not one. Hold onto this picture: the entire second half of this talk will be one instruction, keep your data on the die.

Speaker: This is the one fact the whole rest of the talk hangs on. Registers are instant, shared memory is on-chip and fast, HBM is far away and slow. At the intensity a naive kernel runs, the cores want about 78 terabytes a second and the memory hands them 1.6. So the cores starve. Any kernel that is bottlenecked here runs at single digit percent of peak no matter how fast the math unit is.

Speaker: Switch examples here. Add was memory bound and honest, but every byte was used exactly once, so there is nothing to optimize. Matmul is the opposite: each row of A feeds an entire row of outputs, each column of B feeds a column. The naive kernel throws that reuse away. One thread computes one output by reading a whole row and a whole column from DRAM, and the thread next to it re-reads that same row all over again. Same math as a good kernel, but drowning in redundant loads.

Speaker: Here is the fix. Instead of each thread reaching all the way to DRAM for its own row and column, the whole block cooperates: load one tile of A and one tile of B into shared memory, once. Barrier. Now every thread does its multiplies out of that fast on-chip scratchpad, and each loaded value serves TILE different threads before it is thrown away. Then slide the tiles along K and repeat. The DRAM traffic for this block just fell by a factor of TILE.

Speaker: Yes, it is another DSL. I know exactly what that sounds like in 2026, so bear with me for one slide, because the bill is smaller than it looks: it is still Python, it is still one file you can pip install and run, and there is exactly one new idea in it. Here is the whole model in one slide. In CUDA you write the code for one thread and reason about thousands running in lockstep. In Triton you write the code for one block, using array operations. That is the one idea. Line by line: program_id asks which block am I, the only parallelism you get. arange builds the vector of indices this block owns, so x plus y operates on the whole block at once. mask guards the tail because the length rarely divides evenly. load and store are explicit on purpose: moving bytes is where GPU performance lives, so it is a line you write. BLOCK is a compile time constant so Triton can specialize and unroll. And the deal underneath: you own the block algorithm and the memory movement, the compiler owns mapping the block onto warps, coalescing the loads, and allocating shared memory. That middle layer is exactly what you hand-write in CUDA C++.

Speaker: The decorator does nothing at definition time. The first time you call the kernel, Triton grabs the function's source with inspect and runs Python's own ast.parse on it, the same module you can import. Your function is never executed as Python. What you see here is the real tree for the store line: a Call node, and inside its arguments two BinOps, pointer plus offsets and x plus y. That inner BinOp is the one we chase. Your def is not code to Triton, it is data, and the next two slides follow it down the pipeline at the top.

Speaker: Two stages, one before the hardware and one after. Triton IR first: x plus y is a single addf on a tensor of 1024 floats. The whole block is one value. No threads, no warps, no lanes; this is exactly the mental model you wrote in, load, add, store, block at a time. Then the lowering to Triton-GPU IR, and the only thing that changes on this line is the type. The tensor grows a layout attribute: sizePerThread one, threadsPerWarp thirty-two, warpsPerCTA four. That suffix is the hardware arriving. It says: these 1024 elements will be carried by 128 threads in 4 warps. And notice you never picked 128 anywhere. BLOCK equals 1024 is elements, not threads; how many threads carry them is the compiler's call.

Speaker: Last two stages. LLVM IR is per-thread code, so this is the moment the block-wide add shatters: eight scalar fadds, because each of the 128 threads owns eight of the 1024 elements. Everything block-shaped is gone; from here down it is one thread's program. PTX is those same eight adds in the GPU's assembly dialect, with every store predicated on the mask, and ptxas takes it the last step into SASS, the actual machine code. And the punchline of the whole pipeline: all of this happens once, on the first call, then it is cached. The key is dtypes plus constexpr values plus GPU arch, so passing fp16 instead of fp32, or changing BLOCK, compiles a brand new kernel. After that, calling the function costs only a launch. Next: let's actually read that PTX.

Speaker: Same file as the last slide, now reading the body. Top: the entry point takes your four pointers, and reqntid is the contract: launch this with exactly 128 threads or the driver refuses. First the kernel asks where it is: ctaid is the block index, program_id in Triton, and there is tid, the per-thread index you never wrote; the compiler introduced it when it split the block. pid times BLOCK is a shift left by ten, the constexpr payoff in one opcode. The or chain builds the 8 offsets each thread owns, 128 apart, and that spacing is deliberate: at every step the 32 threads of a warp read 32 neighbouring floats, one clean coalesced transaction, exactly the DRAM-friendly access the memory pyramid demanded. Then the mask: setp writes 8 predicate bits, and every load and store is prefixed with at-p: the instruction runs on all threads, the hardware mutes the lanes where the predicate is false. No if, no branch, no divergence; masked lanes just keep the zero. Eight adds, eight predicated stores, ret. There is no loop in this kernel and no branch either. And the loc lines mean the toolchain can walk every instruction back to your Python source line, which is exactly what a profiler does.

Speaker: I do not want you to take any of the last four slides on faith, because none of it is hidden. The first time you call a Triton kernel, the compiler writes every stage it just built to a folder under tilde slash dot triton slash cache, one folder per kernel, and the folder name is the cache key. Look at the listing: dot source, ttir, ttgir, llir, ptx, cubin. That is the exact breadcrumb strip from the last four slides, sitting on your filesystem as files you can open in vim. Bottom left is the cache key made concrete: I compiled the same add kernel for three architectures and got three folders, because arch is part of the key, exactly as the block-shatters slide claimed. The json tells you what it keyed on and what the compiler chose for you: four warps, three stages. And the knobs are worth memorising. Cache dir puts it somewhere you can watch. Always compile ignores the cache when you are iterating. And kernel dump gives you every intermediate pass plus add dot sass, which is the one rung this deck never showed you, the real machine code that ptxas produced, as readable text. That is the whole point: this is not a black box you pray to. It is a compiler that leaves its homework on the floor.

Speaker: Here is the payoff of the whole pipeline arc, and the real reason to reach for Triton. This is the exact same kernel from the last three slides. I compiled it twice on this laptop, no GPU plugged in, changing one thing: the target. On the left, NVIDIA, it lowers to PTX, the ld-add-st you already read. On the right, the identical Python lowers to AMD GCN for an MI300: global-load, v-add-f32, global-store. Look at how much is different underneath. Different assembly dialect, different register model, and even the warp is a different size, 32 lanes on NVIDIA, 64 on AMD, so Triton picks a different block shape to match. You wrote none of that. In raw CUDA C++ or HIP you would maintain two codebases; here it is one .py and a target flag. And it does not stop at these two: NVIDIA and AMD are in the wheel today, and there are out-of-tree backends for Intel GPUs, for CPUs, and an experimental Apple-silicon path. The bet Triton makes is the same bet MLX, XLA and torch.compile make: describe the computation once, let a compiler specialize it to whatever silicon you land on. That is what "you don't program the hardware, you program the compiler" actually buys you.

Speaker: One more axis of the same portability story: not across vendors this time, but across time. Compute capability is NVIDIA's version number for a GPU generation: 8.6 is Ampere, from 2020; 12.0 is Blackwell, from 2025. In between the machine underneath changed completely: more SMs, more bandwidth, new tensor cores, new instructions. Here is the same add kernel compiled for both, and the honest surprise: the PTX differs in exactly two lines, the header. PTX is a virtual instruction set, like bytecode for GPUs. The last translation, ptxas down to SASS, the real machine code, is where the generations diverge, and the compute capability is how ptxas knows which machine to build for. Do not over-generalize the two-line diff though: this add is simple. The moment a kernel touches tensor cores, a tl.dot, the PTX itself changes per generation, because each one has its own mma instructions and the compiler picks them for you. That is the real point of the slide: the cc appeared in the cache key three slides ago, and you never wrote it. Triton reads it off whatever GPU the process starts on and specializes everything downstream. We ran this deck's kernels on both generations, unmodified.

Speaker: One loose thread from the whole Triton arc: where did BLOCK equals 1024 come from? I picked it. And the honest answer to "what is the best value" is: it depends, on the generation, on the memory system, on the size of n. Triton's answer is to stop pretending you know. autotune hands the decorator a menu of configs instead of one answer. Each config is a set of constexprs, and you already know what that means: each one compiles to its own kernel through the cache we saw. On the first call with a new key, here the size bucket of n, Triton simply runs the race: compile each config, time it on your real GPU with your real data, cache the winner for that key. Every call after that is just the winning launch. Notice what is not in your code: no table of architectures, no heuristics, no model of the hardware at all. The measurement is the oracle. That closes the loop on the last two slides: the same file ran on AMD and on two NVIDIA generations five years apart, and with a menu on top it is not just correct on each of them, it re-derives its own fastest shape on each of them. For our add the stakes are small, but hold this thought for the matmul half: there the knobs are BLOCK_M, BLOCK_N, BLOCK_K, warps, stages, they interact, and the space is genuinely unguessable by hand. Every serious production Triton kernel ships with an autotune menu on top.

Speaker: This is the canonical kernel, and notice how little changed. The multiply-add in the middle is exactly the naive kernel. The whole trick is the two shared arrays and the loop that stages tiles into them. TILE is a compile time constant, which is why the compiler can unroll the inner loop and keep the accumulator in a register. The two barriers are the price of using shared memory: the hardware gives you the SRAM but not the coherence, so you do it by hand. Forget the first barrier and threads compute on half loaded tiles; forget the second and you overwrite a tile someone is still reading. (ty, tx abbreviate threadIdx here for space.)

Speaker: Step back from the mechanics for one slide, because this is the whole reason Triton exists. Everything we have built is one trade. On the left, NumPy and CuPy: you write c equals a plus b and you are done, but you only get the operations someone already wrote, and the moment your problem is not in the library you fall off the cliff we hit earlier. On the right, raw CUDA or HIP C++: total control, the highest ceiling, but you manage every thread by hand, you carry a separate toolchain, and you maintain two codebases if you care about both vendors. Triton is the middle tier, and it is not a compromise, it is the sweet spot: you keep the two things that actually decide GPU performance, the block algorithm and where the bytes move, and you hand the compiler the parts you should never hand-write, mapping blocks to warps, coalescing, shared memory, register allocation. You write it in Python, it JITs from your interpreter, and the same source runs on NVIDIA and AMD as we just saw. And to kill the idea that this is a lecture-hall toy: PyTorch's torch.compile lowers your model to Triton kernels under the hood, and a lot of the fused attention and normalization kernels in modern LLM stacks are written in it. This is the production path, and it is the one you can actually read. One first-half habit left to bank, the asynchronous launch, and then the second half: what owning the memory buys on a real problem.

Speaker: Consequence of writing GPU code from Python, line by line. Launch A: the call returns in about ten microseconds, kernel A starts on stream 0. Launch B: queued behind A, the stream is a FIFO. Meanwhile Python keeps going, CPU and GPU overlap for free. The only stop is d.cpu(), the first sync: you asked for the bytes. Which also means a naive time.time() around the launch measures nothing, the kernel has not run yet. Remember this when we look at real numbers later.

Speaker: Same operation, same tiling, but look what disappeared. There is no __shared__ declaration and no __syncthreads. You loop over K in TILE-sized phases, load a tile of A and a tile of B with block-shaped index math, and tl.dot multiplies the two tiles, which is exactly the As-times-Bs inner loop of the CUDA kernel run across the whole block. The staging into shared memory and the two barriers still happen on the hardware, but the compiler emits them. This is the whole thesis of the Triton half: you kept the block and memory control that makes tiling fast, and you gave up the per-thread bookkeeping that makes CUDA C error prone. The math is identical, the data races are gone.

Speaker: This is the whole story in one picture. The naive kernel sits down in the corner at 0.25 flop per byte, pinned to two percent of peak. Every time you double the tile you double the reuse, which doubles arithmetic intensity, which slides the attainable performance up the diagonal bandwidth roof. You keep climbing until you hit the flat compute roof near the ridge point. In real numbers on a 4096 cubed matmul, tiling with a 32 wide tile cuts DRAM traffic from about 550 gigabytes to 17, and the time spent waiting on memory from about 354 milliseconds to 11, which is now less than the 7 milliseconds of actual math. The kernel went from hopelessly memory bound to basically balanced. That is tiling. Next: fusion, the other half of respecting the hardware.

Speaker: One practical closer, and you do not need an external profiler for it. Proton ships inside Triton. It hooks the vendor profiling layer, CUPTI on NVIDIA, roctracer on AMD, and uses a cheap shadow pass so every kernel is charged back to the scope you named, not some mangled kernel symbol. The result is a call tree you can slice. The real trick is the last two lines: if you tell a scope how many flops and bytes its work costs, the viewer derives throughput, and best of all a util number, which is just the kernel's position on the roofline, the max of your compute fraction and your bandwidth fraction. So the same which-wall question we reasoned about analytically two slides ago, Proton answers with a measurement, per kernel. Next slide: let's actually run it.

Verified against Triton 3.7.0/3.7.1 Proton: scope metrics {flops,bytes} are stored on the scope node (probed on the GPU box); viewer derives flop/s, gflop/s, tflop/s, byte/s, gbyte/s, tbyte/s and util = max(sum(flops)/peak_flops_time, sum(bytes)/peak_bw_time) (triton/profiler/viewer.py). Backends cupti/roctracer/instrumentation, data tree|trace, all real. NOTE: hook="triton" does NOT auto-populate flops/bytes; you must attach them to the scope as shown.

Speaker: So let's actually run it. Left is the whole thing: import the profiler, start it, wrap each launch in a named scope, finalize. That is the entire diff to your script. Then run it under the proton command and read it back with proton-viewer. Right is what came back, and I want to be clear these are real: same matmul, same GPU, same four thousand cube of true fp32, and the only thing that changed is the tile size. The sixteen wide tile throws reuse away and sits at a fifth of peak. The one twenty eight wide tile keeps it and more than doubles the throughput, forty odd percent, off the memory wall. That is exactly the climb the roofline slide predicted, except now nobody is taking my word for it. Respecting the hardware is not a vibe, it is a measurement, and the tool is one import away.

These numbers are REAL: captured on a Blackwell GPU (dc-03-node21) with Triton 3.7.0 + Proton, true fp32 (input_precision="ieee"), 4096^3. Per-scope GPU time straight from proton.hatchet: small_tile 26.36ms/5=5.27ms, large_tile 11.60ms/5=2.32ms. TFLOP/s = 2*4096^3 / time. Peak basis ~136 TFLOP/s fp32 = 170 SM * 128 * 2 * 3.12GHz (max boost); at the card's 2.57GHz recorded clock the peaks read ~23%/53% instead. Repro: research/proton/matmul_triton.py.

Speaker: To close, one real benchmark, the kind of workload my day job is full of. Here is the input: Krakow's main square, 2.8 megapixels. The job: convert it to grey and blur it with an 11 by 11 gaussian, which means every output pixel is a weighted sum of its 121 neighbours. Two contenders. torch, where the whole thing is a few library calls with cuDNN underneath, and Triton, where we write the kernel ourselves with everything this talk taught. Same GPU, same math, and every number that follows is measured, not guessed.

Speaker: The two programs side by side. Torch on the left: grayscale as a tensor expression, pad, one conv2d, cuDNN underneath, and torch.compile gets its best shot too. Triton on the right: everything this talk taught, one fused kernel, about forty lines. Each program owns a 16 by 16 tile of the output, loops over the 121 neighbours, converts each to luma on the fly, so the RGB is read once and the result written once. Round 1, measured: a tie, 1.52 against 1.61 milliseconds. That is an honest result worth internalizing: convolutions are among the most optimized code paths on the planet, and when your workload is exactly the shape the library was tuned for, a hand-written kernel buys you nothing. But this fight was on the library's turf, so we change the ground: change the algorithm.

Speaker: Before we optimize it, look at the kernel actually working. This is the Main Square, one image, cut into five hundred and twelve tiles. Watch the wavefront: sharp colour turning to soft grey, tile by tile. That is the exact filter on the last slide, grayscale plus gaussian blur, and every tile is one program, the pid_x pid_y you just read in the code. Not one fast worker crawling over the photo, but hundreds of small ones each owning a tile, the same cheap kernel on all of them at once. I have shown you it in a grid of little boxes all talk; this is the same picture at the size of a real photo. Now, it is a tie against torch, so let us make it faster.

Speaker: Here is the optimization, and it is an algorithm change, not a code tweak. The blur is a weighted sum, and the weights form a grid. A gaussian's grid is special: it is an outer product, every entry is its row weight times its column weight. The three-tap example shows it, the centre four is two times two. And because the weights factor, the sum factors, and I do not want you to take that on faith, so watch it happen to a real tile of the Rynek. Top path: the full 11 by 11 window, 121 reads per pixel, that is the target. Bottom path, pass 1, 11 taps along x only, and look at it: smeared sideways, the ledges still sharp, the blur is visibly half done. Then pass 2 runs 11 taps down y, and the result lands right under the target. Identical. Not close, identical: we measured the two images, the biggest pixel difference is zero out of 255. So 121 reads become 22, five and a half times less arithmetic, before we touch a single line of GPU code.

Speaker: So why not keep it fused, one kernel that does both passes on-chip? Because of who owns the data. Pass 2 blurs vertically: one output pixel needs eleven rows of pass 1's output. Look at who wrote those rows: three different programs. A program's tile lives in its own registers and shared memory, no program can see a neighbour's, and Triton cannot re-index an on-chip tile at shifted offsets either, tensor slicing is unsupported. So the horizontally-blurred image has to be written out to DRAM and read back by a fresh kernel. The fast path is three launches: grayscale, blur along x, blur along y, DRAM between each. On the last slide we fused because kernel round trips felt wasteful, and now I am telling you to take three of them. And it is still 5.2 times faster than the fused kernel, because the algorithm cut the arithmetic five-fold and a launch costs microseconds. Doing less work beats doing fewer launches.

Speaker: The full scoreboard, measured on the same GPU. Bars are milliseconds per call, shorter is better. Top group: the 2D algorithm, where Triton loses to everybody, even plain cuDNN. Bottom group: the separable algorithm, and torch gets the same trick, two conv1d passes, so this is identical math on both sides. Hand-written Triton: 3.3 times faster than torch, and still 2.3 times faster than torch.compile. That compile row is the interesting one: Inductor fused the grayscale into a Triton kernel it generated, so the bottom two bars are generated Triton against purpose-built Triton. Why does hand-written win? cuDNN's convolution is built for deep learning shapes, big batches, many channels, and a single-channel gaussian pays for all that generality; three tight kernels that do exactly this job beat it. One honest footnote: on Blackwell the same code still wins, but by 1.7x, because newer libraries and four times the bandwidth hide more sins. So the closing thought: when your workload is not the shape the library was tuned for, Triton lets you write the kernel that fits, without leaving Python.

These numbers are REAL: measured 2026-07-16 on the Ampere box (driver 580.159.03, torch 2.13.0+cu130, triton 3.7.1) via research/blur-kernel/blur_gray.py --gpu. do_bench per call: torch 2D 1.521 / torch sep 1.011 / tc 2D 1.205 / tc sep 0.714 / triton fused 1.607 / triton sep 0.307 ms (3.29x / 2.33x). Bar widths = time/1.61 rounded: 94/75/100/63/44/19%. Blackwell validation (see CLAUDE.md): triton sep 0.097 ms, 1.67x vs torch sep, 1.21x vs torch.compile. Wall clock 100 passes: torch sep 0.100 s vs triton sep 0.029 s.

Speaker: That's the whole talk. Same one line of Python, four ways, down to the metal, then respect the hardware and the numbers pay off. Thanks for listening, happy to take questions.