diff --git a/content/posts/rocm-aiter/images/aiter-architecture.png b/content/posts/rocm-aiter/images/aiter-architecture.png
new file mode 100644
index 0000000..c36452a
Binary files /dev/null and b/content/posts/rocm-aiter/images/aiter-architecture.png differ
diff --git a/content/posts/rocm-aiter/images/aiter-block-diagram.png b/content/posts/rocm-aiter/images/aiter-block-diagram.png
new file mode 100644
index 0000000..add32ff
Binary files /dev/null and b/content/posts/rocm-aiter/images/aiter-block-diagram.png differ
diff --git a/content/posts/rocm-aiter/images/aiter-github-banner.png b/content/posts/rocm-aiter/images/aiter-github-banner.png
new file mode 100644
index 0000000..a043598
Binary files /dev/null and b/content/posts/rocm-aiter/images/aiter-github-banner.png differ
diff --git a/content/posts/rocm-aiter/images/aiter-throughput-chart.png b/content/posts/rocm-aiter/images/aiter-throughput-chart.png
new file mode 100644
index 0000000..f9af25c
Binary files /dev/null and b/content/posts/rocm-aiter/images/aiter-throughput-chart.png differ
diff --git a/content/posts/rocm-aiter/images/backend-abstraction.png b/content/posts/rocm-aiter/images/backend-abstraction.png
new file mode 100644
index 0000000..9c63992
Binary files /dev/null and b/content/posts/rocm-aiter/images/backend-abstraction.png differ
diff --git a/content/posts/rocm-aiter/images/backend-decision.png b/content/posts/rocm-aiter/images/backend-decision.png
new file mode 100644
index 0000000..f28989a
Binary files /dev/null and b/content/posts/rocm-aiter/images/backend-decision.png differ
diff --git a/content/posts/rocm-aiter/images/framework-integration.png b/content/posts/rocm-aiter/images/framework-integration.png
new file mode 100644
index 0000000..aa76335
Binary files /dev/null and b/content/posts/rocm-aiter/images/framework-integration.png differ
diff --git a/content/posts/rocm-aiter/images/jit-pipeline.png b/content/posts/rocm-aiter/images/jit-pipeline.png
new file mode 100644
index 0000000..ebdd21c
Binary files /dev/null and b/content/posts/rocm-aiter/images/jit-pipeline.png differ
diff --git a/content/posts/rocm-aiter/images/mi300x-launch.jpg b/content/posts/rocm-aiter/images/mi300x-launch.jpg
new file mode 100644
index 0000000..f774a4f
Binary files /dev/null and b/content/posts/rocm-aiter/images/mi300x-launch.jpg differ
diff --git a/content/posts/rocm-aiter/images/mi300x-platform.jpg b/content/posts/rocm-aiter/images/mi300x-platform.jpg
new file mode 100644
index 0000000..940be21
Binary files /dev/null and b/content/posts/rocm-aiter/images/mi300x-platform.jpg differ
diff --git a/content/posts/rocm-aiter/images/mla-architecture.png b/content/posts/rocm-aiter/images/mla-architecture.png
new file mode 100644
index 0000000..d1ded03
Binary files /dev/null and b/content/posts/rocm-aiter/images/mla-architecture.png differ
diff --git a/content/posts/rocm-aiter/images/mla-performance-comparison.png b/content/posts/rocm-aiter/images/mla-performance-comparison.png
new file mode 100644
index 0000000..bba1daf
Binary files /dev/null and b/content/posts/rocm-aiter/images/mla-performance-comparison.png differ
diff --git a/content/posts/rocm-aiter/images/mla-speedup-chart.png b/content/posts/rocm-aiter/images/mla-speedup-chart.png
new file mode 100644
index 0000000..604be27
Binary files /dev/null and b/content/posts/rocm-aiter/images/mla-speedup-chart.png differ
diff --git a/content/posts/rocm-aiter/index.en.md b/content/posts/rocm-aiter/index.en.md
new file mode 100644
index 0000000..e331652
--- /dev/null
+++ b/content/posts/rocm-aiter/index.en.md
@@ -0,0 +1,342 @@
+---
+date: '2026-03-17T10:00:00+09:00'
+draft: false
+title: 'AITER Analysis: How AMD Doubled ROCm Inference Performance'
+cover:
+  image: "images/mi300x-launch.jpg"
+  alt: "AMD Instinct MI300X launch event"
+  caption: "AMD Instinct MI300X Launch Event"
+  relative: true
+authors: [Minho Park]
+tags: [AMD, ROCm, AITER, inference, kernel, MI300X, GPU]
+categories: [report]
+summary: "An analysis of AITER (AI Tensor Engine for ROCm), which boosts inference performance on AMD GPUs."
+description: "An analysis of AMD's AITER kernel library that doubled ROCm inference performance. We examine its four kernel backend strategies (Triton, CK, HIP, ASM), JIT compile pipeline, and the architecture that achieved ~2× throughput improvement on DeepSeek R1."
+comments: true
+---
+
+# AITER Analysis: How AMD Doubled ROCm Inference Performance
+
+Hello, I'm Minho Park from the HyperAccel ML team.
+
+Semi Analysis is a well-known semiconductor research firm. They run the [InferenceX](https://inferencex.semianalysis.com) benchmark, which measures and compares inference performance of major GPUs.
+
+According to the [InferenceX v2](https://newsletter.semianalysis.com/p/inferencex-v2-nvidia-blackwell-vs) report released in February 2026, AMD MI300X SGLang throughput improved by **nearly 2×** between December 2025 and January 2026. At the center of this improvement was a kernel library called **AI Tensor Engine for ROCm (AITER)**.
+
+![AMD Instinct MI300X launch event](./images/mi300x-launch.jpg)
+
+That news made me curious. What is AITER, and how could software optimization alone nearly double hardware performance? In this post I summarize AITER’s architecture, kernel backend strategy, and the principles behind the speedup. It may be useful if you:
+
+1. Want to understand the current state of AMD GPU inference performance
+2. Are interested in GPU kernel optimization techniques
+3. Are curious about AI accelerator ecosystems beyond NVIDIA
+
+---
+
+## What is AITER?
+
+![AITER project logo](./images/aiter-github-banner.png)
+
+**AI Tensor Engine for ROCm (AITER)** is AMD’s open high-performance AI operator repository. In short, it is a **collection of kernels** for accelerating AI workloads on AMD GPUs. In the ROCm ecosystem, AITER plays a role analogous to cuDNN in the NVIDIA ecosystem.
+
+| Item | Description |
+| --- | --- |
+| **Repository** | [ROCm/aiter](https://github.com/ROCm/aiter) |
+| **License** | MIT |
+| **Languages** | Python (63.6%), CUDA/HIP (25.8%), C++ (9.7%) |
+| **Supported GPUs** | AMD Instinct MI300X, MI325X, MI350 |
+| **Inference Frameworks** | vLLM, SGLang |
+
+AITER’s main value is **drop-in replacement**. In inference frameworks such as vLLM or SGLang, turning on a single environment variable causes existing operators to be automatically replaced with AITER’s optimized kernels. You get better performance without changing application code.
+
+---
+
+## What Operations Does It Support?
+
+AITER supports a wide range of **Large Language Model (LLM)** inference operations.
+
+| Category | Operations | Backends |
+| --- | --- | --- |
+| **Attention** | Flash Attention, **Multi-head Latent Attention (MLA)**, Paged Attention | CK, ASM, Triton |
+| **Fused MoE** | Top-K routing, MoE sorting, BlockScale FP8 FFN | HIP, CK, ASM |
+| **GEMM** | FP8 per-token/channel, block-scale FP8, INT8, pre-shuffle | CK, ASM |
+| **Normalization** | RMSNorm, LayerNorm (including fused quantization) | Triton, CK |
+| **Embedding** | **Rotary Position Embedding (RoPE)** forward/backward | Triton |
+| **Quantization** | BF16/FP16 → FP8/INT4 conversion | CK, Triton |
+| **Communication** | AllReduce, reduce-scatter, all-gather | Triton, HIP |
+
+Notable here is the backend column: a single operation can have multiple backends (CK, ASM, Triton, etc.). This reflects AITER’s core design philosophy—a **multi-backend strategy**—which we’ll cover in more detail later.
+
+---
+
+## Performance Benchmarks: AITER by the Numbers
+
+Seeing is believing. Here are the performance gains AITER achieves on MI300X.
+
+| Kernel/Workload | Throughput Improvement |
+| --- | --- |
+| Block-scale **General Matrix Multiplication (GEMM)** | **2×** |
+| Block-scale Fused **Mixture of Experts (MoE)** | **3×** |
+| MLA Decode | **17×** |
+| **Multi-Head Attention (MHA)** Prefill | **14×** |
+| DeepSeek V3/R1 throughput (SGLang) | **2×** (6,485 → 13,704 tok/s) |
+| DeepSeek R1 prefill latency | **↓52%** (3.13s → 1.51s) |
+| DeepSeek R1 decode latency | **↓47%** (0.053s → 0.028s) |
+| vs. NVIDIA H200 (DeepSeek R1) | **2–5×** higher throughput at same latency |
+
+![DeepSeek throughput before and after AITER — over 2× improvement (Source: AMD ROCm Blog)](./images/aiter-throughput-chart.png)
+
+MLA Decode 17× and MHA Prefill 14× are not numbers you get from simple tuning. They are possible because AITER writes kernels at the assembly level.
+
+---
+
+## Architecture
+
+### High-Level Architecture
+
+![AITER overall architecture diagram](./images/aiter-architecture.png)
+
+AITER’s architecture is organized into five main layers.
+
+1. **User API**: Python (torch-compatible) and C++ APIs
+2. **High-level operators**: Operator orchestrators such as `tuned_gemm.py`, `fused_moe.py`, `mla.py`
+3. **Operator wrappers**: Wrapping layer based on the `@compile_ops` decorator
+4. **JIT compilation**: Compiles kernels on first use and caches them as `.so` files
+5. **Kernel backends**: Four backends—Triton, **Composable Kernel (CK)**, HIP, and ASM
+
+From the user’s perspective, you call PyTorch-compatible functions via `import aiter`. AITER handles selecting, compiling, and caching the best kernel.
+
+### JIT Compilation Pipeline
+
+![JIT compilation pipeline diagram](./images/jit-pipeline.png)
+
+All AITER operators follow the `@compile_ops` decorator pattern.
+
+```python
+@compile_ops("module_gemm_a8w8", fc_name="gemm_a8w8",
+             gen_func=cmdGenFunc, gen_fake=fake_shape_fn)
+def gemm_a8w8(XQ, WQ, x_scale, w_scale, Out):
+    ...  # Body is not executed — replaced by compiled C++ at runtime
+```
+
+The flow is:
+
+1. The decorated function is called from Python.
+2. The module cache is checked for an already-compiled `.so` file.
+3. If missing, module configuration is loaded from `optCompilerConfig.json`.
+4. The code is compiled with `hipcc` and Ninja to produce a `.so` file.
+5. The module is loaded with `importlib.import_module()` and cached.
+6. Subsequent calls use the cached kernel.
+
+Compilation cost is paid only on first run; afterward you get native-level performance. Eighty-four compile modules are managed this way.
+
+---
+
+## Four Kernel Backends
+
+AITER’s distinctive trait is that it **does not rely on a single kernel language**. It uses four backends, chosen per operator.
+
+![Kernel backend abstraction levels diagram](./images/backend-abstraction.png)
+
+### Triton
+
+Triton is a Python-based GPU programming **Domain-Specific Language (DSL)** developed by OpenAI and ported to ROCm by AMD. It uses a block-level programming model: you write the algorithm at the tile level and the compiler handles optimization.
+
+In AITER, **the largest number of kernels (100+)** are written in Triton. They are used mainly for utility-style operations such as RMSNorm, RoPE, quantization, and MoE sorting.
+
+```python
+# Triton RMSNorm kernel example
+@triton.jit
+def _rms_norm_fwd_kernel(
+    X_ptr, W_ptr, Out_ptr,
+    stride_x_row, N, eps,
+    BLOCK_N: tl.constexpr,
+):
+    row_idx = tl.program_id(0)          # Each program handles one row
+    cols = tl.arange(0, BLOCK_N)
+    mask = cols < N
+
+    x = tl.load(X_ptr + row_idx * stride_x_row + cols, mask=mask, other=0.0)
+    x_sq = x * x
+    mean_sq = tl.sum(x_sq, axis=0) / N
+    rrms = tl.rsqrt(mean_sq + eps)      # reciprocal square root
+
+    w = tl.load(W_ptr + cols, mask=mask)
+    out = x * rrms * w
+    tl.store(Out_ptr + row_idx * stride_x_row + cols, out, mask=mask)
+```
+
+Triton’s strength is development speed: you write GPU kernels in Python, and the compiler handles memory coalescing and shared-memory management. On the other hand, compiler limitations make it hard to reach more than ~95% of theoretical bandwidth.
+
+### Composable Kernel (CK)
+
+CK is AMD’s C++ template-based high-performance kernel library, analogous to NVIDIA’s CUTLASS. As the name “composable kernels” suggests, you build complex operators by combining reusable tile-operation building blocks.
+
+CK’s focus is **operation fusion**. For example, you can fuse scaling of GEMM results inline, without a separate kernel.
+
+```cpp
+// Fuse dequantization inline with GEMM result
+template <typename AccDataType, typename DDataType, typename EDataType>
+struct RowwiseScale {
+    __host__ __device__ constexpr void operator()(
+        EDataType &e, const AccDataType &c,
+        const DDataType &d0,    // weight scale
+        const DDataType &d1     // activation scale
+    ) const {
+        const F32 x = ck::type_convert<F32>(c)
+                    * ck::type_convert<F32>(d0)
+                    * ck::type_convert<F32>(d1);
+        e = ck::type_convert<EDataType>(x);
+    }
+};
+// → GEMM computation + dequantization run in a single kernel
+```
+
+In AITER, CK is used for matrix-centric kernels: GEMM (A8W8, A4W4, block-scale), MoE 2-stage, and some Attention. A heuristic dispatcher chooses among 100+ pre-tuned instances based on M, N, K shapes.
+
+### HIP
+
+**Heterogeneous-computing Interface for Portability (HIP)** is AMD’s C++ runtime API and kernel language. It is the AMD counterpart to CUDA; most CUDA kernel code runs on HIP with little more than name changes.
+
+```text
+CUDA concept            →  HIP equivalent
+─────────────────────────────────────────
+cudaMalloc()             →  hipMalloc()
+cudaMemcpy()             →  hipMemcpy()
+__global__ void kernel   →  __global__ void kernel  (same!)
+__shared__               →  __shared__              (same!)
+cudaStream_t             →  hipStream_t
+nvcc                     →  hipcc
+```
+
+AITER uses HIP for general-purpose operations such as Paged Attention, KV cache, TopK, and AllReduce. Thirty-four `.cu` files under `csrc/kernels/` are written this way.
+
+### Assembly (ASM)
+
+<figure style="background-color: white; padding: 16px; border-radius: 8px;">
+
+![MLA layer structure — key target of AITER’s ASM optimization (Source: AMD ROCm Blog)](./images/mla-architecture.png)
+
+</figure>
+
+The ASM backend consists of kernels written directly in **AMDGCN Instruction Set Architecture (ISA)**, the machine code of AMD GPUs. All compiler abstraction is bypassed; the developer has full control over GPU registers, instruction scheduling, and memory access patterns.
+
+Why go this low? High-level compilers generate safe code but miss optimization opportunities. ASM can fully exploit the GPU’s **Matrix Fused Multiply-Add (MFMA)** instructions, special registers, and instruction pipelining. The result is MLA Decode **17×** and MHA Prefill **14×** improvements.
+
+```text
+Register types:
+  SGPR (Scalar GPR)  — control flow, constants (shared across threads)
+  VGPR (Vector GPR)  — data operations (per-thread values)
+  AccVGPR            — MFMA accumulator (stores matrix multiply result)
+
+Key instructions:
+  v_mfma_f32_32x32x8_f16  — 32×32 matrix multiplication (FP16→FP32)
+  v_mfma_f32_16x16x32_fp8 — 16×16 matrix multiplication (FP8→FP32)
+  buffer_load_dwordx4      — Load 128 bits from global memory
+  ds_read_b128             — Load 128 bits from LDS
+```
+
+AITER’s `hsa/` directory contains **354+ precompiled ASM kernels (.co files)**. So many instances are needed because ASM kernels fix all parameters at compile time. Every viable combination is precompiled with no runtime branching.
+
+```text
+head_dim:  {64, 128, 256}      → 3 options
+dtype:     {fp16, bf16, fp8}   → 3 options
+causal:    {true, false}       → 2 options
+→ Combinatorial explosion → 162+ FMHA instances
+```
+
+### Backend Selection Criteria
+
+![Kernel backend selection decision tree](./images/backend-decision.png)
+
+AITER’s criteria for choosing a backend per operation can be summarized as follows.
+
+| Criterion | Triton | CK | HIP | ASM |
+| --- | --- | --- | --- | --- |
+| **Development speed** | ★★★★★ | ★★☆☆☆ | ★★★☆☆ | ★☆☆☆☆ |
+| **Peak performance** | ★★★☆☆ | ★★★★☆ | ★★★☆☆ | ★★★★★ |
+| **Portability** | ★★★★★ | ★★★☆☆ | ★★★★☆ | ★☆☆☆☆ |
+| **Maintainability** | ★★★★★ | ★★★☆☆ | ★★★★☆ | ★☆☆☆☆ |
+| **Kernel count** | 100+ | ~20 | 34 | 354+ (.co) |
+
+You might ask: “Why not write everything in ASM?” ASM is very hard to develop and maintain, and when the GPU architecture changes, you have to rewrite from scratch. AITER uses ASM only on the hottest paths where performance matters most, and keeps development efficiency with Triton and CK elsewhere.
+
+---
+
+## Framework Integration: Enable with One Environment Variable
+
+![Framework integration flow diagram](./images/framework-integration.png)
+
+AITER integrates without touching existing code. You only need to set environment variables.
+
+```bash
+# Enable AITER in vLLM
+VLLM_USE_AITER_MOE=1 VLLM_USE_AITER_BLOCK_GEMM=1 \
+  vllm serve deepseek-ai/DeepSeek-V3 --tensor-parallel-size 8
+
+# Enable AITER in SGLang
+CK_BLOCK_GEMM=1 SGLANG_ROCM_AITER_BLOCK_MOE=1 \
+  python3 -m sglang.launch_server --model deepseek-ai/DeepSeek-V3 --tp 8
+```
+
+Frameworks read these variables and conditionally dispatch to AITER operators. That makes A/B testing straightforward: turn the variables on or off to compare performance with and without AITER.
+
+On first run, JIT compilation is triggered and `.so` files are created under `~/.cache/aiter/`. Later runs reuse the cached kernels, so there is no extra cost.
+
+---
+
+## Summary of Key Design Patterns
+
+Here are the design patterns that stood out while reviewing AITER.
+
+### Multi-Backend Kernel Dispatch
+
+AITER does not commit to a single kernel language. It uses ASM for decode attention, CK for GEMM, and Triton for MoE sorting—a different approach from a unified-compiler strategy.
+
+### FP8 Block-Scale Quantization
+
+Per-token (1×128) activation scale plus per-weight (128×128) scale enables efficient mixed-precision operations. This is especially effective for MoE architectures like DeepSeek.
+
+### CSV-Based Auto-Tuning
+
+Kernel parameters are stored in CSV files per (M, N, K, cu_num) shape. Tuning per model is possible without recompilation. Pre-tuned settings exist for 21 models including DeepSeek V3, Qwen3, and LLaMA 405B.
+
+### Hardware-Specific Optimizations
+
+![AMD Instinct MI300X 8-GPU OAM platform (Source: ServeTheHome)](./images/mi300x-platform.jpg)
+
+AITER includes hardware-level optimizations such as custom AllReduce tailored to MI300X inter-GPU topology and weight pre-shuffling (16×16 layout) for better memory access patterns.
+
+---
+
+## Conclusion
+
+AITER is AMD’s kernel library to narrow the AI inference performance gap with NVIDIA. Instead of depending on one kernel language, it uses four backends—Triton, CK, HIP, and ASM—and selects the right one per operator.
+
+As noted in Semi Analysis’s InferenceX v2 report, AITER’s optimizations nearly doubled MI300X SGLang performance. At the same time, as Semi Analysis pointed out, individual kernel performance is strong, but when many optimizations—FP4, disaggregated serving, expert parallelism, etc.—are **combined**, composability still lags behind NVIDIA.
+
+![AITER MLA kernel performance improvement by batch size (Source: AMD ROCm Blog)](./images/mla-speedup-chart.png)
+
+What AITER demonstrates is clear: software optimization alone can double performance on the same hardware. For us, working on NPU development, it’s a reminder of how important it is to turn theoretical hardware performance into real-world performance.
+
+---
+
+## References
+
+- [GitHub: ROCm/aiter](https://github.com/ROCm/aiter)
+- [Semi Analysis: InferenceX v2 — NVIDIA Blackwell Vs AMD vs Hopper](https://newsletter.semianalysis.com/p/inferencex-v2-nvidia-blackwell-vs)
+- [Semi Analysis: AMD 2.0 — New Sense of Urgency](https://newsletter.semianalysis.com/p/amd-2-0-new-sense-of-urgency-mi450x-chance-to-beat-nvidia-nvidias-new-moat)
+- [AMD Blog: AITER — AI Tensor Engine For ROCm](https://rocm.blogs.amd.com/software-tools-optimization/aiter-ai-tensor-engine/README.html)
+- [AMD Blog: Accelerate DeepSeek-R1 with AITER + SGLang](https://rocm.blogs.amd.com/artificial-intelligence/aiter-intergration-s/README.html)
+
+---
+
+## P.S.
+
+### HyperAccel Careers
+
+Analyzing AITER drove home how much difference kernel-level optimization can make. At our NPU company we are also building a software stack that can fully utilize hardware performance. With our LPU ASIC launch ahead, we are looking for teammates to work on everything from kernel optimization to inference framework integration.
+
+**Careers**: https://hyperaccel.career.greetinghr.com/ko/guide
+
+If you’re interested, we’d love to hear from you!
diff --git a/content/posts/rocm-aiter/index.ko.md b/content/posts/rocm-aiter/index.ko.md
new file mode 100644
index 0000000..bbb7c83
--- /dev/null
+++ b/content/posts/rocm-aiter/index.ko.md
@@ -0,0 +1,342 @@
+---
+date: '2026-03-17T10:00:00+09:00'
+draft: false
+title: 'AITER 분석: AMD가 ROCm inference 성능을 2배로 올린 방법'
+cover:
+  image: "images/mi300x-launch.jpg"
+  alt: "AMD Instinct MI300X 출시 행사"
+  caption: "AMD Instinct MI300X Launch Event"
+  relative: true
+authors: [Minho Park]
+tags: [AMD, ROCm, AITER, inference, kernel, MI300X, GPU]
+categories: [report]
+summary: "AMD GPU의 inference 성능을 끌어올린 AITER(AI Tensor Engine for ROCm)를 분석합니다."
+description: "AMD가 ROCm inference 성능을 2배로 올린 AITER 커널 라이브러리를 분석합니다. 4가지 커널 백엔드 전략(Triton, CK, HIP, ASM), JIT 컴파일 파이프라인, 그리고 DeepSeek R1에서 2배 throughput 향상을 달성한 구조를 살펴봅니다."
+comments: true
+---
+
+# AITER 분석: AMD가 ROCm inference 성능을 2배로 올린 방법
+
+안녕하세요? HyperAccel ML팀 소속 박민호입니다.
+
+Semi Analysis 는 반도체 업계에서 유명한 리서치 기관입니다. 이 기관은 주요 GPU 의 inference 성능을 실측 비교하는 [InferenceX](https://inferencex.semianalysis.com) 벤치마크를 운영하고 있습니다.
+
+2026년 2월에 공개된 [InferenceX v2](https://newsletter.semianalysis.com/p/inferencex-v2-nvidia-blackwell-vs) 보고서에 따르면, AMD MI300X 의 SGLang throughput 성능이 2025년 12월에서 2026년 1월 사이 **거의 2배 가까이** 향상되었다고 합니다. 이 성능 향상의 중심에 **AI Tensor Engine for ROCm(AITER)** 이라는 커널 라이브러리가 있었습니다.
+
+![AMD Instinct MI300X 출시 행사](./images/mi300x-launch.jpg)
+
+저는 이 소식을 접하고 궁금해졌습니다. 도대체 AITER 가 뭐길래, 소프트웨어 최적화만으로 하드웨어의 성능을 2배나 끌어올릴 수 있었을까? 이번 포스트에서는 AITER 의 아키텍처, 커널 백엔드 전략, 그리고 성능 향상의 원리를 정리해 봤습니다. 아래 목록에 해당되는 분들께 도움이 될 것 같습니다:
+
+1. AMD GPU 의 inference 성능 현황이 궁금하신 분
+2. GPU 커널 최적화 기법에 관심이 있으신 분
+3. NVIDIA 외의 AI 가속기 생태계가 궁금하신 분
+
+---
+
+## AITER란 무엇인가?
+
+![AITER 프로젝트 로고](./images/aiter-github-banner.png)
+
+**AI Tensor Engine for ROCm(AITER)** 는 AMD 가 공개한 고성능 AI operator 통합 저장소입니다. 한 마디로 요약하면 AMD GPU 에서 AI 워크로드를 가속하기 위한 **커널 모음집** 입니다. NVIDIA 생태계에서 cuDNN 이 하는 역할을 ROCm 생태계에서 AITER 가 수행한다고 보면 됩니다.
+
+| 항목 | 내용 |
+| --- | --- |
+| **저장소** | [ROCm/aiter](https://github.com/ROCm/aiter) |
+| **라이선스** | MIT |
+| **언어 구성** | Python (63.6%), CUDA/HIP (25.8%), C++ (9.7%) |
+| **지원 GPU** | AMD Instinct MI300X, MI325X, MI350 |
+| **Inference Framework** | vLLM, SGLang |
+
+AITER 의 핵심 가치는 **드롭인 교체** 에 있습니다. vLLM 이나 SGLang 같은 inference framework에서 환경 변수 하나만 켜면 기존 operator가 AITER 의 최적화된 커널로 자동 교체됩니다. 코드 수정 없이 성능 향상을 얻을 수 있다는 뜻입니다.
+
+---
+
+## 어떤 operation을 지원하는가?
+
+AITER 는 **Large Language Model(LLM)** inference 의 핵심 operation을 폭넓게 지원합니다.
+
+| 카테고리 | operation | 백엔드 |
+| --- | --- | --- |
+| **Attention** | Flash Attention, **Multi-head Latent Attention(MLA)**, Paged Attention | CK, ASM, Triton |
+| **Fused MoE** | Top-K routing, MoE sorting, BlockScale FP8 FFN | HIP, CK, ASM |
+| **GEMM** | FP8 per-token/channel, block-scale FP8, INT8, pre-shuffle | CK, ASM |
+| **Normalization** | RMSNorm, LayerNorm (fused quantization 포함) | Triton, CK |
+| **Embedding** | **Rotary Position Embedding(RoPE)** forward/backward | Triton |
+| **Quantization** | BF16/FP16 → FP8/INT4 변환 | CK, Triton |
+| **Communication** | AllReduce, reduce-scatter, all-gather | Triton, HIP |
+
+여기서 눈에 띄는 것은 백엔드 열입니다. 하나의 operation에도 CK, ASM, Triton 등 여러 백엔드가 병기되어 있습니다. 이것이 AITER 의 핵심 설계 철학인 **멀티 백엔드 전략** 입니다. 이에 대해서는 뒤에서 자세히 다루겠습니다.
+
+---
+
+## 성능 벤치마크: 숫자로 보는 AITER
+
+백문이 불여일견이라고 하죠. AITER 가 MI300X 에서 달성한 성능 향상 수치를 먼저 보겠습니다.
+
+| 커널/워크로드 | Throughput 향상 |
+| --- | --- |
+| Block-scale **General Matrix Multiplication(GEMM)** | **2배** |
+| Block-scale Fused **Mixture of Experts(MoE)** | **3배** |
+| MLA Decode | **17배** |
+| **Multi-Head Attention(MHA)** Prefill | **14배** |
+| DeepSeek V3/R1 throughput (SGLang 기준) | **2배** (6,485 → 13,704 tok/s) |
+| DeepSeek R1 prefill latency | **↓52%** (3.13s → 1.51s) |
+| DeepSeek R1 decode latency | **↓47%** (0.053s → 0.028s) |
+| NVIDIA H200 대비 (DeepSeek R1) | 동일 latency에서 **2-5배** 높은 throughput |
+
+![AITER 적용 전후 DeepSeek throughput 비교 — 2배 이상 향상 (출처: AMD ROCm Blog)](./images/aiter-throughput-chart.png)
+
+MLA Decode 17배, MHA Prefill 14배는 단순한 튜닝으로 나올 수 있는 수치가 아닙니다. AITER 가 어셈블리 수준까지 내려가서 커널을 직접 작성했기 때문에 가능한 결과입니다.
+
+---
+
+## 아키텍처 구조
+
+### 상위 레벨 아키텍처
+
+![AITER 전체 아키텍처 다이어그램](./images/aiter-architecture.png)
+
+AITER 의 아키텍처는 크게 5개 계층으로 구성됩니다.
+
+1. **사용자 API**: Python(torch 호환) 및 C++ API
+2. **상위 레벨 operator**: `tuned_gemm.py`, `fused_moe.py`, `mla.py` 등 operator 오케스트레이터
+3. **operator 래퍼**: `@compile_ops` 데코레이터 기반의 operator 래핑 계층
+4. **JIT 컴파일**: 첫 호출 시 커널을 컴파일하고 `.so` 파일로 캐시
+5. **커널 백엔드**: Triton, **Composable Kernel(CK)**, HIP, ASM 4가지 백엔드
+
+사용자 입장에서는 `import aiter` 로 PyTorch 호환 함수를 호출하면 됩니다. 나머지는 AITER 가 알아서 최적의 커널을 선택하고, 컴파일하고, 캐싱합니다.
+
+### JIT 컴파일 파이프라인
+
+![JIT 컴파일 파이프라인 다이어그램](./images/jit-pipeline.png)
+
+AITER 의 모든 operator는 `@compile_ops` 데코레이터 패턴을 따릅니다.
+
+```python
+@compile_ops("module_gemm_a8w8", fc_name="gemm_a8w8",
+             gen_func=cmdGenFunc, gen_fake=fake_shape_fn)
+def gemm_a8w8(XQ, WQ, x_scale, w_scale, Out):
+    ...  # 본문은 실행되지 않음 — 런타임에 컴파일된 C++로 대체
+```
+
+이 패턴의 동작 원리는 다음과 같습니다.
+
+1. Python 코드에서 데코레이터 함수를 호출합니다
+2. 모듈 캐시에서 이미 컴파일된 `.so` 파일이 있는지 확인합니다
+3. 없으면 `optCompilerConfig.json` 에서 모듈 설정을 로드합니다
+4. `hipcc` 로 컴파일하고 Ninja 빌드 시스템으로 `.so` 를 생성합니다
+5. `importlib.import_module()` 로 로드하여 캐시에 저장합니다
+6. 이후 호출에서는 캐시된 커널을 바로 사용합니다
+
+첫 실행 시에만 컴파일 비용이 발생하고, 이후에는 네이티브 수준의 성능을 얻습니다. 84개의 컴파일 모듈이 이 방식으로 관리됩니다.
+
+---
+
+## 4가지 커널 백엔드
+
+AITER 의 특징은 **하나의 커널 언어를 고집하지 않는다** 는 점입니다. operator 특성에 따라 4가지 백엔드를 골라 씁니다.
+
+![커널 백엔드 추상화 수준 다이어그램](./images/backend-abstraction.png)
+
+### Triton
+
+Triton 은 OpenAI 가 개발하고 AMD 가 ROCm 용으로 포팅한 Python 기반 GPU 프로그래밍 **Domain-Specific Language(DSL)** 입니다. 블록 단위 프로그래밍 모델로, 개발자가 타일 수준에서 알고리즘을 작성하면 컴파일러가 자동으로 최적화합니다.
+
+AITER 에서 **가장 많은 커널(100개 이상)** 이 Triton 으로 작성되었습니다. RMSNorm, RoPE, quantization, MoE sorting 등 유틸리티성 operation에 주로 사용됩니다.
+
+```python
+# Triton RMSNorm 커널 예시
+@triton.jit
+def _rms_norm_fwd_kernel(
+    X_ptr, W_ptr, Out_ptr,
+    stride_x_row, N, eps,
+    BLOCK_N: tl.constexpr,
+):
+    row_idx = tl.program_id(0)          # 각 프로그램이 하나의 행 처리
+    cols = tl.arange(0, BLOCK_N)
+    mask = cols < N
+
+    x = tl.load(X_ptr + row_idx * stride_x_row + cols, mask=mask, other=0.0)
+    x_sq = x * x
+    mean_sq = tl.sum(x_sq, axis=0) / N
+    rrms = tl.rsqrt(mean_sq + eps)      # reciprocal square root
+
+    w = tl.load(W_ptr + cols, mask=mask)
+    out = x * rrms * w
+    tl.store(Out_ptr + row_idx * stride_x_row + cols, out, mask=mask)
+```
+
+Triton 의 강점은 개발 속도입니다. Python 문법으로 GPU 커널을 작성할 수 있고, 메모리 coalescing 이나 shared memory 관리를 컴파일러가 자동 처리합니다. 반면 컴파일러의 한계로 이론 대역폭의 95% 이상에 도달하기는 어렵습니다.
+
+### Composable Kernel(CK)
+
+CK 는 AMD 가 개발한 C++ 템플릿 기반 고성능 커널 라이브러리입니다. NVIDIA 의 CUTLASS 와 유사한 위치에 있습니다. "조합 가능한 커널" 이라는 이름처럼, 재사용 가능한 타일 operation 빌딩 블록을 조합하여 복잡한 operator를 만듭니다.
+
+CK 의 핵심은 **operation fusion** 입니다. 예를 들어 GEMM 결과에 스케일링을 별도 커널 없이 인라인으로 fusion할 수 있습니다.
+
+```cpp
+// GEMM 결과에 dequantization를 인라인으로 fusion
+template <typename AccDataType, typename DDataType, typename EDataType>
+struct RowwiseScale {
+    __host__ __device__ constexpr void operator()(
+        EDataType &e, const AccDataType &c,
+        const DDataType &d0,    // weight 스케일
+        const DDataType &d1     // activation scale
+    ) const {
+        const F32 x = ck::type_convert<F32>(c)
+                    * ck::type_convert<F32>(d0)
+                    * ck::type_convert<F32>(d1);
+        e = ck::type_convert<EDataType>(x);
+    }
+};
+// → GEMM 계산 + dequantization가 하나의 커널에서 수행됩니다
+```
+
+AITER 에서는 GEMM(A8W8, A4W4, block-scale), MoE 2-stage, 일부 Attention 등 matrix operation 중심 커널에 CK 를 사용합니다. M, N, K 형상에 따라 100개 이상의 pre-tuned된 인스턴스 중 최적의 것을 선택하는 휴리스틱 디스패치도 구현되어 있습니다.
+
+### HIP
+
+**Heterogeneous-computing Interface for Portability(HIP)** 는 AMD 의 C++ 런타임 API 이자 커널 언어입니다. CUDA 의 AMD 대응물로, 대부분의 CUDA 커널 코드가 이름만 바꾸면 HIP 에서 동작합니다.
+
+```text
+CUDA 개념              →  HIP 대응물
+─────────────────────────────────────────
+cudaMalloc()           →  hipMalloc()
+cudaMemcpy()           →  hipMemcpy()
+__global__ void kernel →  __global__ void kernel  (동일!)
+__shared__             →  __shared__              (동일!)
+cudaStream_t           →  hipStream_t
+nvcc                   →  hipcc
+```
+
+AITER 에서는 Paged Attention, KV 캐시, TopK, AllReduce 등 general-purpose operation에 HIP 를 사용합니다. `csrc/kernels/` 디렉토리에 34개의 `.cu` 파일이 이 방식으로 작성되어 있습니다.
+
+### Assembly(ASM)
+
+<figure style="background-color: white; padding: 16px; border-radius: 8px;">
+
+![MLA 레이어 구조 — AITER가 ASM으로 최적화한 핵심 대상 (출처: AMD ROCm Blog)](./images/mla-architecture.png)
+
+</figure>
+
+ASM 백엔드는 AMD GPU 의 machine code인 **AMDGCN Instruction Set Architecture(ISA)** 로 직접 작성된 커널입니다. 컴파일러의 모든 추상화를 우회하고 GPU 레지스터, instruction scheduling, 메모리 접근 패턴을 개발자가 100% 제어합니다.
+
+왜 여기까지 내려가야 할까요? 고수준 컴파일러는 안전한 코드를 생성하는 대신 최적화 기회를 놓칩니다. ASM 은 GPU 의 **Matrix Fused Multiply-Add(MFMA)** instruction, 특수 레지스터, instruction pipelining을 100% 활용할 수 있습니다. 그 결과가 MLA Decode **17배**, MHA Prefill **14배** 향상입니다.
+
+```text
+레지스터 타입:
+  SGPR (Scalar GPR)  — control flow, constant (모든 스레드 공유)
+  VGPR (Vector GPR)  — 데이터 operation (각 스레드별 독립 값)
+  AccVGPR            — MFMA accumulator (matrix multiplication 결과 저장)
+
+핵심 instruction:
+  v_mfma_f32_32x32x8_f16  — 32×32 matrix multiplication (FP16→FP32)
+  v_mfma_f32_16x16x32_fp8 — 16×16 matrix multiplication (FP8→FP32)
+  buffer_load_dwordx4      — Global 메모리에서 128비트 로드
+  ds_read_b128             — LDS에서 128비트 로드
+```
+
+AITER 의 `hsa/` 디렉토리에는 **354개 이상의 사전 컴파일된 ASM 커널(.co 파일)** 이 있습니다. 이렇게 많은 인스턴스가 필요한 이유는 ASM 커널은 컴파일 타임에 모든 파라미터가 고정되기 때문입니다. 런타임 분기 없이 모든 가능한 조합을 사전 컴파일합니다.
+
+```text
+head_dim:  {64, 128, 256}      → 3가지
+dtype:     {fp16, bf16, fp8}   → 3가지
+causal:    {true, false}       → 2가지
+→ 조합 폭발 → 162개+ FMHA 인스턴스
+```
+
+### 백엔드 선택 기준
+
+![커널 백엔드 선택 의사결정 트리](./images/backend-decision.png)
+
+AITER 가 각 operation에 특정 백엔드를 선택하는 기준을 정리하면 다음과 같습니다.
+
+| 기준 | Triton | CK | HIP | ASM |
+| --- | --- | --- | --- | --- |
+| **개발 속도** | ★★★★★ | ★★☆☆☆ | ★★★☆☆ | ★☆☆☆☆ |
+| **최대 성능** | ★★★☆☆ | ★★★★☆ | ★★★☆☆ | ★★★★★ |
+| **이식성** | ★★★★★ | ★★★☆☆ | ★★★★☆ | ★☆☆☆☆ |
+| **유지보수** | ★★★★★ | ★★★☆☆ | ★★★★☆ | ★☆☆☆☆ |
+| **커널 수** | 100+ | ~20 | 34 | 354+ (.co) |
+
+"그러면 다 ASM 으로 짜면 되지 않나?" 라고 생각할 수 있는데, ASM 은 개발 난이도가 매우 높고 GPU 아키텍처가 바뀌면 처음부터 다시 작성해야 합니다. AITER 는 성능이 가장 중요한 핫 패스에만 ASM 을 쓰고, 나머지는 Triton 이나 CK 로 개발 효율을 확보합니다.
+
+---
+
+## Framework Integration: 환경 변수 하나로 활성화
+
+![프레임워크 통합 흐름 다이어그램](./images/framework-integration.png)
+
+AITER 는 기존 코드를 건드리지 않고 통합할 수 있습니다. 환경 변수만 설정하면 됩니다.
+
+```bash
+# Enable AITER in vLLM
+VLLM_USE_AITER_MOE=1 VLLM_USE_AITER_BLOCK_GEMM=1 \
+  vllm serve deepseek-ai/DeepSeek-V3 --tensor-parallel-size 8
+
+# Enable AITER in SGLang
+CK_BLOCK_GEMM=1 SGLANG_ROCM_AITER_BLOCK_MOE=1 \
+  python3 -m sglang.launch_server --model deepseek-ai/DeepSeek-V3 --tp 8
+```
+
+프레임워크가 환경 변수를 확인하고 조건부로 AITER operator로 디스패치합니다. 덕분에 A/B 테스트도 간단합니다. 환경 변수를 켜고 끄는 것만으로 AITER 적용 전후 성능을 비교할 수 있습니다.
+
+첫 실행 시에는 JIT 컴파일이 트리거되어 `~/.cache/aiter/` 에 `.so` 파일이 생성됩니다. 이후 실행에서는 캐시된 커널을 재사용하므로 추가 비용이 없습니다.
+
+---
+
+## 핵심 설계 패턴 정리
+
+AITER 를 살펴보면서 눈에 띄었던 설계 패턴들을 정리합니다.
+
+### 멀티 백엔드 커널 디스패치
+
+하나의 커널 언어를 고집하지 않습니다. decode attention 에는 ASM, GEMM 에는 CK, MoE sorting 에는 Triton 을 사용합니다. unified compiler 전략과는 다른 접근입니다.
+
+### FP8 Block-scale Quantization
+
+토큰별 (1x128) activation scale + weight별 (128x128) 스케일로 효율적인 mixed precision operation을 가능하게 합니다. 특히 DeepSeek 과 같은 MoE 아키텍처에서 큰 효과를 발휘합니다.
+
+### CSV 기반 auto-tuning
+
+커널 파라미터가 (M, N, K, cu_num) 형상별로 CSV 파일에 저장됩니다. 재컴파일 없이 모델별 튜닝이 가능합니다. DeepSeek V3, Qwen3, LLaMA 405B 등 21개 모델 설정이 pre-tuned되어 있습니다.
+
+### 하드웨어 특화 최적화
+
+![AMD Instinct MI300X 8-GPU OAM 플랫폼 (출처: ServeTheHome)](./images/mi300x-platform.jpg)
+
+MI300X 인터 GPU 토폴로지에 특화된 커스텀 AllReduce, 최적 메모리 접근 패턴을 위한 weight pre-shuffling(16x16 레이아웃) 등 하드웨어 수준의 최적화가 포함되어 있습니다.
+
+---
+
+## 결론
+
+AITER 는 AMD 가 NVIDIA 와의 AI inference 성능 격차를 줄이기 위해 만든 커널 라이브러리입니다. 단일 커널 언어에 의존하지 않고 Triton, CK, HIP, ASM 4가지 백엔드를 operator별로 골라 쓰는 구조가 특징입니다.
+
+Semi Analysis 의 InferenceX v2 보고서에서도 언급되었듯이, AITER 의 최적화는 MI300X 의 SGLang 성능을 거의 2배로 끌어올렸습니다. 다만 Semi Analysis 가 지적한 것처럼, 개별 커널 성능은 뛰어나지만 FP4, disaggregated serving, expert parallelism 등 여러 최적화를 **동시에 조합** 했을 때의 성능(composability)은 아직 NVIDIA 대비 부족한 부분이 있습니다.
+
+![AITER MLA 커널 적용 시 배치 크기별 성능 향상 (출처: AMD ROCm Blog)](./images/mla-speedup-chart.png)
+
+그래도 AITER 가 보여준 것은 분명합니다. 소프트웨어 최적화만으로 같은 하드웨어에서 2배 성능을 뽑아낼 수 있다는 사실입니다. NPU 를 개발하고 있는 저희 입장에서도, 하드웨어의 이론적 성능을 실제 성능으로 전환하는 것이 얼마나 중요한지 다시 한번 느끼게 되었습니다.
+
+---
+
+## 참고 자료
+
+- [GitHub: ROCm/aiter](https://github.com/ROCm/aiter)
+- [Semi Analysis: InferenceX v2 — NVIDIA Blackwell Vs AMD vs Hopper](https://newsletter.semianalysis.com/p/inferencex-v2-nvidia-blackwell-vs)
+- [Semi Analysis: AMD 2.0 — New Sense of Urgency](https://newsletter.semianalysis.com/p/amd-2-0-new-sense-of-urgency-mi450x-chance-to-beat-nvidia-nvidias-new-moat)
+- [AMD 블로그: AITER — AI Tensor Engine For ROCm](https://rocm.blogs.amd.com/software-tools-optimization/aiter-ai-tensor-engine/README.html)
+- [AMD 블로그: Accelerate DeepSeek-R1 with AITER + SGLang](https://rocm.blogs.amd.com/artificial-intelligence/aiter-intergration-s/README.html)
+
+---
+
+## 추신
+
+### HyperAccel 채용 안내
+
+AITER 를 분석하면서 커널 레벨 최적화가 얼마나 큰 차이를 만드는지 체감했습니다. NPU 회사인 저희도 하드웨어 성능을 제대로 끌어내는 소프트웨어 스택을 만들어가고 있습니다. LPU 의 ASIC 출시를 앞두고, 커널 최적화부터 inference framework integration까지 함께할 팀원을 기다리고 있습니다.
+
+**채용 사이트**: https://hyperaccel.career.greetinghr.com/ko/guide
+
+혹시 관심이 있으시다면 언제든지 연락 주세요!