Building production-grade LLM inference runtimes in C++/CUDA/PTX and Huawei CANN/Ascend C. Focused on heterogeneous PD disaggregation and H200-class serving infrastructure.

Focused on high-performance LLM inference, GPU/NPU kernel optimization, and heterogeneous distributed runtime design across NVIDIA CUDA and Huawei Ascend CANN ecosystems.

Based in Curitiba, Brazil (GMT-3). Open to remote (US/EU) and on-site (China) engagements. Z-Visa-ready for China relocation.

Production-oriented serving and inference runtime projects.

OpenAI-Compatible Serving Runtime

• Continuous batching, SSE streaming, PagedAttention KV cache management.
• Prometheus metrics, graceful shutdown, health checks.
• 33/33 unit tests + 6/6 end-to-end scenarios passing.

PagedAttention + Scheduling Runtime (Single-GPU Study)

• Custom CUDA implementation of block-based PagedAttention (logical to physical KV mapping).
• 2.89M tok/s kernel microbenchmark (small-head configuration, RTX 2070).
• 29x kernel optimization gain (v1 to v3): warp reduction + online softmax + shared memory tiling + half2 vectorization.
• 205-397x vs naive PyTorch attention (non-Flash baseline, attention-layer only).
• 66.7% memory reduction via block-based KV allocation.
• Includes continuous batching, swap-first preemption, and cost-based victim policy.

Distributed LLM Runtime

• Multi-process runtime implementing Megatron-style Tensor Parallelism and Pipeline Parallelism.
• GPipe-style fill-drain micro-batching.
• KV cache sharding across processes.
• Bit-identical validation on Qwen2-0.5B weights (controlled setup).

Model Compression Research Implementation

• From-scratch AWQ and GPTQ implementations.
• Validated on Qwen2-0.5B with 26.05 perplexity at 4-bit.
• Explores Hessian-based compensation vs activation-aware weighting.

Heterogeneous PD Serving: NVIDIA Blackwell x Huawei Ascend 910B2

• Disaggregated Prefill-Decode serving engine connecting NVIDIA RTX 5060 (Blackwell CC 12.0) and Huawei Ascend 910B2 (DaVinci V300 / CANN 8.6) across a live intercontinental WAN bridge.
• 32 empirical JSON artifacts and 39 high-resolution microarchitectural figures from physical silicon characterization via the OpenI / Peng Cheng CloudBrain 3 research platform.
• Validated end-to-end between Curitiba (Brazil) and Peng Cheng CloudBrain 3 (China): TTFT 35.43 ms, TPOT 0.65 ms/token.
• Empirical characterization of the DaVinci V300 microarchitecture: Fractal-16 flat-latency law, L1 buffer spill at 1024 KB, 32 AI Core scaling at 3.23M tok/s, HBM2e stride collapse (-99.1% at s=256).
• Proved Ascend 910B2 achieves 83.5% silicon efficiency (317.16 / 380 TFLOPS) — the ~60% gap vs NVIDIA H100 is fabrication node (SMIC 7nm DUV vs TSMC 4N EUV), not architecture.
• Advanced algorithm benchmarks: DeepSeek MLA (93% VRAM cut), DeepSeekMoE E=64/K=8 (2.85x), PagedAttention B=16 sweet spot, Zero-Copy CoW Forking (286x), FlashAttention Online Softmax (17 MB at 64k tokens), INT4 quantization (4.00x).
• High-level architecture and telemetry artifacts are open-source; production kernels are maintained separately for enterprise deployment.
• Repository: github.com/JohnScheuer/hybrid-inference-runtime

LayerNorm Kernel: CUDA C++ vs Hand-Written PTX vs Ascend C

• Implemented the same memory-bound LayerNorm kernel across three abstraction layers: CUDA C++ (SM-120), hand-written PTX (SM-120), and Ascend C (CANN 8.x).
• Key finding: Hand-written PTX with ld.global.nc was 4.3% slower than nvcc -O3 auto-vectorized code. Root cause: .nc bypasses L1 cache, which hurts multi-pass kernels (LayerNorm reads data 3 times). This demonstrates the importance of knowing when not to hand-tune.
• Demonstrated the architectural differences between NVIDIA's warp-synchronous model and Huawei's explicit triple-pipe (MTE/Vector/Cube) architecture.
• Repository: github.com/JohnScheuer/cross-arch-kernel-porting

Empirical Proof of WDDM Paging Collapse (up to 1013x)

• Quantified the catastrophic throughput collapse when a GPU's physical VRAM is oversubscribed on Windows. Forced WDDM to page memory to system RAM (and eventually to the SSD) over PCIe Gen4.
• Measured on RTX 5060 (8 GB):
  • Run 1 (RAM available): 68.4× collapse (P99: 3.01 ms → 169.49 ms, TFLOPS: 9.13 → 0.13).
  • Run 2 (RAM under pressure): 1013.1× collapse (P99: 2.41 ms → 2797.32 ms, TFLOPS: 9.17 → 0.01).
• Proved that the penalty is not just the PCIe vs GDDR7 bandwidth ratio (14×). It depends on the entire OS memory hierarchy: when VRAM overflows and RAM is tight, the OS pages to the SSD, turning a memory-bound kernel into a disk-bound kernel.
• Justifies strict physical-VRAM Admission Control in LLM runtimes (vLLM, TGI) before OS-level paging destroys tail latency.
• Repository: github.com/JohnScheuer/wddm-paging-simulator

Edge-Native Inference & Statistical Cost Benchmarking

• Benchmarked FreeToken (edge-native MoE engine from UC Berkeley/MIT) on consumer hardware (RTX 5060 8GB) using Throttle (statistical measurement tool).
• Dense Baseline: Achieved $0.81/M tokens (38% cost reduction via concurrency tuning).
• MoE Stress Test: Served a 21B-parameter MoE (gpt-oss-20b) on an 8GB GPU, rigorously measuring the $19.17/M token cost caused by CPU-GPU offloading bottlenecks.
• Validated that while consumer GPUs can serve dense models cost-effectively, frontier MoE models require ≥16GB VRAM for viable production economics.
• Repository: github.com/JohnScheuer/freetoken-throttle-benchmarks

Full C++/CUDA Runtime

• Transformer inference engine in pure C++/CUDA (no PyTorch execution path).
• 32.4k tok/s peak, achieving 87% of INT8 Tensor Core ceiling on RTX 2070.
• Custom Flash Attention, cuBLASLt INT8, CUDA Graphs, Megatron-style TP.
• Runs TinyStories 110M and TinyLlama 1.1B end-to-end.

PTX-Level GEMM Engineering

• Manual PTX instruction emission for Tensor Core GEMM.
• Outperforms cuBLAS on specific 2048x2048x2048 configuration (SM75 tuned).
• Achieves 60-65% of cuBLAS performance at 8192^3.
• Auto-tuning over WARPS / STAGE_K / ILP_DEPTH configuration space.

INT4 CUDA Kernel + LLM Integration

• Bit-accurate INT4 weight-only quantization kernel using raw PTX assembly.
• 1.000 cosine similarity vs FP16 baseline (layer-level validation).
• End-to-end Qwen2-0.5B integration: -14.7% VRAM usage with coherent outputs.
• Documents SM75 structural limits (no cp.async, no native INT4 MMA operands).

Compiler MVP with CUDA Codegen

• Full pipeline: IR to CUDA codegen to nvcc to GPU execution.
• 9 Transformer fusion patterns (GEMM+Bias+GeLU, LayerNorm fusion, etc.).
• 10.5 TFLOPS peak on RTX 2070 (37% of SM75 theoretical peak).
• Templated CUDA generation with type/shape specialization.

• flash-attention-sm75-register-fused — Register-fused FlashAttention (5.0 TFLOPS peak, roofline-validated).
• flash-attention-sm75 — FlashAttention v1 forward pass in CUDA (SM75).
• fused-backward-kernels-sm75 — Fused GEMM+Bias+GeLU backward kernels (3.1x vs autograd at M=1024).
• speculative-decoding-runtime — 1.41x speedup via draft-model verification.
• npu-performance-model — Cycle-approximate DaVinci V300 microarchitectural emulator (Ascend 910B2).
• vllm-multinode-diagnostics — Multi-node telemetry harness with coordinated-omission-safe latency reporting.

AI Infrastructure Engineer (Contract) — Confidential Leading Chinese LLM Company (Under NDA) | Sep 2025 – Mar 2026

• Engineered production LLM inference infrastructure with Rust/Tokio async pipelines and backpressure-safe concurrency.
• Worked across driver ↔ runtime ↔ framework boundaries to debug cross-layer performance issues (memory, latency, throughput).
• Collaborated with distributed teams on GPU memory profiling and long-context KV cache allocation.

Software Architecture Lead — ZBrasil | Apr 2025 – Aug 2025

• Led backend architecture for high-concurrency, cross-border data processing systems in C++/Python.
• Reduced redundant compute overhead by ~25% via batch coalescing, kernel fusion, and memory restructuring.
• Coordinated distributed teams across UTC-3 / UTC+8 timezones.

Contributed to LLM Inference Handbook (llm-inference-at-scale) — module on observability signal coverage for LLM serving.

LinkedIn: linkedin.com/in/joaofelipescheuer Email: [email protected] WeChat: JohnScheuer7 Location: Curitiba, Brazil — Open to remote (US/EU) and on-site (China)