- ptxas has enough good heuristic

- in average you can reduce ~3% of stall counts

- overall speed up is not equivalent to the number of optimized stall counts

Kernel engineers are not obsolete. But asking a general-purpose coding agent to rediscover years of CUDA and Triton engineering knowledge every time it writes a kernel probably should be.

After months of writing, debugging, and optimizing kernels, I turned the reasoning patterns I kept using into an open-source skill library for AI coding agents:

npm install u/krxgu/kernel-skills

This is not a collection of vague prompts saying “make this CUDA kernel faster.”

Each skill is a detailed engineering playbook that forces the agent to think about:

• Exact shapes, dtypes, layouts, and target hardware before writing code
• Coalescing, tiling, bank conflicts, occupancy, and register pressure
• Numerical stability and non-power-of-two boundary conditions
• Correctness tests across adversarial shapes and dtypes
• Whether a custom kernel should exist at all
• When to stop being clever and use cuBLAS, CUTLASS, or an existing primitive

The library currently covers CUDA, Triton, INT8 and FP8 quantization, kernel fusion, CUDA to Triton and HIP portability, and inference hot paths including RMSNorm, fused add plus RMSNorm, RoPE, sampling, paged KV-cache append, dequantization, prefill versus decode, and vLLM custom-op integration.

I also did not want this to become prompt-engineering theatre, so the repository includes before-and-after proof runs using the same model and task, with the skill file being the only difference:

• Softmax: naive output failed on adversarial and larger shapes. Skill-guided output had 0 failures across 16 tests and reached within 1.2% of torch.softmax bandwidth
• Reduction: 2.6 to 3.5x faster than the naive agent output
• GEMM: 7.7 to 8.6x faster
• LayerNorm: 1.9 to 3.2x faster
• Triton softmax: fixed crashes at dimensions above 16,384 and worked up to 131,072
• Triton attention: fixed the common GQA failure where H_q != H_kv

To be completely clear, those speedups are against the naive agent-generated kernels, not against cuBLAS or other vendor-tuned libraries. In fact, the GEMM skill explicitly tells the agent not to write a custom kernel when cuBLAS or CUTLASS already solves the problem.

Example:

kernel-skills bundle \
  triton.write-triton-layernorm-kernel \
  patterns.write-numerically-stable-kernel \
  patterns.write-kernel-test-plan \
  > bundle.md

Give that bundle to Claude Code, Cursor, ChatGPT, Gemini CLI, or another coding agent before asking it to touch the kernel.

The spicy thesis is simple:

Models are increasingly interchangeable. The accumulated engineering judgment surrounding them is not.

Everything is open source and MIT licensed:

https://github.com/tensormux/kernel-skills

I would especially love kernel engineers to tear this apart.

Which skill is missing? Which technical rule is wrong? Where can an agent still produce something that looks convincing but quietly fails on real hardware?

initially opened the Kimi K3 release, probably like everyone else on the planet, just because of the obvious headline: 2.8 trillion parameters.

After spending a few days digging through the architecture, I actually think that is the least interesting number.

The interesting question that sparked into my mind is: what happens when you scale model capacity while aggressively keeping active computation under control?

K3 activates 16 of 896 experts per token. That means the model is storing a massive amount of capacity, but only a small fraction participates in each forward pass. this actually shifts the bottleneck

The million-token context is also interesting because it changes the economics. Long context is usually presented as a capability problem, but at this scale it is really a memory and serving problem.

The question I keep coming back to over and over: are frontier models moving from being primarily "large neural networks" toward being "large distributed systems with neural networks inside"?

Curious what people working on MoE systems think. Is expert routing/load balancing becoming the real scaling bottleneck?

(I wrote a longer analysis with some simulations, but I'm mostly interested in the discussions here.)

https://open.substack.com/pub/softwarefrontier/p/kimi-k3-28-trillion-parameters-four?r=3c7w5a&utm_medium=ios

I’ve open-sourced REAP CUDA: a CUDA-focused implementation of Router-weighted Expert Activation Pruning for MoE models.

GitHub: Link

The core idea is simple: don’t prune experts by how often they are selected alone. Measure their router-weighted activation contribution over calibration data, accumulate a saliency score, then remove the least useful experts structurally.

For LFM2.5-8B-A1B, the pipeline is:

calibration batches

→ MoETransformerObserver

→ observe router activations

→ accumulate expert saliency

→ rank experts

→ slice expert tensors

→ update model config

→ save a valid pruned checkpoint

On Liquid AI’s LFM2.5-8B-A1B, we used it to cut every MoE layer from 32 experts to 16 across 22 MoE layers:

- 8.47B / 16.94 GB base checkpoint

- 4.59B / 8.57 GB after REAP-50 pruning

- No retraining required

- Recorded MATH500: 88.76% base → 77.0% pruned

- Recorded BFCLv3 single-turn: 64.79% base → 59.07% pruned

We also tested the pruned checkpoint in a separate, external AWQ INT4 stage. That produced a 2.79 GB packed artifact—but to be clear, AWQ is not implemented by REAP CUDA, and the quality/vLLM measurements for that stage used an AWQ-scaled BF16 evaluation derivative rather than direct packed-INT4 serving.

The part I’m most excited about is that this is structural pruning, not masking: the actual expert tensors are sliced, the config is updated, and the result is a smaller checkpoint that can be inspected and deployed.

Would love feedback from people working on MoE routing, expert redundancy, calibration strategies, and pruning criteria—especially on where router-weighted activation saliency breaks down versus more expensive second-order approaches.

Put together a post on profiling CUDA for inference serving. The usual dev-time approach (full kernel traces, Nsight captures) doesn’t really work once you’re running real load, especially without root in containers/K8s.

We built something for this: low-overhead GPU profiling grouped by what the engine is actually doing (attention, matmul, comms, KV cache, etc.) instead of dumping every kernel, plus tying it to vLLM/SGLang traces and GPU metrics. Works with graphsignal-run vllm serve ...

Curious how people here profile vLLM/SGLang (or similar) when something’s off with Nsight, PyTorch profiler, or just guessing from throughput?

Hey everyone. I’m a 15-year-old solo systems architect and AI-native researcher operating out of a small industrial employee-town in Iran, coming from an ordinary working-class family.

Due to extreme geopolitical embargoes, severe internet blackouts, and absolutely zero access to global cloud compute (no AWS, no GCP), I am essentially locked in a digital cage. But logic has no borders. My ultimate vision is to break out of this geopolitical prison, secure an O-1A Visa sponsorship, relocate to the US, and dedicate my life to AI Safety, Deep-Tech, and pushing the boundaries of human biology.

To prove my capability from the ground up, I orchestrated Google Gemini 3.1 Pro (via 6 Million tokens of pure adversarial prompt engineering and internal safety-filter bypasses) to help me architect the "Zero Cancer Reactor".

It is a brutally optimized, GPU-bound mathematical simulator modeling the Tumor Microenvironment (TME) at a 1:1 local tissue scale (70 Million autonomous cells) at a locked 60.0 FPS.

I am looking for brutal, senior-level feedback on my architecture from C++ veterans, AI researchers, and computational biologists:

• Zero Fake-IFs (Empirical Cytokine & Endocrine Networks): There are no hardcoded death triggers or pre-scripted events. Every single interaction is driven by real-time non-linear differential equations, Michaelis-Menten kinetics, and stochastic Gaussian noise. The system actively simulates a highly complex cytokine mesh: Macrophages dynamically polarize between M1 (aggressive killers driven by IL-12 and IFN-gamma) and M2 (executing efferocytosis and tissue repair via IL-4/IL-13 and VEGF secretion). CD8+ T-Cells launch apoptosis via the Perforin/Granzyme and Fas/FasL pathways, actively calculating MHC-I antigen presentation rates. However, under chronic antigen exposure, the engine mathematically calculates T-Cell Exhaustion, upregulating TIM-3 and LAG-3 receptors to deplete lymphocyte metabolism. The tumor defends itself by expressing PD-L1 to blind T-Cells and secretes TGF-beta to summon Regulatory T-Cells (Tregs), creating a suppressive stromal aura that physically halts CXCL8 chemotaxis. Concurrently, the Endocrine axis reacts dynamically—cortisol spikes under systemic stress to suppress immunity, while the liver executes the Cori Cycle to convert toxic tumor lactate back into brain-accessible glucose, preventing lethal lactic acidosis.
• The Transhumanist Exploit (Weaponizing Exosomes for Immortality): The core objective of this reactor is not to eradicate cancer, but to hack its neoplastic nature for chronological age-reversal. Human cells biologically age and die due to the progressive shortening of their telomeres after each division (the Hayflick limit). Cancer possesses the unique biological flaw of infinite replication via TP53 inhibition and hTERT (Telomerase) overdrive. By implementing an advanced biological PID controller (Proportional-Integral-Derivative tuning), the engine forces the tumor into a strict Lotka-Volterra Predator-Prey equilibrium with the immune system (The Z-Tumor). Once mathematically stabilized, this cybernetic tumor acts as a biological factory, secreting highly targeted, telomerase-rich exosomes. These exosomes function as biological payloads; they are absorbed by aging healthy cells, physically lengthening their shortened, aged telomeres back to their prime state via localized hTERT mRNA translation and bypassing senescence without initiating Yamanaka OSKM factor-induced dedifferentiation. The algorithm precisely calculates the exact base-pair length required, throttling the exosome secretion to biologically lock the host at exactly 25.0 years old, forever, without triggering uncontrolled systemic neoplasia.
• The "Future is Now" Time-Warp Simulation: Think about the implications of this. I have successfully simulated 1,300 years of biological interactions, cellular mutations, and immune responses in a matter of hours (72,000 computational ticks). The future is already calculated. We can now feed this massive, multi-million-row telemetry log dataset into modern LLMs and Predictive AI models to foresee unknown biological behaviors, discover novel bio-markers, or invent new targeted drugs long before physical biological PID controllers are even created in real-world labs. The Zero Cancer Reactor is the prototype of the future. We can solve tomorrow's transhumanist problems today.
• VRAM Squeezing (Hardware Hack): To fit 70M live biological agents into a single 12GB VRAM buffer without crashing, I strictly locked my Cell struct to exactly 64-bytes (alignas(64)), perfectly matching Ampere's L1/L2 cache lines to prevent memory fragmentation and ensure a 100% cache-hit rate across 3,584 CUDA cores.
• Asynchronous Decoupling: Used cudaStreamCreateWithPriority to decouple the massive CUDA compute thread from my ImGui/DX11 rendering thread. The biological chaos runs at a constant 93% GPU load, but the UI never drops a single frame.
• HPC Architecture & Asynchronous VRAM DMA Transfer (The Compute Engine): To sustain high-throughput calculations across 70M cells, the engine bypasses standard CUDA overhead using advanced High-Performance Computing (HPC) patterns. All telemetry buffers (h_telemetry_pinned_, h_counters_pinned_, and h_success_pinned_) are allocated using page-locked, pinned host memory (cudaMallocHost). This enables direct, zero-copy DMA memory transfers over PCIe via cudaMemcpyAsync, completely bypassing host page-table translation overhead and eliminating PCIe bus bottlenecks. To maximize memory coalescing, block-level parallel reductions are executed within CUDA Shared Memory (__shared__), consolidating warp metrics before committing centralized updates to global device memory using atomic operations (atomicAdd, atomicMin). Finally, thread synchronization is managed via a lock-free, atomic triple-buffering handshake (snapshots_[3] controlled via std::atomic<int> state flags), allowing the heavy DirectX 11 thread to fetch read-only telemetry snapshots asynchronously without stalling the active CUDA compute pipeline.
• The 1,300-Year Log & Proof: In its final stress test, the reactor autonomously processed over 9.2 Billion targeted exosome interactions. It maintained 0.0% systemic inflammation and effectively halted biological aging, proving the mathematical concept of immortality over 1,300 years of real-time homeostasis.

I am completely self-taught. I built this in the dark at 4:30 AM while fighting internet blocks. If you understand the scale of what is happening here, I need you to look at my logic and tell me where I can optimize or what I did wrong.

Here's a 1 minute, 54 second video run of the reactor, showing telemetry, real-time biological phasing, and hardware load:

https://reddit.com/link/1v2pusz/video/37o26v5wbmeh1/player

What you see here is merely a tiny, insignificant fraction of the massive Zero Cancer Reactor project. For the complete, highly detailed biological, mathematical, and scientific data, the raw C++ engine, the 72K-Tick CSV dataset, and my broader cyber-arsenal, visit my GitHub. It is fully open for your research and audit:
🔗 https://github.com/Z-E-7-0-7-R-O/Zero-Ai-Native

Logic is the only law. 💎

I have an onsite round coming up for Google’s GPU performance track.

ML performance, domain knowledge of fundamentals and ML, GPU concepts, techniques and applications are some things I’m focusing on.

Are there any topics in this field which are more important to know or read up on?

Any help is appreciated!

Edit: I am interviewing for L4 role and have 3 yoe total.

A question we get a lot: if you're not using NVIDIA MIG, how do you slice a single GPU across multiple workloads — and how does that work on AMD? Here's the short version.

The MIG tradeoff

MIG (Multi-Instance GPU) partitions a card into isolated instances at the hardware level. It's great for hard isolation, but it's rigid: fixed slice profiles (you pick from a preset menu, not an arbitrary size), supported only on newer top-end data-center cards, reconfiguring usually means a config/firmware change and a node drain + reboot, and it's NVIDIA-only. So if your workload needs \~30% of a card, you round up to the nearest profile and strand the rest.

How we approach it (PodVirt)

Our slicing is software-defined rather than hardware-partitioned. PodVirt sits above the hardware and slices a GPU from 12.5% to 100%: any slice size (not a fixed menu), resized dynamically without reprovisioning the node, working across both NVIDIA and AMD with no vendor SDK lock-in. Each tenant is metered per-minute, so you pay for the slice you actually use.

Because it isn't tied to MIG's firmware path, it runs on a much wider range of hardware — we've tested it across most NVIDIA and AMD GPUs, and even AI PC-class silicon like NVIDIA's GB10.

Why it matters economically

Whole-card rental on long commitments means paying for VRAM you never touch. Sub-card slicing plus per-minute billing turns idle VRAM into usable (and, for datacenters, sellable) capacity. It's the same reason DC operators license the underlying stack to run their own neocloud instead of just renting out whole cards.

Happy to go deeper on the scheduling and isolation side in the comments. And curious — what are you all using today for sub-card utilization: MIG, MPS, time-slicing, or something custom?

Hey everyone,

I’ve been working on a runtime adaptation framework for local models. While traditional fine-tuning locks a model's weights into place, Test-Time Adaptation (TTA) allows the model to adjust dynamically to incoming data streams.

I built TTA-Torch to bring real-time, confidence-gated TTA to LLMs using dynamic LoRA tracking in PyTorch.

Core Highlights:

• Confidence-gating mechanism to determine when adaptation is necessary.

• Low-overhead dynamic LoRA adjustments during the inference pass.

• Clean, modular PyTorch implementation.

I’d love to know if anyone here has experimented with test-time evaluation strategies for local setups, or if you have any feedback on handling memory/kv-cache overhead during dynamic steps!

Repo link: https://github.com/Griffith-7/TTA-Torch

I'm trying to bundle a Python app that uses cudf-polars into a standalone .exe using PyInstaller, but no matter what I do, cudf-polars refuses to come along for the ride.

What I've tried:

• Adding --hidden-import=cudf_polars (and various submodules) to the PyInstaller command
• Adding it to hiddenimports in the .spec file
• Using --collect-all=cudf_polars / --collect-all=cudf / --collect-all=rmm

What happens:
The build completes, but at runtime it either throws ModuleNotFoundError/ImportError on cudf-polars or the GPU engine silently fails to register with Polars

Does anyone face this problem?

I would like to share my latest open source local Unsloth (GGUF) LLM inference engine and applications. It supports many models from Unsloth, like Gemma4, DiffusionGemma, Qwen3.6 with multi-modal (image, vision, audio), Qwen Image Edit, reasoning and function tool. It can run on Windows/MacOS/Linux and fully leverage GPU's capability(Nvidia, Apple, AMD, Intel and others supported by Vulkan, CUDA and Metal). The API is completely compatible with OpenAI and Ollama interface. It has on par performance than llama.cpp Here is the benchmark results in overall:

Performance ratio — TensorSharp vs reference engines

Geomean of TensorSharp's per-scenario speedup over each reference engine on the same backend, across every scenario both engines ran (single-stream, MTP-off). A value > 1.0× means TensorSharp is faster (for decode / prefill throughput) or lower-latency (for TTFT); — = no overlapping cells. Per-scenario ratios are in each model's section below.

This project is not just a C# wrapper of llama.cpp. It implemented the entire LLM inference engine from bottom to top. If you use CPU backend, it's 100% pure C# code execution. Besides CPU backend, I also implmented CUDA, MLX and GGML backend. The GGML backend refer GGML project as external project, and I build a few fusion operation at higher level.

I learned a lot from other projects and apply them for TensorSharp, such as paged KV cache and continuous batching from vLLM, SSD based cache for MoE model from oMLX, GGUF quanztized from llama.cpp and other optimizations for prefill and decode.

Any feedback and comments are welcome. If you like it, it would be really appreciated if you can get this project a star in GitHub. Thanks in advance.

Project Github: GitHub - zhongkaifu/TensorSharp: A native .NET LLM inference engine for GGUF models. TensorSharp provides a console application, a web-based chatbot interface, and Ollama/OpenAI-compatible HTTP APIs for programmatic access. It supports Windows/MacOS/Linux with full GPU capability · GitHub

Space on Huggingface: TensorSharp Chat hosting a Gemma-4 E2B uncensored model (It may be in sleep, so may need to wait for a while to get it waked up)

I work at a social science research institute where we maintain a record linkage library. It connects records across administrative datasets (i.e connecting "Arnold Smitherson" in dataset A with "Arnol Smtherson" in dataset B?). Matching is done with Jaro-Winkler similarity, which scores string pairs on character position and transpositions, so it catches typos, OCR/scanning errors, and spelling variants.

https://github.com/ajwinters/crosswalk

In two datasets with 1 million records, you can't compare all N×M pairs (a trillion comparison), so the pipeline uses blocking by finding exact matches on key fields and removing them from the comparison pool. That still leaves ~20 billion candidate pairs, because pair count grows quadratically.

~95% mapped onto RAPID. cuDF for the blocking joins, cuPy for the Fellegi-Sunter arithmetic. The one thing RAPIDS doesn't have is a string-similarity primitive, so I wrote the Jaro-Winkler kernel with Numba CUDA (one thread per candidate pair, names encoded into a fixed-width uint8 buffer + length array so the kernel can index them).

20B pairs × 7 comparison columns won't fit in 16 GB . So it streams and uploads record data once, then walk blocking keys, predicting each chunk's output size so chunks stay under a VRAM budget, scoring each chunk fully on-device and keeping only links above a threshold. Peak VRAM stayed flat, the memory is bounded by chunk size, not pair count.

Speedup grows with scale small launches starve the SMs (Numba warns about low occupancy), while at 1M throughput climbed to ~84M pairs/sec. For context, this 1M job historically took ~14 hours on a distributed CPU platform; it's now ~4 minutes on my 4070. Different systems, so not a controlled comparison, but that's the practical before/after.

Things for consideration:

• JW is branchy (variable-length strings are divergent loop trip counts within a warp). I'm eating divergence and still winning on raw parallelism. Bucketing pairs by string length so warps stay uniform seems like the obvious next step. Has anyone found that worth it for string kernels?
• The GPU doesn't change the complexity. It's still O(N²), just with a ~200× smaller constant. At 10M records you're back to needing smarter blocking, not faster hardware.

Hello guys. how are you? I kept seeing "MLIR is LLVM for AI" and wanted something more concrete for a CUDA audience. This walks one neural-net layer from high-level tensor IR down to actual PTX: parallel loops, gpu.launch, kernel outlining, the nvvm boundary where gpu.block_id becomes nvvm.read.ptx.sreg.ctaid.x, and finally the NVPTX backend output.

The interesting bit is a redundant st.global.f32 in the inner loop: the accumulator gets stored every k iteration because at that level nothing can prove the output buffer doesn't alias the inputs. Four levels up that was free (tensors can't alias). It is a clean illustration of why the optimization has to happen high.

NVIDIA's CuTe DSL also JITs through MLIR now, which is part of why I think this matters for CUDA folks. Everything is reproducible with the LLVM 20 tools. Would love corrections from people who live in ptxas and SASS.

Hey everyone,

I’ve been working on optimizing the memory overhead that comes with 4-bit inference. I wrote a custom NF4 dequantization kernel using Triton to see if I could eliminate the C++ dispatch bottlenecks found in current baselines.

🚀 Key Results:

• Up to 1.41x speedup compared to the standard bitsandbytes implementation across various tensor shapes.

• Written completely in Python/Triton, making it super easy to inspect, customize, or drop directly into your PyTorch compilation pipelines.

• Passes the Unsloth AI founding engineer challenge requirements (14/14 points).

I'd love to hear the community's feedback, especially if anyone wants to run their own benchmarks on different GPU architectures or suggest further optimization tricks!

Source code & full implementation:

https://github.com/Griffith-7/nf4-triton-kernel

Hi all, first time posting here. I've been building a CUDA profiling tool called GPUFlight and wanted to share it with you all.

Basically you hook it into your GPU workload, either by running your program through the gpufl CLI or embedding it in your code, and it collects all the profiling data while it runs (kernel timeline, occupancy, memory access, SASS, stall reasons, that kind of thing), ships it to a server automatically, and you look at all of it in the browser. The client is open source, and the backend and frontend are the hosted part, which I run as a SaaS.

While I was building the profiler I kept thinking about how to let people try it without installing anything, so I made Performance Lab. It's a small in-browser editor where you write CUDA. You start from a naive kernel they give you, try to optimize it, and if your numbers beat the metrics the problem is asking for, you pass.

If you've got some time, I'd really appreciate it if you gave it a try and left any feedback. Here's the Performance Lab page:
https://gpuflight.com/try/performance-lab
(running a kernel needs a free account)

And the open-source client:
https://github.com/gpu-flight/gpufl-client

Thank you!

I've got some experience in C though and a decent amount in Java, both are rusty (C a little more) but I think I can regain my confidence quick, and I want to start learning GPU programming. CUDA seems like the obvious entry point but I'm not sure what's changed recently or what's actually worth learning first in 2026.

A few questions for anyone who's been down this road:

• Is NVIDIA's own docs still the best starting point, or has something better come along?
• Given my Oop background, should I just dive straight into CUDA C/C++, or is there value in going through something like CUDA Python first to get the fundamentals down before dealing with memory management and kernel launches directly in C? Also curious whether the industry actually takes CUDA Python seriously, or if it's mostly seen as a stepping stone/learning tool rather than something used in production.
• Any free courses, books, or YouTube series that are very relevant you'd recommend?
• Is there a decent low-cost/free way to actually run and test CUDA code without owning an NVIDIA GPU? I have a rtx 5070 so I am just using that right now, set up the toolkit and ran the first adding vectors function lol.
• Are there any textbooks you'd recommend to learn from?

I want to get into the deep learning side as well of course, but also just get to know more since this is new territory, currently I am familiarizing myself with some C++ fundamentals, if you have tips for that also I'd appreciate it.
Thank you! and apologies if this is asked a lot it is just I am really lost, if this is the wrong subreddit for that let me know!

Hey guys I need some help related running open source image generation ai model locally I m lacking the essential hardware. I need a setup of gpu with high vram especially 20-25 gb vram.