Skip to content
View danila-permogorskii's full-sized avatar

Block or report danila-permogorskii

Block user

Prevent this user from interacting with your repositories and sending you notifications. Learn more about blocking users.

You must be logged in to block users.

Maximum 250 characters. Please don’t include any personal information such as legal names or email addresses. Markdown is supported. This note will only be visible to you.
Report abuse

Contact GitHub support about this user’s behavior. Learn more about reporting abuse.

Report abuse

Danila Permogorsky

GPU kernel and inference engineer. ROCm / HIP, CDNA matrix cores, SGLang and vLLM in production. AMD Instinct MI300X and MI350X alongside NVIDIA A100 — CUDA-to-ROCm portability is routine work here, not a claim.

Almaty Region, Kazakhstan (UTC+5) · bogdanna.dev · LinkedIn · permogorsky.danila@gmail.com


Upstream

ROCm/rocm-libraries #9627 FP64 16×16×4 MFMA warp GEMM for ck_tile (gfx90a / gfx942 / gfx950). Warp attribute, type aliases, dispatcher specialisations and a unit test against the CPU reference. C-tile lane distribution corrected against the MI300 and MI350 ISAs (§7.1.3.4) after the FP32-derived layout failed on 75% of output elements. Runtime-verified on MI300X.
ROCm/rocm-libraries #10381 Two precision-agnostic gaps in shared ck_tile infrastructure that no 8-byte element type had exercised: missing numeric_traits<double>, and a compile-time divide-by-zero in the universal GEMM pipeline's LDS descriptor at sizeof(ADataType)=8.
ROCm/rocm-libraries #7934 The originating FP64 GEMM feature proposal. Scope of the 16×16×4 MVP agreed directly with an AMD engineer.
sgl-project/sglang #28141 DeepSeek-V4-Flash on ROCm / MI350X still enters a CUDA-linked DeepGEMM dependency during weight post-processing. Reproduced across ten vendor images with full environment.
Solo5 v0.12.1 Listed release contributor.

Own work

mfma — a progressive CDNA kernel curriculum written from scratch on MI300X. Wavefront execution and lane divergence, LDS bank conflicts and cross-lane operations, MFMA register mechanics across VGPR / SGPR / AGPR, a tiled GEMM built on MFMA, XCD awareness on the 8-die package, and separate trees for reading the assembly amdclang++ actually emits. Each module is mapped onto its Composable Kernel equivalent.

hip — HIP and ck_tile experiments on gfx942 and gfx950. Tile distributions, coordinate systems, descriptors and transforms; a toolchain smoke test that measures the HBM bandwidth an SR-IOV virtual function actually delivers rather than the spec-sheet figure. Full dev loop in the open — rocgdb, rocprofv3, rocprof-compute.

inference-ops — 20+ production serving deployments for SGLang, vLLM and llama.cpp across six hardware classes: A100 80GB, MI300X (192 GB HBM3, CDNA3), MI350X (288 GB HBM3E, CDNA4), 4× RTX A4000, and single consumer GPUs with CPU offload. Three-tier observability on AMD.

system-dev — unikernel and kernel work targeting solo5-spt from scratch.

markov-quantum-walks — quantum-walk simulation in Julia on AMDGPU. State-vector simulation is one of the workloads that made the FP64 GEMM gap worth closing.

In production

Two TP=1 SGLang replicas behind a cache-aware router on 2× A100 80GB, serving Qwen3.6-27B for a research institute — data-parallel chosen over TP=2 after establishing the host had no GPU peer-to-peer. Served context raised from 160,000 to 169,000 tokens at zero VRAM cost. Configuration is public: qwen3.6-27b-2-A100.

A 35B mixture-of-experts model served across 4× RTX A4000 for a regional-government platform in Kazakhstan, in state testing since September 2025, entirely inside the customer's perimeter under data-residency law.

Member of the AMD AI Developer Community's highest cohort, with MI350X-class access.


A note on process: I use an LLM assistant for review and environment bootstrapping, running on our own infrastructure. The kernels are mine; where a commit was co-authored, the trailer says so.

Pinned Loading

  1. mfma mfma Public

    Progressive CDNA kernel curriculum on MI300X — wavefronts, LDS bank conflicts, MFMA register mechanics (VGPR/SGPR/AGPR), tiled GEMM, XCD awareness, generated assembly

    C++ 2

  2. hip hip Public

    HIP and ck_tile experiments on MI300X/MI350X — tile distributions, coordinate systems, descriptors and transforms, measured HBM bandwidth on SR-IOV virtual functions

    HIP

  3. bogdannadev/inference-ops bogdannadev/inference-ops Public

    Production LLM serving configurations — SGLang, vLLM and llama.cpp across A100, MI300X, MI350X, RTX A4000 and consumer GPUs, with three-tier observability on AMD

    Python 2

  4. system-dev system-dev Public

    Unikernel and kernel work targeting solo5-spt from scratch

    LLVM

  5. markov-quantum-walks markov-quantum-walks Public

    Quantum-walk simulation in Julia on AMDGPU. State-vector workloads on AMD Instinct hardware

    Julia

  6. quantum-foundations quantum-foundations Public

    Julia