diff --git a/.claude/skills/well-writing/SKILL.md b/.claude/skills/well-writing/SKILL.md new file mode 100644 index 0000000..647fa38 --- /dev/null +++ b/.claude/skills/well-writing/SKILL.md @@ -0,0 +1,230 @@ +--- +name: Well-Writing +description: This skill should be used when reviewing and editing markdown blog posts for the HyperAccel technical blog. It provides style guide compliance checking and readability improvements. +--- + +# Well-Writing + +This skill helps review and edit markdown blog posts for the HyperAccel technical blog by ensuring style guide compliance and improving readability. + +## When to Use This Skill + +Use this skill when: +- Reviewing or editing markdown blog posts (`.md` files) +- Checking style guide compliance +- Improving readability and technical writing quality +- Validating frontmatter and metadata + +## Workflow + +### 1. Load Style Guide Reference + +To review a blog post, first load the style guide: +- Load `references/styleguide.md` for detailed rules and examples +- Load `references/checklist.md` for quick reference during review + +The style guide contains: +- Critical Issues (Must Fix): Bold rendering, acronym usage, capitalization +- Content Quality (High Priority): Readability, technical writing +- Markdown Formatting (Medium Priority): Headings, code blocks, images, lists +- Frontmatter and Metadata: Required fields and format + +### 2. Review Critical Issues + +Check for these critical issues first (must fix before merging): + +**Bold Rendering**: +- Search for pattern: `\*\*[^*]+\*\*[가-힣]` (bold text immediately followed by Korean) +- Ensure space between closing `**` and Korean text +- Example: `**GPU** 는` (correct) vs `**GPU**는` (incorrect) + +**Acronym Usage**: +- Verify all acronyms are introduced with full terms first: `**FullTerm(Acronym)**` +- Check that acronyms in title/summary are introduced in first paragraph +- Common acronyms: GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX, Pallas, Triton + +**Capitalization**: +- Ensure consistent capitalization throughout +- Acronyms uppercase: GPU, TPU, SM +- Common nouns lowercase: loop, thread, warp, core +- Brand names correct: NVIDIA, Google, TensorFlow, PyTorch + +### 3. Review Content Quality + +Check readability and technical writing: + +**Readability**: +- Sentences should be within 50 characters (flexible for technical terms) +- Paragraphs should be within 10 sentences or 200 characters +- Break complex explanations into multiple sentences +- Use blank lines to separate distinct ideas + +**Technical Writing**: +- Technical terms explained when first introduced +- Code examples have comments or explanations +- Images have meaningful alt text +- Content has logical structure with clear sections + +### 4. Review Markdown Formatting + +Check formatting consistency: + +- Headings: Proper hierarchy (H2 → H3, maximum H3 level, no H4 or deeper) +- Code blocks: All have language tags (` ```python`, ` ```bash`, etc.) +- Images: Descriptive alt text, meaningful file names +- Links: Descriptive link text +- Lists: Consistent formatting (`-` for unordered) +- Frontmatter: All required fields present and correct + +### 5. Images from Confluence + +When a blog post references images from Confluence: + +- **Images must be downloaded manually** from the Confluence page +- The agent should inform the user that images need to be manually downloaded +- Provide guidance on where to find the images: + 1. Identify the corresponding Confluence page (if the page ID or URL is known) + 2. Open the Confluence page in a browser + 3. Right-click on each image and select "Save image as..." or "Download" + 4. Save images to the appropriate `images/` directory relative to the markdown file + 5. **Rename images with sequential numbers**: Images should be named with a two-digit prefix (e.g., `01-`, `02-`, `03-`) followed by a descriptive name + - Format: `{number}-{descriptive-name}.{ext}` + - Example: `01-tpuv1-arch.png`, `02-pallas-lowering.png`, `03-tpuv7-arch.png` + - Numbers should reflect the order in which images appear in the blog post + 6. Ensure the saved file names match the image references in the markdown file + +**Note**: Automatic image download from Confluence is not supported. All images must be manually downloaded and placed in the correct directory. + +### 6. Verify and Insert Images in Blog Posts + +When reviewing a blog post, the agent should verify that images are properly referenced and placed: + +1. **Check the images directory**: + - Locate the `images/` directory relative to the markdown file (typically `content/posts/{post-name}/images/`) + - List all image files in the directory + - Note the file names and formats (`.png`, `.jpg`, `.webp`, etc.) + - **Verify naming convention**: All image files should follow the format `{number}-{descriptive-name}.{ext}` where: + - `{number}` is a two-digit sequential number (01, 02, 03, etc.) + - `{descriptive-name}` is a lowercase, hyphen-separated descriptive name + - Example: `01-tpuv1-arch.png`, `02-pallas-lowering.png`, `03-tpuv7-arch.png` + +2. **Read and analyze the markdown file**: + - Search for all image references using the pattern: `!\[.*?\]\(images/.*?\)` + - Extract: + - Image file paths (e.g., `images/tpu-v1-architecture.png`) + - Alt text (the text between `![` and `]`) + - Caption text (if present in alt text or nearby context) + +3. **Match images with references**: + - For each image file in the directory, check if it's referenced in the markdown + - For each image reference in the markdown, verify the file exists + - Identify: + - **Orphaned images**: Files in directory but not referenced in markdown + - **Missing images**: References in markdown but files don't exist + - **Mismatched names**: References that don't match actual file names + +4. **Verify image placement**: + - Check that images are placed in logical positions: + - **Images should be placed BEFORE the content they illustrate** (recommended) + - This allows readers to see the visual context before reading the explanation + - Images should appear near the text that describes them + - Images should not interrupt the flow of reading + - Verify image context: + - The surrounding text should mention or describe what the image shows + - Images should support the narrative, not just be decorative + +5. **Check image metadata**: + - **Alt text**: Should be descriptive and meaningful (not just the filename) + - Good: `![TPU v1 아키텍처 다이어그램](images/01-tpuv1-arch.png)` + - Bad: `![image](images/01-tpuv1-arch.png)` or `![01-tpuv1-arch.png](images/01-tpuv1-arch.png)` + - **File names**: Must follow the sequential numbering format `{number}-{descriptive-name}.{ext}` + - Numbers should be two digits (01, 02, 03, ..., 10, 11, etc.) + - Numbers should reflect the order images appear in the blog post + - Descriptive name should be lowercase with hyphens (e.g., `tpuv1-arch`, `pallas-lowering`) + - **Image format**: Prefer `.webp` for smaller file sizes, but `.png` is acceptable for diagrams + +6. **Suggest improvements**: + - If images are missing, suggest adding them with appropriate alt text and sequential numbering + - If images are in wrong locations, suggest moving them to better positions + - If alt text is missing or poor, suggest improvements + - If file names don't follow the numbering convention, suggest renaming: + - Determine the correct order based on where images appear in the markdown + - Rename files to follow `{number}-{descriptive-name}.{ext}` format + - Update all references in the markdown file to match the new names + - If file names don't match references, suggest renaming files or updating references + - If orphaned images exist, ask if they should be removed or referenced + - If sequential numbers are missing or incorrect, suggest renumbering all images in order + +7. **When inserting images**: + - **Place images BEFORE the content they illustrate** (recommended) + - This allows readers to see the visual context before reading the explanation + - Use descriptive alt text that explains what the image shows + - Consider the reading flow - images should enhance understanding, not interrupt it + - **File naming**: Before inserting, ensure the image file follows the sequential numbering format: + - Determine the correct sequence number based on where this image appears in the post + - If it's the first image, use `01-`, second image use `02-`, etc. + - Format: `{number}-{descriptive-name}.{ext}` (e.g., `01-tpuv1-arch.png`) + - Format: `![Descriptive alt text](images/{number}-{descriptive-name}.{ext})` + - If the image needs a caption, include it in the alt text or add a caption below + - **Renaming existing images**: If inserting an image requires renumbering existing ones: + - Rename all affected image files + - Update all references in the markdown file to match the new names + - Maintain sequential order throughout the post + +**Example workflow**: +``` +1. Read markdown file: content/posts/pallas-programming-model/index.md +2. Check images directory: content/posts/pallas-programming-model/images/ +3. Find all image references in markdown (in order of appearance) +4. Verify all image files follow naming convention: {number}-{descriptive-name}.{ext} +5. Match files with references +6. Verify each image is in the right place contextually +7. Check alt text quality +8. Verify sequential numbering matches the order images appear in the post +9. Report any issues and suggest fixes (renaming, renumbering, etc.) +``` + +**Image naming examples**: +- First image in post: `01-tpuv1-arch.png` +- Second image in post: `02-pallas-lowering.png` +- Third image in post: `03-tpuv7-arch.png` +- Tenth image in post: `10-common-hw-model.png` + +## Providing Feedback + +When reviewing a post: + +1. **Prioritize by severity**: + - Critical: Must fix (rendering issues, major readability problems) + - High: Significantly impacts readability or understanding + - Medium: Formatting or consistency issues + +2. **Be specific**: + - Point to exact locations (line numbers if possible) + - Provide before/after examples + - Reference specific rules from `references/styleguide.md` + +3. **Suggest fixes**: + - Show correct format + - Explain why the fix is needed + - Reference relevant sections of the style guide + +4. **Use the checklist**: + - Refer to `references/checklist.md` for quick verification + - Check off items as you review + - Focus on critical and high-priority items first + +## Common Patterns + +Based on existing blog posts: + +- **Series posts**: Use `series: ["시리즈 이름"]` in frontmatter +- **Code languages**: Python, CUDA, Bash, C, Go, F# are commonly used +- **Image formats**: `.webp`, `.png`, `.jpg` are used +- **Technical terms**: GPU, TPU, SM, CUDA, HBM, Systolic Array, Tensor Core, etc. + +## Notes + +- **Context matters**: Some rules may be flexible for technical content (e.g., longer sentences for complex explanations) +- **Consistency first**: If a pattern is used consistently throughout a post, it may be acceptable even if not ideal +- **Be constructive**: When pointing out issues, suggest specific fixes and explain the reasoning diff --git a/.claude/skills/well-writing/references/checklist.md b/.claude/skills/well-writing/references/checklist.md new file mode 100644 index 0000000..8bb0047 --- /dev/null +++ b/.claude/skills/well-writing/references/checklist.md @@ -0,0 +1,84 @@ +# Quick Reference Checklist + +This is a quick reference checklist for reviewing blog posts. For detailed explanations, see `styleguide.md`. + +--- + +## Critical Issues (Must Fix) + +### Bold Rendering +- [ ] No `**text**한글` patterns without spaces +- [ ] Search pattern: `\*\*[^*]+\*\*[가-힣]` +- [ ] All bold markers followed by Korean text have a space + +### Acronym Usage +- [ ] All acronyms introduced with full terms first: `**FullTerm(Acronym)**` +- [ ] Acronyms in title/summary are introduced in first paragraph +- [ ] Only full term is bolded, not the acronym +- [ ] Common acronyms: GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX, Pallas, Triton + +### Capitalization +- [ ] Consistent capitalization throughout document +- [ ] Acronyms are uppercase: GPU, TPU, SM, CUDA +- [ ] Common nouns are lowercase: loop, thread, warp, core +- [ ] Brand names correct: NVIDIA, Google, TensorFlow, PyTorch + +--- + +## High Priority + +### Readability +- [ ] Sentences within 50 characters (flexible for technical terms) +- [ ] Paragraphs within 10 sentences +- [ ] Paragraphs within 200 characters +- [ ] Complex explanations broken into multiple sentences +- [ ] Blank lines separate distinct ideas + +### Technical Writing +- [ ] Technical terms explained when first introduced +- [ ] Code examples have comments or explanations +- [ ] Images have meaningful alt text +- [ ] Content has logical structure with clear sections +- [ ] Transitions connect sections smoothly + +--- + +## Medium Priority + +### Markdown Formatting +- [ ] H1 only for page title +- [ ] Heading hierarchy not skipped (H2 → H3 → H4) +- [ ] All code blocks have language tags +- [ ] Inline code uses single backticks +- [ ] List markers are consistent (`-` for unordered) +- [ ] Images have descriptive alt text +- [ ] Link text is descriptive + +### Frontmatter +- [ ] All required fields present +- [ ] Date format: `'2026-01-03T17:20:16+09:00'` +- [ ] Authors match entries in `content/authors/` +- [ ] Tags are relevant and consistent +- [ ] Summary is concise and descriptive + +--- + +## Quick Search Patterns + +Use these patterns to find common issues: + +- **Bold + Korean**: `\*\*[^*]+\*\*[가-힣]` +- **Acronyms**: Search for common acronyms (GPU, TPU, SM, etc.) and verify first use +- **Capitalization**: Search for variations (Loop/loop, GPU/gpu) +- **Long sentences**: Look for sentences with multiple commas or conjunctions +- **Code blocks**: Search for ` ``` ` without language tags + +--- + +## Common Acronyms Reference + +**Must follow FullTerm(Acronym) rule**: +- GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX, Pallas, Triton, VLIW, DMA, ALU, SIMD, SPMD + +**May skip full form**: +- CPU, AI, ML, API, URL, HTTP, HTTPS diff --git a/.claude/skills/well-writing/references/styleguide.md b/.claude/skills/well-writing/references/styleguide.md new file mode 100644 index 0000000..4ff268a --- /dev/null +++ b/.claude/skills/well-writing/references/styleguide.md @@ -0,0 +1,553 @@ +# HyperAccel Blog Style Guide + +This document defines style rules for writing markdown posts in the HyperAccel technical blog. Follow these rules to ensure consistency, readability, and proper rendering across all blog posts. + +--- + +## Table of Contents + +1. [Critical Issues (Must Fix)](#1-critical-issues-must-fix) + - [Markdown Bold Rendering](#11-markdown-bold-rendering) + - [Acronym Usage](#12-acronym-usage) + - [Term Capitalization Consistency](#13-term-capitalization-consistency) +2. [Content Quality (High Priority)](#2-content-quality-high-priority) + - [Readability](#21-readability) + - [Technical Writing](#22-technical-writing) +3. [Markdown Formatting (Medium Priority)](#3-markdown-formatting-medium-priority) + - [Headings](#31-headings) + - [Code Blocks](#32-code-blocks) + - [Images and Links](#33-images-and-links) + - [Lists and Quotes](#34-lists-and-quotes) +4. [Frontmatter and Metadata](#4-frontmatter-and-metadata) +5. [Complete Examples](#5-complete-examples) +6. [Review Checklist](#6-review-checklist) + +--- + +## 1. Critical Issues (Must Fix) + +These issues cause rendering problems or significantly impact readability. They must be fixed before merging. + +### 1.1 Markdown Bold Rendering + +**Problem**: When Korean text immediately follows bold text, markdown rendering may break, causing asterisks to appear as literal characters in the rendered output. + +**Rules**: +- Always add a space when Korean particles (은, 는, 이, 가, 을, 를, 으로, 로, 의, 에, 에서, 와, 과, 등) immediately follow bold markers `**` +- Pay special attention when Korean text follows bold text ending with parentheses +- This applies to all bold text, not just acronyms +- The space should be between the closing `**` and the Korean text + +**Incorrect Examples**: +```markdown +**SM(Streaming Multiprocessor)**으로 구성되어 있습니다. +**TPU(Tensor Processing Unit)**는 이와 대조적입니다. +**대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)**하는 방식입니다. +**GPU**는 병렬 연산에 최적화되어 있습니다. +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)** 으로 구성되어 있습니다. +**Tensor Processing Unit(TPU)** 는 이와 대조적입니다. +**대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)** 하는 방식입니다. +**GPU** 는 병렬 연산에 최적화되어 있습니다. +``` + +**How to Check**: +- Search for pattern: `\*\*[^*]+\*\*[가-힣]` (bold text immediately followed by Korean) +- Visual inspection: Render the markdown and check if asterisks appear as literal characters + +**Checkpoints**: +- [ ] No pattern like `**text(acronym)**한글` exists +- [ ] No pattern like `**text**한글` exists +- [ ] All bold markers followed by Korean text have a space +- [ ] Rendered output shows no literal asterisks + +--- + +### 1.2 Acronym Usage + +**Problem**: Using acronyms before introducing the full term confuses readers. Incorrect formatting makes terms hard to understand and breaks the reading flow. + +**Rules**: +- **First use**: Always write `FullTerm(Acronym)` format +- **Subsequent use**: Only the acronym can be used +- **Bold formatting**: Bold only the full term, not the acronym: `**FullTerm(Acronym)**` +- If an acronym appears in the title or summary, introduce it in the first paragraph +- For well-known acronyms (e.g., CPU, GPU, AI), you may skip the full form if it's extremely common + +**Incorrect Examples**: +```markdown +SM은 GPU의 핵심 구성 요소입니다. (Acronym used before full term) +**SM(Streaming Multiprocessor)**으로 구성되어 있습니다. (Acronym is bolded) +이 글에서는 TPU에 대해 설명합니다. (TPU not introduced) +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 구성 요소입니다. +이후 문장에서는 SM으로 지칭할 수 있습니다. + +이 글에서는 **Tensor Processing Unit(TPU)**에 대해 설명합니다. +TPU는 구글이 개발한 AI 가속기입니다. +``` + +**Common Acronyms in This Blog** (must follow the rule): +- GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX, Pallas, Triton, VLIW, DMA, ALU, SIMD, SPMD + +**Well-Known Acronyms** (may skip full form): +- CPU, AI, ML, API, URL, HTTP, HTTPS + +**How to Check**: +- Search for acronyms in the document +- Verify that each acronym appears with its full term before standalone use +- Check that acronyms in title/summary are introduced in the first paragraph + +**Checkpoints**: +- [ ] Acronym is presented with full term when it first appears +- [ ] Acronym does not appear before the full term +- [ ] Only the full term is bolded, not the acronym +- [ ] Acronyms in title/summary are introduced in first paragraph + +--- + +### 1.3 Term Capitalization Consistency + +**Problem**: Inconsistent capitalization of the same term throughout a document confuses readers and looks unprofessional. + +**Rules**: +- Decide capitalization when a technical term first appears, then maintain it consistently throughout the document +- Acronyms should be all uppercase (e.g., GPU, TPU, SM, CUDA, HBM, SMEM) +- Common nouns should be lowercase (e.g., loop, thread, warp, core, unit, array) +- Proper nouns and brand names follow their official capitalization (e.g., NVIDIA, Google, TensorFlow, PyTorch) +- When in doubt, check official documentation or websites + +**Incorrect Examples**: +```markdown +이 문서에서는 Loop와 loop를 혼용하고 있습니다. +GPU와 gpu가 섞여 있습니다. +Tensorflow와 TensorFlow가 혼용됩니다. +``` + +**Correct Examples**: +```markdown +이 문서에서는 loop를 사용합니다. +GPU는 Graphics Processing Unit의 약자입니다. +TensorFlow는 구글이 개발한 프레임워크입니다. +``` + +**Common Terms and Their Correct Capitalization**: +- **Lowercase**: loop, thread, warp, core, unit, array, buffer, memory, register +- **Uppercase (Acronyms)**: GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX +- **Brand Names**: NVIDIA, Google, TensorFlow, PyTorch, JAX, Pallas, Triton + +**How to Check**: +- Search for variations of the same term (e.g., "Loop", "loop", "LOOP") +- Verify consistency throughout the document +- Check brand names against official sources + +**Checkpoints**: +- [ ] Same term uses consistent capitalization within the document +- [ ] All acronyms are unified in uppercase +- [ ] Brand names use correct capitalization +- [ ] Common nouns are lowercase + +--- + +## 2. Content Quality (High Priority) + +These rules ensure blog posts are readable, well-structured, and technically accurate. + +### 2.1 Readability + +**Problem**: Long sentences and paragraphs make technical content hard to follow. Readers lose focus and miss important information. + +**Rules**: + +**Sentence Length**: +- A sentence should be **within 50 characters** (recommended, flexible for technical terms) +- Break complex explanations into multiple sentences +- Split long sentences connected with commas (`,`) or conjunctions (`그리고`, `또한`, `또한`) +- Use periods (`.`) more liberally to create shorter, clearer sentences +- One idea per sentence when possible + +**Paragraph Length**: +- A paragraph should be approximately **5-7 sentences** +- Consider splitting if a paragraph has **10 or more sentences** +- Consider splitting if a paragraph exceeds **200 characters** +- Use blank lines to separate distinct ideas +- Each paragraph should focus on one main idea + +**Incorrect Example**: +```markdown +TPU는 AI 연산에 특화된 하드웨어로, 구글이 2016년부터 개발해온 것으로, 행렬 곱셈 연산에 최적화되어 있으며, Systolic Array라는 특별한 연산 유닛을 사용하며, 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**Correct Example**: +```markdown +TPU는 AI 연산에 특화된 하드웨어입니다. 구글이 2016년부터 개발해온 TPU는 행렬 곱셈 연산에 최적화되어 있습니다. + +TPU는 Systolic Array라는 특별한 연산 유닛을 사용합니다. 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**How to Check**: +- Count characters in each sentence (approximate) +- Count sentences in each paragraph +- Look for sentences with multiple commas or conjunctions +- Check for paragraphs that span multiple screen heights + +**Checkpoints**: +- [ ] No sentence exceeds 50 characters (flexible for technical terms) +- [ ] No paragraph has 10 or more sentences +- [ ] No paragraph exceeds 200 characters +- [ ] Complex explanations are broken into multiple sentences +- [ ] Blank lines separate distinct ideas + +--- + +### 2.2 Technical Writing + +**Problem**: Technical content without proper explanations, context, or structure is hard to understand. + +**Rules**: + +**Term Definitions**: +- Briefly explain technical terms when they first appear +- Provide context: why is this term important? How does it relate to the topic? +- Follow the "Acronym Usage Rules" above for acronyms +- Use examples or analogies when helpful + +**Code Examples**: +- Use code examples that actually work (test them if possible) +- Add explanations through comments in the code +- Explain complex code step by step before or after the code block +- Use appropriate language tags for syntax highlighting +- If code is pseudocode or conceptual, mention it in the text + +**Diagrams and Images**: +- Use diagrams to explain complex concepts (architecture, flow, relationships) +- Add meaningful alt text that describes the content, not just "image" or "diagram" +- Use captions when helpful for context +- Name image files meaningfully (e.g., `tpu-architecture.png`, not `image1.png`) + +**Structure**: +- Start with an introduction that sets context +- Use headings to organize content logically +- End with a summary or conclusion +- Use transitions between sections +- Guide readers from general to specific concepts + +**Incorrect Examples**: +```markdown +SM은 GPU의 구성 요소입니다. (No explanation of what SM does or why it matters) + +```python +def compute(x, y): + return x * y +``` (No explanation of what this code does or why it's relevant) +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 연산 단위입니다. 각 SM은 수백 개의 CUDA 코어를 포함하며, 독립적으로 스레드 블록을 실행합니다. 이는 GPU의 병렬 처리 능력의 기반이 됩니다. + +다음 코드는 행렬 곱셈을 수행합니다: + +```python +def matrix_multiply(A, B): + """두 행렬 A와 B를 곱합니다.""" + return A @ B # NumPy의 행렬 곱셈 연산자 사용 +``` + +이 함수는 NumPy의 최적화된 행렬 곱셈을 사용하여 효율적으로 연산을 수행합니다. +``` + +**Checkpoints**: +- [ ] Technical terms are explained when first introduced +- [ ] Code examples have comments or explanations +- [ ] Images have meaningful alt text +- [ ] Content has logical structure with clear sections +- [ ] Transitions connect sections smoothly + +--- + +## 3. Markdown Formatting (Medium Priority) + +These rules ensure proper markdown rendering and consistent formatting. + +### 3.1 Headings + +**Rules**: +- Use H1(`#`) only for the page title (usually in frontmatter, not in body) +- Start body sections with H2(`##`) +- Use H3(`###`) for subsections +- Use H4(`####`) for sub-subsections if needed +- Do not skip heading hierarchy (e.g., do not use H4 after H2) +- Use descriptive headings that summarize the section content +- Headings should form a logical outline of the content + +**Incorrect Examples**: +```markdown +# Main Title +## Section 1 +#### Subsection (Skipped H3) +``` + +**Correct Examples**: +```markdown +## TPU Architecture +### Systolic Array +#### Data Flow +``` + +**Checkpoints**: +- [ ] H1 is only used for page title +- [ ] Heading hierarchy is not skipped +- [ ] Headings are descriptive +- [ ] Headings form a logical outline + +--- + +### 3.2 Code Blocks + +**Rules**: +- Code blocks must specify language tags: ` ```python`, ` ```bash`, ` ```cuda`, etc. +- Use single backticks for inline code: `` `variable_name` `` +- For file paths or commands, use inline code: `` `hugo server` `` +- Add comments in code to explain complex logic +- If code is pseudocode or conceptual, mention it in the text +- Format code properly (indentation, spacing) + +**Incorrect Examples**: +```markdown +``` +def hello(): + print("world") +``` +(No language tag) +``` + +**Correct Examples**: +```python +def hello(): + """Print hello world.""" + print("world") +``` + +**Common Languages in This Blog**: +- `python`, `cuda`, `bash`, `c`, `go`, `fsharp`, `javascript`, `yaml`, `toml` + +**Checkpoints**: +- [ ] All code blocks have language tags +- [ ] Inline code uses single backticks +- [ ] Complex code has comments or explanations +- [ ] Code is properly formatted + +--- + +### 3.3 Images and Links + +**Rules**: + +**Images**: +- Always include alt text: `![descriptive alt text](path/to/image.png)` +- Alt text should describe the content, not just say "image" or "diagram" +- Use relative paths for images in the same post directory +- Use meaningful file names (e.g., `tpu-architecture.png`, not `img1.png`) +- Prefer `.webp` format for better compression + +**Links**: +- Use descriptive link text: `[descriptive text](URL)` +- Avoid generic text like "click here" or "link" +- For external links, consider adding `{:target="_blank"}` if needed (Hugo-specific) +- Link to authoritative sources when possible + +**Incorrect Examples**: +```markdown +![image](tpu.png) (No descriptive alt text) +[여기](https://example.com)를 클릭하세요. (Generic link text) +``` + +**Correct Examples**: +```markdown +![TPU architecture diagram showing Systolic Array](tpu-architecture.png) +[TPU 논문](https://arxiv.org/pdf/1704.04760)을 참고하세요. +``` + +**Checkpoints**: +- [ ] All images have descriptive alt text +- [ ] Link text is descriptive +- [ ] Image file names are meaningful +- [ ] Images use appropriate formats + +--- + +### 3.4 Lists and Quotes + +**Rules**: + +**Lists**: +- Use `-` for unordered lists (consistent throughout) +- Use `1.` for ordered lists when order matters +- Maintain consistency between list items (same structure, same level of detail) +- Nested lists are recommended up to 2 levels +- Add blank lines before and after lists for readability +- Use parallel structure in list items + +**Block Quotes**: +- Use `>` for block quotes +- Use quotes for important statements, definitions, or callouts +- Keep quotes concise +- Use quotes to highlight key concepts or definitions + +**Incorrect Examples**: +```markdown +- Item 1 +* Item 2 (Inconsistent list markers) +``` + +**Correct Examples**: +```markdown +- Item 1 +- Item 2 + - Sub-item 2.1 + - Sub-item 2.2 +``` + +**Checkpoints**: +- [ ] List markers are consistent +- [ ] List items have consistent structure +- [ ] Quotes are used appropriately +- [ ] Lists have proper spacing + +--- + +## 4. Frontmatter and Metadata + +**Rules**: +- All posts must have frontmatter with required fields +- `date` should be in ISO 8601 format: `'2026-01-03T17:20:16+09:00'` +- `draft: false` for published posts +- `title` should be descriptive and match the H1 heading +- `tags` should be relevant and consistent (use existing tags when possible) +- `authors` must match entries in `content/authors/` +- `summary` should be a concise one-line description +- `description` can be longer (used for SEO) +- For series posts, include `series: ["시리즈 이름"]` +- `cover` image is optional but recommended + +**Required Fields**: +```yaml +--- +date: '2026-01-03T17:20:16+09:00' +draft: false +title: 'Post Title' +authors: ["Author Name"] +tags: ["Tag1", "Tag2"] +categories: ["Category"] +summary: "One-line summary" +description: "Longer description for SEO" +--- +``` + +**Optional Fields**: +```yaml +series: ["Series Name"] +cover: + image: "path/to/image.png" + alt: "Image description" + caption: "Image caption" + relative: true +comments: true +``` + +**Checkpoints**: +- [ ] All required frontmatter fields are present +- [ ] Date format is correct +- [ ] Authors match existing author entries +- [ ] Tags are relevant and consistent +- [ ] Summary is concise and descriptive + +--- + +## 5. Complete Examples + +### 5.1 Before and After: Acronym and Bold Rendering + +**Before (Incorrect)**: +```markdown +SM은 GPU의 핵심 구성 요소입니다. **SM(Streaming Multiprocessor)**으로 구성되어 있으며, 각 SM은 수백 개의 CUDA 코어를 포함합니다. +``` + +**After (Correct)**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 구성 요소입니다. GPU는 수백 개의 SM 으로 구성되어 있으며, 각 SM은 수백 개의 CUDA 코어를 포함합니다. +``` + +### 5.2 Before and After: Readability + +**Before (Incorrect)**: +```markdown +TPU는 AI 연산에 특화된 하드웨어로, 구글이 2016년부터 개발해온 것으로, 행렬 곱셈 연산에 최적화되어 있으며, Systolic Array라는 특별한 연산 유닛을 사용하며, 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**After (Correct)**: +```markdown +TPU는 AI 연산에 특화된 하드웨어입니다. 구글이 2016년부터 개발해온 TPU는 행렬 곱셈 연산에 최적화되어 있습니다. + +TPU는 Systolic Array라는 특별한 연산 유닛을 사용합니다. 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +--- + +## 6. Review Checklist + +When reviewing a blog post, check the following items in order of priority: + +### Critical (Must Fix) +- [ ] **Bold rendering**: No `**text**한글` patterns without spaces +- [ ] **Acronyms**: All acronyms introduced with full terms first +- [ ] **Capitalization**: Consistent capitalization throughout + +### High Priority +- [ ] **Readability**: Sentences within 50 characters, paragraphs within 10 sentences +- [ ] **Technical terms**: All technical terms explained when first introduced +- [ ] **Code examples**: All code blocks have language tags and explanations +- [ ] **Images**: All images have descriptive alt text + +### Medium Priority +- [ ] **Headings**: Proper hierarchy (H2 → H3 → H4, no skipping) +- [ ] **Links**: Descriptive link text +- [ ] **Lists**: Consistent formatting +- [ ] **Frontmatter**: All required fields present and correct + +### Additional Checks +- [ ] Content flows logically with clear transitions +- [ ] No typos or grammatical errors +- [ ] Examples are accurate and tested (if applicable) +- [ ] Images are properly placed and referenced + +--- + +## Notes for Reviewers + +- **Severity Levels**: + - **Critical**: Causes rendering issues or major readability problems + - **High**: Significantly impacts readability or understanding + - **Medium**: Formatting or consistency issues + +- **Be Constructive**: When pointing out issues, suggest specific fixes +- **Context Matters**: Some rules may be flexible for technical content (e.g., longer sentences for complex explanations) +- **Consistency First**: If a pattern is used consistently throughout a post, it may be acceptable even if not ideal + +--- + +## Common Patterns in This Blog + +Based on existing posts, here are common patterns to follow: + +- **Series Posts**: Use series tag for series like `series: ["지피지기면 백전불태"]` +- **Code Languages**: Python, CUDA, Bash, C, Go, F# are commonly used +- **Image Formats**: `.webp`, `.png`, `.jpg` are used +- **Technical Terms**: GPU, TPU, SM, CUDA, HBM, Systolic Array, Tensor Core, etc. diff --git a/.gemini/config.yaml b/.gemini/config.yaml new file mode 100644 index 0000000..95f8a9d --- /dev/null +++ b/.gemini/config.yaml @@ -0,0 +1,34 @@ +# Gemini Code Assist Configuration +# 블로그 포스트 리뷰를 위한 설정 + +have_fun: true # 블로그 글은 비공식적(유머 등)이어도 좋습니다. + +memory_config: + disabled: false + +# 리뷰에서 제외할 파일/디렉토리 패턴 +ignore_patterns: + - "themes/**" # 테마 파일은 제외 + - "*.lock" # 빌드 락 파일 + - ".git/**" # Git 메타데이터 + - "public/**" # 빌드 결과물 + - "static/**" # 정적 파일 (이미지, favicon 등) + - "layouts/**" # HTML 템플릿 파일 + - "archetypes/**" # Hugo 템플릿 파일 + - "assets/**" # CSS 등 에셋 파일 (선택적) + +# 코드 리뷰 설정 +code_review: + disable: false + comment_severity_threshold: MEDIUM # MEDIUM 이상의 이슈만 코멘트 + max_review_comments: -1 # 최대 코멘트 수 제한 + pull_request_opened: + help: false + summary: true + code_review: true + include_drafts: true # 초안 PR에서 제미나이로 수정 많이 해주세요 + +# 마크다운 파일에 집중 +file_patterns: + - "content/posts/**/*.md" + - "content/**/*.md" diff --git a/.gemini/styleguide.md b/.gemini/styleguide.md new file mode 100644 index 0000000..fae4978 --- /dev/null +++ b/.gemini/styleguide.md @@ -0,0 +1,621 @@ +# HyperAccel Blog Style Guide + +This document defines style rules for writing markdown posts in the HyperAccel technical blog. Follow these rules to ensure consistency, readability, and proper rendering across all blog posts. + +--- + +## Table of Contents + +1. [Critical Issues (Must Fix)](#1-critical-issues-must-fix) + - [Markdown Bold Rendering](#11-markdown-bold-rendering) + - [Acronym Usage](#12-acronym-usage) + - [Term Capitalization Consistency](#13-term-capitalization-consistency) +2. [Content Quality (High Priority)](#2-content-quality-high-priority) + - [Readability](#21-readability) + - [Technical Writing](#22-technical-writing) +3. [Markdown Formatting (Medium Priority)](#3-markdown-formatting-medium-priority) + - [Headings](#31-headings) + - [Code Blocks](#32-code-blocks) + - [Images and Links](#33-images-and-links) + - [Lists and Quotes](#34-lists-and-quotes) +4. [Frontmatter and Metadata](#4-frontmatter-and-metadata) +5. [Complete Examples](#5-complete-examples) +6. [Review Checklist](#6-review-checklist) + +--- + +## 1. Critical Issues (Must Fix) + +These issues cause rendering problems or significantly impact readability. They must be fixed before merging. + +### 1.1 Markdown Bold Rendering + +**Problem**: When Korean text immediately follows bold text, markdown rendering may break, causing asterisks to appear as literal characters in the rendered output. + +**Rules**: +- Always add a space when Korean particles (은, 는, 이, 가, 을, 를, 으로, 로, 의, 에, 에서, 와, 과, 등) immediately follow bold markers `**` +- Pay special attention when Korean text follows bold text ending with parentheses +- This applies to all bold text, not just acronyms +- The space should be between the closing `**` and the Korean text + +**Incorrect Examples**: +```markdown +**SM(Streaming Multiprocessor)**으로 구성되어 있습니다. +**TPU(Tensor Processing Unit)**는 이와 대조적입니다. +**대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)**하는 방식입니다. +**GPU**는 병렬 연산에 최적화되어 있습니다. +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)** 으로 구성되어 있습니다. +**Tensor Processing Unit(TPU)** 는 이와 대조적입니다. +**대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)** 하는 방식입니다. +**GPU** 는 병렬 연산에 최적화되어 있습니다. +``` + +**How to Check**: +- Search for pattern: `\*\*[^*]+\*\*[가-힣]` (bold text immediately followed by Korean) +- Visual inspection: Render the markdown and check if asterisks appear as literal characters + +**Checkpoints**: +- [ ] No pattern like `**text(acronym)**한글` exists +- [ ] No pattern like `**text**한글` exists +- [ ] All bold markers followed by Korean text have a space +- [ ] Rendered output shows no literal asterisks + +--- + +### 1.2 Acronym Usage + +**Problem**: Using acronyms before introducing the full term confuses readers. Incorrect formatting makes terms hard to understand and breaks the reading flow. + +**Rules**: +- **First use**: Always write `FullTerm(Acronym)` format +- **Subsequent use**: Only the acronym can be used +- **Bold formatting**: Bold only the full term, not the acronym: `**FullTerm(Acronym)**` +- If an acronym appears in the title or summary, introduce it in the first paragraph +- For well-known acronyms (e.g., CPU, GPU, AI), you may skip the full form if it's extremely common + +**Incorrect Examples**: +```markdown +SM은 GPU의 핵심 구성 요소입니다. (Acronym used before full term) +**SM(Streaming Multiprocessor)**으로 구성되어 있습니다. (Acronym is bolded) +이 글에서는 TPU에 대해 설명합니다. (TPU not introduced) +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 구성 요소입니다. +이후 문장에서는 SM으로 지칭할 수 있습니다. + +이 글에서는 **Tensor Processing Unit(TPU)**에 대해 설명합니다. +TPU는 구글이 개발한 AI 가속기입니다. +``` + +**Common Acronyms in This Blog** (must follow the rule): +- GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX, Pallas, Triton, VLIW, DMA, ALU, SIMD, SPMD + +**Well-Known Acronyms** (may skip full form): +- CPU, AI, ML, API, URL, HTTP, HTTPS + +**How to Check**: +- Search for acronyms in the document +- Verify that each acronym appears with its full term before standalone use +- Check that acronyms in title/summary are introduced in the first paragraph + +**Checkpoints**: +- [ ] Acronym is presented with full term when it first appears +- [ ] Acronym does not appear before the full term +- [ ] Only the full term is bolded, not the acronym +- [ ] Acronyms in title/summary are introduced in first paragraph + +--- + +### 1.3 Term Capitalization Consistency + +**Problem**: Inconsistent capitalization of the same term throughout a document confuses readers and looks unprofessional. + +**Rules**: +- Decide capitalization when a technical term first appears, then maintain it consistently throughout the document +- Acronyms should be all uppercase (e.g., GPU, TPU, SM, CUDA, HBM, SMEM) +- Common nouns should be lowercase (e.g., loop, thread, warp, core, unit, array) +- Proper nouns and brand names follow their official capitalization (e.g., NVIDIA, Google, TensorFlow, PyTorch) +- When in doubt, check official documentation or websites + +**Incorrect Examples**: +```markdown +이 문서에서는 Loop와 loop를 혼용하고 있습니다. +GPU와 gpu가 섞여 있습니다. +Tensorflow와 TensorFlow가 혼용됩니다. +``` + +**Correct Examples**: +```markdown +이 문서에서는 loop를 사용합니다. +GPU는 Graphics Processing Unit의 약자입니다. +TensorFlow는 구글이 개발한 프레임워크입니다. +``` + +**Common Terms and Their Correct Capitalization**: +- **Lowercase**: loop, thread, warp, core, unit, array, buffer, memory, register +- **Uppercase (Acronyms)**: GPU, TPU, SM, CUDA, HBM, SMEM, VMEM, MXU, VPU, TCS, XLA, JAX +- **Brand Names**: NVIDIA, Google, TensorFlow, PyTorch, JAX, Pallas, Triton + +**How to Check**: +- Search for variations of the same term (e.g., "Loop", "loop", "LOOP") +- Verify consistency throughout the document +- Check brand names against official sources + +**Checkpoints**: +- [ ] Same term uses consistent capitalization within the document +- [ ] All acronyms are unified in uppercase +- [ ] Brand names use correct capitalization +- [ ] Common nouns are lowercase + +--- + +## 2. Content Quality (High Priority) + +These rules ensure blog posts are readable, well-structured, and technically accurate. + +### 2.1 Readability + +**Problem**: Long sentences and paragraphs make technical content hard to follow. Readers lose focus and miss important information. + +**Rules**: + +**Sentence Length**: +- A sentence should be **within 50 characters** (recommended, flexible for technical terms) +- Break complex explanations into multiple sentences +- Split long sentences connected with commas (`,`) or conjunctions (`그리고`, `또한`, `또한`) +- Use periods (`.`) more liberally to create shorter, clearer sentences +- One idea per sentence when possible + +**Paragraph Length**: +- A paragraph should be approximately **5-7 sentences** +- Consider splitting if a paragraph has **10 or more sentences** +- Consider splitting if a paragraph exceeds **200 characters** +- Use blank lines to separate distinct ideas +- Each paragraph should focus on one main idea + +**Incorrect Example**: +```markdown +TPU는 AI 연산에 특화된 하드웨어로, 구글이 2016년부터 개발해온 것으로, 행렬 곱셈 연산에 최적화되어 있으며, Systolic Array라는 특별한 연산 유닛을 사용하며, 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**Correct Example**: +```markdown +TPU는 AI 연산에 특화된 하드웨어입니다. 구글이 2016년부터 개발해온 TPU는 행렬 곱셈 연산에 최적화되어 있습니다. + +TPU는 Systolic Array라는 특별한 연산 유닛을 사용합니다. 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**How to Check**: +- Count characters in each sentence (approximate) +- Count sentences in each paragraph +- Look for sentences with multiple commas or conjunctions +- Check for paragraphs that span multiple screen heights + +**Checkpoints**: +- [ ] No sentence exceeds 50 characters (flexible for technical terms) +- [ ] No paragraph has 10 or more sentences +- [ ] No paragraph exceeds 200 characters +- [ ] Complex explanations are broken into multiple sentences +- [ ] Blank lines separate distinct ideas + +--- + +### 2.2 Technical Writing + +**Problem**: Technical content without proper explanations, context, or structure is hard to understand. + +**Rules**: + +**Term Definitions**: +- Briefly explain technical terms when they first appear +- Provide context: why is this term important? How does it relate to the topic? +- Follow the "Acronym Usage Rules" above for acronyms +- Use examples or analogies when helpful + +**Code Examples**: +- Use code examples that actually work (test them if possible) +- Add explanations through comments in the code +- Explain complex code step by step before or after the code block +- Use appropriate language tags for syntax highlighting +- If code is pseudocode or conceptual, mention it in the text + +**Diagrams and Images**: +- Use diagrams to explain complex concepts (architecture, flow, relationships) +- Add meaningful alt text that describes the content, not just "image" or "diagram" +- Use captions when helpful for context +- Name image files meaningfully (e.g., `tpu-architecture.png`, not `image1.png`) + +**Structure**: +- Start with an introduction that sets context +- Use headings to organize content logically +- End with a summary or conclusion +- Use transitions between sections +- Guide readers from general to specific concepts + +**Incorrect Examples**: +```markdown +SM은 GPU의 구성 요소입니다. (No explanation of what SM does or why it matters) + +```python +def compute(x, y): + return x * y +``` (No explanation of what this code does or why it's relevant) +``` + +**Correct Examples**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 연산 단위입니다. 각 SM은 수백 개의 CUDA 코어를 포함하며, 독립적으로 스레드 블록을 실행합니다. 이는 GPU의 병렬 처리 능력의 기반이 됩니다. + +다음 코드는 행렬 곱셈을 수행합니다: + +```python +def matrix_multiply(A, B): + """두 행렬 A와 B를 곱합니다.""" + return A @ B # NumPy의 행렬 곱셈 연산자 사용 +``` + +이 함수는 NumPy의 최적화된 행렬 곱셈을 사용하여 효율적으로 연산을 수행합니다. +``` + +**Checkpoints**: +- [ ] Technical terms are explained when first introduced +- [ ] Code examples have comments or explanations +- [ ] Images have meaningful alt text +- [ ] Content has logical structure with clear sections +- [ ] Transitions connect sections smoothly + +--- + +## 3. Markdown Formatting (Medium Priority) + +These rules ensure proper markdown rendering and consistent formatting. + +### 3.1 Headings + +**Rules**: +- Use H1(`#`) only for the page title (usually in frontmatter, not in body) +- Start body sections with H2(`##`) +- Use H3(`###`) for subsections +- **Maximum heading level is H3** - Do not use H4(`####`) or deeper levels +- Do not skip heading hierarchy (e.g., do not use H3 after H1) +- Use descriptive headings that summarize the section content +- Headings should form a logical outline of the content +- If you need more levels of organization, use lists or paragraphs instead of deeper headings + +**Incorrect Examples**: +```markdown +# Main Title +## Section 1 +#### Subsection (Skipped H3, and H4 not allowed) + +## TPU Architecture +### Systolic Array +#### Data Flow (H4 not allowed - maximum is H3) +``` + +**Correct Examples**: +```markdown +## TPU Architecture +### Systolic Array + +데이터 흐름은 다음과 같습니다: (Use paragraph instead of H4) + +## TPU Architecture +### Systolic Array +### Data Flow (Use H3 instead of H4) +``` + +**Checkpoints**: +- [ ] H1 is only used for page title +- [ ] Maximum heading level is H3 (no H4 or deeper) +- [ ] Heading hierarchy is not skipped +- [ ] Headings are descriptive +- [ ] Headings form a logical outline + +--- + +### 3.2 Code Blocks + +**Rules**: +- Code blocks must specify language tags: ` ```python`, ` ```bash`, ` ```cuda`, etc. +- Use single backticks for inline code: `` `variable_name` `` +- For file paths or commands, use inline code: `` `hugo server` `` +- Add comments in code to explain complex logic +- If code is pseudocode or conceptual, mention it in the text +- Format code properly (indentation, spacing) + +**Incorrect Examples**: +```markdown +``` +def hello(): + print("world") +``` +(No language tag) +``` + +**Correct Examples**: +```python +def hello(): + """Print hello world.""" + print("world") +``` + +**Common Languages in This Blog**: +- `python`, `cuda`, `bash`, `c`, `go`, `fsharp`, `javascript`, `yaml`, `toml` + +**Checkpoints**: +- [ ] All code blocks have language tags +- [ ] Inline code uses single backticks +- [ ] Complex code has comments or explanations +- [ ] Code is properly formatted + +--- + +### 3.3 Images and Links + +**Rules**: + +**Images**: +- Always include alt text: `![descriptive alt text](path/to/image.png)` +- Alt text should describe the content, not just say "image" or "diagram" +- Use relative paths for images in the same post directory +- Use meaningful file names (e.g., `tpu-architecture.png`, not `img1.png`) +- Prefer `.webp` format for better compression +- **Image placement**: Place images **BEFORE** the content they illustrate (recommended) + - This allows readers to see the visual context before reading the explanation + - Example: Place a diagram before the paragraph that explains it, not after + +**Links**: +- Use descriptive link text: `[descriptive text](URL)` +- Avoid generic text like "click here" or "link" +- For external links, consider adding `{:target="_blank"}` if needed (Hugo-specific) +- Link to authoritative sources when possible + +**Incorrect Examples**: +```markdown +![image](tpu.png) (No descriptive alt text) +[여기](https://example.com)를 클릭하세요. (Generic link text) + +TPU는 Systolic Array 구조를 사용합니다. +![TPU architecture](tpu-architecture.png) (Image placed after content) +``` + +**Correct Examples**: +```markdown +![TPU architecture diagram showing Systolic Array](tpu-architecture.png) + +TPU는 Systolic Array 구조를 사용합니다. 이 구조는... + +[TPU 논문](https://arxiv.org/pdf/1704.04760)을 참고하세요. +``` + +**Checkpoints**: +- [ ] All images have descriptive alt text +- [ ] Images are placed BEFORE the content they illustrate (recommended) +- [ ] Link text is descriptive +- [ ] Image file names are meaningful +- [ ] Images use appropriate formats + +--- + +### 3.4 Lists and Quotes + +**Rules**: + +**Lists**: +- Use `-` for unordered lists (consistent throughout) +- Use `1.` for ordered lists when order matters +- Maintain consistency between list items (same structure, same level of detail) +- Nested lists are recommended up to 2 levels +- Add blank lines before and after lists for readability +- Use parallel structure in list items + +**Block Quotes**: +- Use `>` for block quotes +- Use quotes for important statements, definitions, or callouts +- Keep quotes concise +- Use quotes to highlight key concepts or definitions + +**Incorrect Examples**: +```markdown +- Item 1 +* Item 2 (Inconsistent list markers) +``` + +**Correct Examples**: +```markdown +- Item 1 +- Item 2 + - Sub-item 2.1 + - Sub-item 2.2 +``` + +**Checkpoints**: +- [ ] List markers are consistent +- [ ] List items have consistent structure +- [ ] Quotes are used appropriately +- [ ] Lists have proper spacing + +--- + +## 4. Frontmatter and Metadata + +**Rules**: +- All posts must have frontmatter with required fields +- `date` should be in ISO 8601 format: `'2026-01-03T17:20:16+09:00'` +- `draft: false` for published posts +- `title` should be descriptive and match the H1 heading +- `tags` should be relevant and consistent (use existing tags when possible) +- `authors` must match entries in `content/authors/` +- `summary` should be a concise one-line description +- `description` can be longer (used for SEO) +- For series posts, include `series: ["시리즈 이름"]` +- `cover` image is optional but recommended + +**Required Fields**: +```yaml +--- +date: '2026-01-03T17:20:16+09:00' +draft: false +title: 'Post Title' +authors: ["Author Name"] +tags: ["Tag1", "Tag2"] +categories: ["Category"] +summary: "One-line summary" +description: "Longer description for SEO" +--- +``` + +**Optional Fields**: +```yaml +series: ["Series Name"] +cover: + image: "path/to/image.png" + alt: "Image description" + caption: "Image caption" + relative: true +comments: true +``` + +**Checkpoints**: +- [ ] All required frontmatter fields are present +- [ ] Date format is correct +- [ ] Authors match existing author entries +- [ ] Tags are relevant and consistent +- [ ] Summary is concise and descriptive + +--- + +## 5. Complete Examples + +### 5.1 Before and After: Acronym and Bold Rendering + +**Before (Incorrect)**: +```markdown +SM은 GPU의 핵심 구성 요소입니다. **SM(Streaming Multiprocessor)**으로 구성되어 있으며, 각 SM은 수백 개의 CUDA 코어를 포함합니다. +``` + +**After (Correct)**: +```markdown +**Streaming Multiprocessor(SM)**은 GPU의 핵심 구성 요소입니다. GPU는 수백 개의 SM 으로 구성되어 있으며, 각 SM은 수백 개의 CUDA 코어를 포함합니다. +``` + +### 5.2 Before and After: Readability + +**Before (Incorrect)**: +```markdown +TPU는 AI 연산에 특화된 하드웨어로, 구글이 2016년부터 개발해온 것으로, 행렬 곱셈 연산에 최적화되어 있으며, Systolic Array라는 특별한 연산 유닛을 사용하며, 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +**After (Correct)**: +```markdown +TPU는 AI 연산에 특화된 하드웨어입니다. 구글이 2016년부터 개발해온 TPU는 행렬 곱셈 연산에 최적화되어 있습니다. + +TPU는 Systolic Array라는 특별한 연산 유닛을 사용합니다. 이는 일반적인 CPU나 GPU와는 다른 접근 방식을 취하고 있습니다. +``` + +### 5.3 Complete Post Example + +**Good Example Structure**: +```markdown +--- +date: '2026-01-03T17:20:16+09:00' +draft: false +title: '지피지기면 백전불태 2편: TPU의 등장과 부상' +authors: ["Jaewon Lim"] +tags: ["TPU", "Google", "Ironwood"] +categories: ["AI Hardware"] +series: ["지피지기면 백전불태"] +summary: "TPU의 등장 배경과 하드웨어/소프트웨어 구조를 분석합니다." +--- + +# 지피지기면 백전불태 2편: TPU의 등장과 부상 + +> **지피지기면 백전불태(知彼知己 百戰不殆)** +> 상대를 알고 나를 알면 백 번 싸워도 위태롭지 않다는 뜻입니다. + +## TPU? + +최근 AI 가속기 시장에서 가장 핫하게 떠오른 키워드 중 하나는 바로 **Tensor Processing Unit(TPU)**일 것입니다. + +## TPU의 등장 + +**Tensor Processing Unit(TPU)** 구조를 이해하기 위해서는 먼저 TPU가 개발된 배경에 대해 알아야 할 필요가 있습니다. + +![TPU 실제 사진](tpu2.webp) + +## TPU 하드웨어 아키텍처 + +### Systolic Array + +TPU는 AI 연산에 특화되어 제작되었습니다. AI 연산 중 가장 큰 특징은 대규모의 행렬 곱셈을 수행한다는 점입니다. + +다음 코드는 행렬 곱셈의 개념을 보여줍니다: + +```python +def matrix_multiply(A, B): + """두 행렬을 곱합니다.""" + return A @ B # NumPy의 행렬 곱셈 +``` + +이러한 Systolic Array는 TPU에서 사용한 대표적인 특징입니다. +``` + +--- + +## 6. Review Checklist + +When reviewing a blog post, check the following items in order of priority: + +### Critical (Must Fix) +- [ ] **Bold rendering**: No `**text**한글` patterns without spaces +- [ ] **Acronyms**: All acronyms introduced with full terms first +- [ ] **Capitalization**: Consistent capitalization throughout + +### High Priority +- [ ] **Readability**: Sentences within 50 characters, paragraphs within 10 sentences +- [ ] **Technical terms**: All technical terms explained when first introduced +- [ ] **Code examples**: All code blocks have language tags and explanations +- [ ] **Images**: All images have descriptive alt text and are placed before the content they illustrate + +### Medium Priority +- [ ] **Headings**: Proper hierarchy (H2 → H3, maximum H3, no skipping) +- [ ] **Links**: Descriptive link text +- [ ] **Lists**: Consistent formatting +- [ ] **Frontmatter**: All required fields present and correct + +### Additional Checks +- [ ] Content flows logically with clear transitions +- [ ] No typos or grammatical errors +- [ ] Examples are accurate and tested (if applicable) +- [ ] Images are properly placed (before content) and referenced + +--- + +## Notes for Reviewers + +- **Severity Levels**: + - **Critical**: Causes rendering issues or major readability problems + - **High**: Significantly impacts readability or understanding + - **Medium**: Formatting or consistency issues + +- **Be Constructive**: When pointing out issues, suggest specific fixes +- **Context Matters**: Some rules may be flexible for technical content (e.g., longer sentences for complex explanations) +- **Consistency First**: If a pattern is used consistently throughout a post, it may be acceptable even if not ideal + +--- + +## Common Patterns in This Blog + +Based on existing posts, here are common patterns to follow: + +- **Series Posts**: Use series tag for series like `series: ["지피지기면 백전불태"]` +- **Code Languages**: Python, CUDA, Bash, C, Go, F# are commonly used +- **Image Formats**: `.webp`, `.png`, `.jpg` are used +- **Technical Terms**: GPU, TPU, SM, CUDA, HBM, Systolic Array, Tensor Core, etc. diff --git a/assets/css/extended/custom.css b/assets/css/extended/custom.css index 93643e4..e29f795 100644 --- a/assets/css/extended/custom.css +++ b/assets/css/extended/custom.css @@ -219,3 +219,34 @@ body .posts-grid-container .post-entry { color: var(--primary); text-decoration: underline; } + +/* Table styling with borders */ +.post-content table { + width: 100% !important; + border-collapse: collapse !important; + margin: 1.5rem 0 !important; +} + +.post-content table thead th, +.post-content table tbody th { + border: 1px solid var(--border) !important; + padding: 16px 24px !important; + text-align: center !important; + vertical-align: middle !important; +} + +.post-content table tbody td { + border: 1px solid var(--border) !important; + padding: 16px 24px !important; + text-align: left !important; + vertical-align: middle !important; +} + +.post-content table thead th { + background-color: var(--code-bg) !important; + font-weight: 600 !important; +} + +.post-content table tbody tr:hover { + background-color: var(--code-bg) !important; +} diff --git a/content/posts/TPU-deep-dive/index.ko.md b/content/posts/TPU-deep-dive/index.ko.md index 5d39d02..5be3d23 100644 --- a/content/posts/TPU-deep-dive/index.ko.md +++ b/content/posts/TPU-deep-dive/index.ko.md @@ -101,9 +101,9 @@ Systolic array는 이 과정을 생략할 수 있는 대안 중 하나인데요. ### TPU Building Block 다음으로는 TPU 개별 칩에서 사용하는 연산 단위에 대해 알아보겠습니다. -앞서 살펴본 GPU는 수백 개의 **SM(Streaming Multiprocessor)**으로 구성되어 있습니다. 하드웨어 스케줄러가 32개 스레드 묶음인 '워프(Warp)' 중 실행 준비가 된 것을 실시간으로 선택해 연산 자원에 투입하는데, 이는 명령어를 무작위로 실행하는 것이 아니라 **대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)**하는 방식입니다. 즉, GPU는 수많은 작은 연산 단위를 촘촘하게 관리하며 전체 처리량을 높이는 데 최적화되어 있습니다. +앞서 살펴본 GPU는 수백 개의 **Streaming Multiprocessor(SM)** 으로 구성되어 있습니다. 하드웨어 스케줄러가 32개 스레드 묶음인 '워프(Warp)' 중 실행 준비가 된 것을 실시간으로 선택해 연산 자원에 투입하는데, 이는 명령어를 무작위로 실행하는 것이 아니라 **대기 시간을 다른 작업으로 메워 효율을 극대화(Latency Hiding)** 하는 방식입니다. 즉, GPU는 수많은 작은 연산 단위를 촘촘하게 관리하며 전체 처리량을 높이는 데 최적화되어 있습니다. -반면, **TPU(Tensor Processing Unit)**는 이와 대조적인 접근 방식을 취합니다. 잘게 쪼개진 수많은 유닛 대신, 거대한 행렬 연산을 한 번에 처리할 수 있는 소수의 강력한 전용 코어를 탑재하고 있습니다. 아울러 여러 인스트럭션을 합쳐둔 더 고수준의 명령어(VLIW, Very Long Instruction Word)를 통해 다량의 데이터를 연산 블록에 한번에 밀어넣는 방식으로 동작합니다. 복잡한 동적 스케줄링의 비중을 줄이는 대신, 인공지능 연산의 핵심인 대규모 행렬 곱셈을 마치 거대한 파이프라인이 흐르듯 단번에 수행하도록 설계된 것이 TPU 아키텍처의 핵심입니다. +반면, **Tensor Processing Unit(TPU)** 는 이와 대조적인 접근 방식을 취합니다. 잘게 쪼개진 수많은 유닛 대신, 거대한 행렬 연산을 한 번에 처리할 수 있는 소수의 강력한 전용 코어를 탑재하고 있습니다. 아울러 여러 인스트럭션을 합쳐둔 더 고수준의 명령어(VLIW, Very Long Instruction Word)를 통해 다량의 데이터를 연산 블록에 한번에 밀어넣는 방식으로 동작합니다. 복잡한 동적 스케줄링의 비중을 줄이는 대신, 인공지능 연산의 핵심인 대규모 행렬 곱셈을 마치 거대한 파이프라인이 흐르듯 단번에 수행하도록 설계된 것이 TPU 아키텍처의 핵심입니다. TPU의 개별 칩에 사용되는 코어는 선형 대수 연산에 특화된 TensorCore와 임베딩 연산에 특화된 SparseCore로 나뉩니다. (GPU에서 사용하는 Tensor Core와 용어는 같지만 역할의 차이가 있습니다.) diff --git a/content/posts/pallas-programming-model/images/00-main-image.png b/content/posts/pallas-programming-model/images/00-main-image.png new file mode 100644 index 0000000..6d04f26 Binary files /dev/null and b/content/posts/pallas-programming-model/images/00-main-image.png differ diff --git a/content/posts/pallas-programming-model/images/01-tpuv1-arch.png b/content/posts/pallas-programming-model/images/01-tpuv1-arch.png new file mode 100644 index 0000000..daeeb78 Binary files /dev/null and b/content/posts/pallas-programming-model/images/01-tpuv1-arch.png differ diff --git a/content/posts/pallas-programming-model/images/02-tpuv7-arch.png b/content/posts/pallas-programming-model/images/02-tpuv7-arch.png new file mode 100644 index 0000000..f95c27c Binary files /dev/null and b/content/posts/pallas-programming-model/images/02-tpuv7-arch.png differ diff --git a/content/posts/pallas-programming-model/images/04-grid-and-block.png b/content/posts/pallas-programming-model/images/04-grid-and-block.png new file mode 100644 index 0000000..0b7a4e4 Binary files /dev/null and b/content/posts/pallas-programming-model/images/04-grid-and-block.png differ diff --git a/content/posts/pallas-programming-model/images/05-grid-block-in-tpu.png b/content/posts/pallas-programming-model/images/05-grid-block-in-tpu.png new file mode 100644 index 0000000..16ad675 Binary files /dev/null and b/content/posts/pallas-programming-model/images/05-grid-block-in-tpu.png differ diff --git a/content/posts/pallas-programming-model/images/06-common-hw-model.png b/content/posts/pallas-programming-model/images/06-common-hw-model.png new file mode 100644 index 0000000..32d74e9 Binary files /dev/null and b/content/posts/pallas-programming-model/images/06-common-hw-model.png differ diff --git a/content/posts/pallas-programming-model/images/07-tpu-hw-model.png b/content/posts/pallas-programming-model/images/07-tpu-hw-model.png new file mode 100644 index 0000000..3c601d5 Binary files /dev/null and b/content/posts/pallas-programming-model/images/07-tpu-hw-model.png differ diff --git a/content/posts/pallas-programming-model/images/08-gpu-hw-model.png b/content/posts/pallas-programming-model/images/08-gpu-hw-model.png new file mode 100644 index 0000000..1d7d6c3 Binary files /dev/null and b/content/posts/pallas-programming-model/images/08-gpu-hw-model.png differ diff --git a/content/posts/pallas-programming-model/images/09-pallas-lowering.png b/content/posts/pallas-programming-model/images/09-pallas-lowering.png new file mode 100644 index 0000000..e8812bb Binary files /dev/null and b/content/posts/pallas-programming-model/images/09-pallas-lowering.png differ diff --git a/content/posts/pallas-programming-model/index.en.md b/content/posts/pallas-programming-model/index.en.md new file mode 100644 index 0000000..69e9e29 --- /dev/null +++ b/content/posts/pallas-programming-model/index.en.md @@ -0,0 +1,377 @@ +--- +date: '2026-01-19T14:50:09+09:00' +draft: false +title: 'Know Your Enemy 2.5: Pallas Programming Model' +cover: + image: "" + # can also paste direct link from external site + # ex. https://i.ibb.co/K0HVPBd/paper-mod-profilemode.png + alt: "Pallas Programming Model" + caption: "" + relative: true # To use relative path for cover image, used in hugo Page-bundles +authors: ["Donghyeon Choi"] # must match with content/authors +tags: ["Pallas", "TPU", "Google", "Ironwood", "Kernel", "Custom Kernel", "Programming Model"] +categories: ["AI Hardware", "Computer Architecture"] +series: ["Know Your Enemy, Know Yourself"] +summary: "Learn about Pallas programming model that enables writing custom kernels on TPU." +comments: true +description: "Building on the TPU architecture explored in part 2, we examine Pallas, a programming model that enables writing custom kernels on TPU." +keywords: [ + "Pallas", "TPU", "Google", "Ironwood", "Custom Kernel", "Kernel Programming", + "AI Hardware", "Programming Model", "TPU Optimization" +] +--- + +# Know Your Enemy 2.5: Pallas Programming Model + +[Know Your Enemy 1: GPU History and Fundamentals](https://hyper-accel.github.io/en/posts/how-gpu-works/) +[Know Your Enemy 2: The Rise of TPU](https://hyper-accel.github.io/en/posts/tpu-deep-dive/) + +> **Know Your Enemy, Know Yourself (知彼知己 百戰不殆)** +> If you know your enemy and know yourself, you need not fear the result of a hundred battles. +> This series aims to deeply understand competitors' hardware for AI accelerator design. +> Part 2.5 covers **Pallas**, a programming model that enables writing custom kernels on TPU. + +In part 2, we explored TPU's hardware architecture and software stack. We discussed how the XLA compiler optimizes operations and its limitations. At the end, we previewed that the next article (part 3) would introduce Groq's LPU. + +However, to truly understand TPU, hardware and software stack alone are not enough. Especially in the latest TPU generation, Ironwood, the **Pallas** programming model plays a crucial role in performance optimization. Pallas is a tool that allows writing custom kernels on TPU, similar to CUDA or Triton. It enables direct control over hardware details while remaining relatively easy to use in a Python environment. + +As mentioned in part 2 regarding the limitations of the XLA compiler, automatic compilers often struggle to optimize cutting-edge algorithms. Pallas is the kernel language Google created to overcome these limitations. When announced alongside Ironwood, it was emphasized with the slogan "Extreme performance: Custom kernels via Pallas," highlighting its important role in maximizing TPU performance. + +Therefore, before diving into Groq's LPU in part 3, we first examine the Pallas programming model in part 2.5 to complete our understanding of TPU. In this article, we'll explore what Pallas is, why it's needed, and how it enhances TPU performance. + +--- + +## Background: TPU Architecture and Why Pallas is Needed + +To understand Pallas, we first need to understand how TPU differs from CPUs or GPUs. This difference is exactly why Pallas is needed. + +### Traditional Architecture vs TPU's Systolic Array + +Traditional CPU or GPU architectures repeatedly fetch and store data from SRAM to register files for each operation. This approach causes massive **memory bandwidth waste** in operations like matrix multiplication that repeatedly use the same data. + +TPU solves this by adopting a structure where **once-loaded data flows between compute units and is continuously calculated**. This is the Systolic Array structure. + +### TPU's Data Flow + +Operations within TPU follow this flow: + +1. **Weight Supply**: Weight data stored in DRAM is loaded into MXU (Matrix Multiply Unit) through Weight FIFO. +2. **Activation Supply**: Input data stored in Unified Buffer (UB) is delivered to MXU. +3. **Matrix Operations**: Multiplication and accumulation (Multiply-Accumulate) operations occur simultaneously through the Systolic Array structure inside MXU. +4. **Post-processing Pipeline**: MXU output values pass through Accumulator, Activation Unit (ReLU, etc.), and Normalize/Pool Unit sequentially to perform additional operations required by AI models. +5. **Result Storage**: Final output after all operation flows returns to Unified Buffer and is stored. + +### Key Difference from GPU: Change in Execution Unit + +This structural difference creates a very significant difference from a programming model perspective. + +- **GPU (SIMT)**: Numerous threads perform calculations independently, managed in groups of 32 threads called Warps. It features fine-grained parallel processing focused on individual data. +- **TPU (Tensor-centric, SPMD)**: A single data load completes the entire operation sequence. That is, rather than individual threads, **'the entire program for one tensor (or tile)' is considered as one minimum execution unit**. + +> **Why Pallas?** Pallas is a language designed to abstract these TPU hardware characteristics (direct Unified Buffer control, MXU scheduling, etc.) while allowing developers to directly optimize this 'tensor-level flow'. + +TPU has evolved through generations, strengthening vector operation units and developing scale-up technologies for large-scale models. The latest architecture, Ironwood, has become a 'monster compute unit' with a massive package adopting chiplet structure and mounting **four 256x256 Systolic Arrays (MXU)**. However, despite the dramatic increase in hardware scale, the core design principle of **'performing as many operations as possible through a single data load'** remains unchanged. + +### Why Pallas is Needed + +When we discussed the XLA compiler in part 2, we mentioned that while XLA is a powerful optimization compiler, it has limitations. When new operation algorithms emerge, it's difficult for manually created custom kernels to match performance until the compiler is updated to a version that can optimize them. + +For example, cutting-edge algorithms like Flash Attention or MoE (Mixture of Experts) often have complex memory access patterns or high data dependencies, making them difficult for automatic compilers to optimize. In such cases, developers need to directly understand the memory hierarchy and finely control how to tile data, when to move data between memory levels, and so on. + +On GPUs, this problem was solved with CUDA or Triton. CUDA allows direct hardware control but has a high barrier to entry, while Triton has a higher level of abstraction, making it relatively easier to use. However, both only work on GPUs. + +What about TPU? Google began providing **Pallas** as an experimental extension of JAX around 2023. Pallas shares a similar philosophy with Triton but differs significantly in that it supports both GPUs and TPUs. + +--- + +## What is Pallas? + +Pallas is a kernel language that enables writing custom kernels within the JAX ecosystem. It works on both GPUs and TPUs, allowing direct control over hardware memory hierarchy, data tiling, and block-level parallelism. + +### Basic Concepts + +Pallas's core idea is simple. It allows developers to control operations that are difficult for high-level automated compilers to handle, at a level closer to hardware. However, unlike CUDA, it doesn't go completely low-level but is abstracted to be relatively easy to use in a Python environment. + +### Key Components + +Pallas consists of several core concepts: + +**1. Grid and Program ID: Parallel Execution Abstraction** + +Pallas models execution units through a **Grid** abstraction. Grid has different meanings depending on hardware but provides a unified interface. + +```python +def kernel(o_ref): + i = pl.program_id(0) # Current program's ID + o_ref[i] = i + +# Grid size specification: (8,) = 8 parallel programs +result = pl.pallas_call( + kernel, + out_shape=jax.ShapeDtypeStruct((8,), jnp.int32), + grid=(8,), # 8 parallel instances +)() +``` + +#### Grid Semantics: GPU vs TPU + +In Pallas, **Grid** is a collection of execution units that divide the entire work, but how hardware processes this Grid has fundamental differences depending on the architecture. + +**Grid on GPU**: In GPU backends (Triton/Mosaic), Grid assumes **fully parallel execution** by hardware schedulers. Each Grid item is mapped to one **Thread Block** and executed independently on individual SMs (Streaming Multiprocessors). Since hardware schedules threads non-deterministically, execution order between Grids is not guaranteed. Therefore, `BlockSpec`'s `index_map` design must strictly manage **race conditions** to prevent different programs from writing to the same HBM location. + +**Grid on TPU**: In TPU backends, Grid is a model that combines parallelism between multiple cores and **sequential pipelining** within a single core. TPU is a very wide SIMD machine, but can be coded like a **single-threaded processor** from a software perspective. The control unit (TCS) provides an intuitive flow that controls entire operations by looping. There are Parallel dimensions that distribute work when multiple TensorCores exist and execute **physically simultaneously**, and Sequential dimensions that guarantee **serial execution** within a single TensorCore. Sequential dimensions are not just for slow execution, but are used as a strategic means to **hide memory latency** by overlapping (Overlap) current operations and next data loading through **Semaphores**. + +**2. BlockSpec: Memory Layout Abstraction** + +Pallas abstracts the process of dividing massive data into chunks that hardware can digest through **BlockSpec**. This goes beyond simply cutting data and defines **data transfer protocols between HBM (Remote) and SRAM (Local)**. + +BlockSpec components: + +- `block_shape`: The size of data that each program instance will place on the Local Memory (SRAM) workbench. Designing this not to exceed hardware SRAM capacity is key to performance optimization. +- `index_map` **function**: Takes `grid` indices (i, j) as input and returns the block start position on HBM. This function is analyzed at compile time and converted to hardware DMA (Direct Memory Access) address calculation logic. +- `memory_space`: Specifies the physical container where fragmented data will reside. If not specified, it's allocated to the default space (`pl.SRAM`) according to backend settings. + +**3. Ref: Memory Reference Abstraction** + +Pallas abstracts complex hardware memory address systems through **Ref** objects. This goes beyond simple pointers to data and provides a **logical view of specific data blocks on SRAM (Local Memory)**. + +Ref's key features: + +- **Local Memory Reference**: `Ref` points to data that has already been loaded or will be loaded on **SRAM (TPU's VMEM/SMEM, GPU's Shared Memory)**, the fastest workbench in hardware, not HBM. +- **Dereferencing**: When using brackets like `x_ref[...]`, actual data loading from **SRAM → Register File** occurs, converting it to a 'Value'. +- **Hardware Abstraction**: Even using the same `Ref` interface, it automatically converts to TPU's Vector/Scalar memory or GPU's Shared Memory access depending on the backend. + +**4. Memory Hierarchy Control** + +Pallas provides memory models that share the same big picture but differ in details for TPU and GPU. Let's first look at the overall structure that's abstracted and provided in common. + +The structure abstracted and provided in both: + +- **Remote Memory (HBM)**: Memory space mainly specified through `pl.ANY`. +- **Multiple "Core" structure**: Independent compute units for parallel processing of grids. + - **Local Memory (SRAM or Cache)**: SRAM memory located inside cores, slightly faster. It's recommended to separate READ and WRITE variables to maintain consistency with HBM data. + - **Register Files in Execution Units**: Small memory that can directly exchange data with compute units inside cores. + +Generally, Local Memory takes 10x the latency of Register File, and Remote Memory takes another 10x the latency of Local Memory. However, in TPU v7 Ironwood adopting chiplet structure, since one chip contains HBM, Remote Memory communication can be 2~5x faster than the previous 10x. + +#### TPU Memory Hierarchy Structure + +Looking at TPU's memory hierarchy model: + +- **Local Memory** consists of Vector Memory and Scalar Memory: + - **Vector Memory (VMEM)**: Memory that stores data for Vector and Matrix related operations. Accessible from Vector/Matrix Units (VPU/MXU/XLU) within the same tensor core. + - **Scalar Memory (SMEM)**: Memory that stores data for scalar operations for logic flow (loops, conditions, etc.). Accessible from Scalar Unit (TCS) within the same tensor core. Data stored on SMEM can also be directly used in vector operations through high-level commands generated by TCS. + +- Each VMEM and SMEM has separate READ/WRITE areas to ensure data sync with HBM. This is because modifying data read as READ without write-back can cause data inconsistency with HBM. + +- When calling kernels, data is loaded from Remote Memory to Local Memory according to the specified BlockSpec. At this time, corresponding data is automatically fetched according to the set Grid number. + +- When calling Ref data inside kernels, data is fetched from Local Memory to Register File for operations. During Local Memory ↔ Register File movement and operation, data loading from Remote Memory → Local Memory for the next Grid automatically overlaps. This can hide memory latency significantly. + +- Writing result data to the output Ref data set in the kernel writes back to Local Memory, and when all programs for that grid finish, the output Ref data is written back to Remote Memory. + +#### GPU Memory Hierarchy Structure + +Looking at GPU's hardware model: + +- **Local Memory is Shared Memory (SMEM)**: Each SM (Streaming Multiprocessor) has independent Shared Memory/L1 Cache space. Unlike TPU, it's a unified workspace without Scalar/Vector distinction, where numerous threads access this space simultaneously for parallel operations. + +- **Data Movement and Pipelining**: Unlike existing methods that rely only on hardware scheduling, Pallas GPU explicitly **overlaps HBM (Remote) → SMEM (Local) data loading and TensorCore operations**. Through `plgpu.emit_pipeline`, etc., while compute units process current data, data for the next grid is prefetched asynchronously to hide memory latency. + +- **Ref Variable Entity and Hardware Mapping**: Ref variables inside kernels are addresses of data chunks that have already been loaded or will be loaded on Local Memory (SMEM), not HBM. Through `memory_space=plgpu.GPUMemorySpace.GMEM` settings, HBM (Remote) space is explicitly specified and can be directly controlled with `copy_gmem_to_smem`. + +**5. Memory Pipelining & Semaphore** + +Data movement between Remote Memory and Local Memory in TPU architecture causes latency of hundreds of cycles. To overcome this, Pallas supports pipelining using hardware-level semaphores. + +**Hardware Role of Semaphore**: Acts as a **traffic light** between data movement (DMA) and compute units. It hardware-checks "Has data loading completed?" or "Has computation finished and space become available?" to ensure two tasks safely overlap execution. + +**Double Buffering**: Allocates two buffers (e.g., Buffer 0, Buffer 1) within Local Memory. **While VPU computes Buffer 0, DMA prefetches next data from HBM to Buffer 1**. As this process repeats, memory latency is completely hidden behind computation time. + +Semaphore is a **synchronization device** that coordinates two different engines—data transfer (DMA) and computation (VPU/TCS)—so they don't collide. Through `pltpu.semaphore`, it prevents problems like 'trying to read before data arrives (Read-before-ready)' or 'overwriting with new data while still computing (Write-over-active)' at the hardware level, realizing **Auto-Overlap**. + +**3. Backend Lowering** + +Kernels written in Pallas are eventually converted to hardware code. On GPUs, they are lowered through Triton or Mosaic GPU backends, while on TPUs, they are lowered through the Mosaic compiler to MLIR form and finally converted to hardware code. + +During this process, operator fusion, tiling automation, and overlapping of data transfer and computation are optimized. The high-level code written by developers is transformed into hardware-optimized code. + +**4. Compatibility with JAX Transforms** + +Pallas kernels are compatible with JAX transforms like `jit`, `vmap`, and `grad`. This is a major advantage, as you can still utilize features like automatic differentiation, mapping, and compilation while writing high-performance kernels. + +--- + +## CUDA vs Pallas: Programming Model Comparison + +In part 1, we explored the CUDA programming model. Now let's compare CUDA and Pallas to clearly understand their characteristics and differences. + +### Abstraction Level and Approach + +**CUDA**: A very low-level programming model close to hardware. Through a clear hierarchical structure of Thread, Warp, Block, and Grid, it allows direct control over hardware details. Developers must directly calculate thread IDs and carefully design memory access patterns. + +**Pallas**: Provides one level higher abstraction than CUDA. Working in a Python environment, you only need to define how to divide data through Grid and BlockSpec, and the hardware automatically handles memory movement and scheduling. Developers can focus on "what data to place where." + +### Difference in Execution Units + +**CUDA's Execution Units**: +- **Thread**: The minimum unit of parallel processing. Each thread has a unique thread ID and executes independently. +- **Warp**: A hardware execution unit bundling 32 consecutive threads. Executes the same instruction simultaneously. +- **Thread Block**: A group bundling up to 1024 threads. Shares Shared Memory and can synchronize with `__syncthreads()`. +- **Grid**: The entire unit of kernel execution. A collection of all thread blocks. + +**Pallas's Execution Units**: +- **Grid**: An abstraction specifying the iteration space for parallel execution. A `(4, 5)` grid means 20 work units. +- **BlockSpec**: Defines how to tile data and which memory space to place it in. The hardware automatically handles data movement. +- **Ref**: A logical view of data blocks on Local Memory (SRAM). Accessing with `x_ref[...]` automatically loads to Register File. + +### Memory Hierarchy Control + +**CUDA**: +- Developers must explicitly specify memory spaces: `__global__`, `__shared__`, `__device__`, etc. +- Must directly manage Shared Memory allocation and usage. +- Requires deep understanding of hardware structure to optimize memory access patterns. + +**Pallas**: +- Specifies memory space through `memory_space` parameter, but hardware automatically optimizes. +- Automatically maps to VMEM/SMEM on TPU and Shared Memory on GPU. +- Can declaratively define data tiling and memory movement through BlockSpec. + +### Hardware Target and Portability + +**CUDA**: +- NVIDIA GPU exclusive. +- Optimization methods may vary depending on GPU architecture (e.g., Hopper, Blackwell). +- However, the CUDA ecosystem is very mature with rich tools and libraries. + +**Pallas**: +- Supports both GPU and TPU. +- Flexibility to run the same kernel definition on multiple hardware platforms. +- However, still in experimental stage, so some features may be limited. + +### Development Convenience and Learning Curve + +**CUDA**: +- Low-level language based on C/C++. +- Must deeply understand hardware structure to achieve optimal performance. +- Steep learning curve but provides complete control. + +**Pallas**: +- Python-based, naturally integrated with the JAX ecosystem. +- Compatible with JAX transforms like `jit`, `vmap`, `grad`. +- Relatively easy to start, but still requires hardware understanding for optimal performance. + +### Practical Usage Example Comparison + +Comparing simple vector addition implemented in CUDA and Pallas clearly shows the difference. + +**CUDA Approach**: +```cuda +__global__ void vector_add(float *A, float *B, float *C, int N) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < N) { + C[idx] = A[idx] + B[idx]; + } +} +``` + +**Pallas Approach**: +```python +def vector_add_kernel(a_ref, b_ref, c_ref): + c_ref[...] = a_ref[...] + b_ref[...] + +result = pl.pallas_call( + vector_add_kernel, + out_shape=jax.ShapeDtypeStruct((N,), dtype), + grid=(N,), +)(a, b) +``` + +In CUDA, you must directly calculate thread IDs and perform boundary checks, while Pallas abstracts at the data block level. + +### When to Use What? + +**Choose CUDA when**: +- Maximum performance is needed on NVIDIA GPUs +- Need fine-grained control over complex custom algorithms +- Want to leverage the rich tools and libraries of the CUDA ecosystem + +**Choose Pallas when**: +- Need to support both GPU and TPU +- Want a workflow integrated with the JAX ecosystem +- Want to quickly prototype in a Python environment +- Want to utilize automatic differentiation or other JAX transforms + +--- + +## How Pallas is Used on TPU + +### Ironwood and Pallas + +In the latest TPU generation, Ironwood, Pallas is emphasized with the slogan "Extreme performance: Custom kernels via Pallas." In the Ironwood stack, developers can directly specify strategies for memory tiling, data movement, and MXU utilization through Pallas definitions, and the Mosaic compiler lowers this to TPU code. + +Through this combination, tiling strategies like input stationary, weight stationary, and output stationary, as well as distributed processing at stride and batch levels, can be efficiently designed. Additionally, it enables designs that minimize scheduling bottlenecks in the entire pipeline by overlapping HBM ↔ on-chip memory data transfer and MXU computation simultaneously. + +### Real-World Use Cases + +Pallas is primarily used in the following cases: + +- **Complex Attention Mechanisms**: Cutting-edge attention algorithms like Flash Attention have complex memory access patterns that are difficult for automatic compilers to optimize. Using Pallas, you can efficiently implement them by directly controlling the memory hierarchy. + +- **MoE (Mixture of Experts)**: In MoE models, expert routing logic has high data dependencies, making it difficult to optimize at compile time. Using Pallas, dynamic routing can be handled efficiently. + +- **Sparse Operations**: For sparse matrix operations or ragged tensors, padding may change dynamically or irregular memory access may be required. Pallas is also useful in these cases. + +--- + +## Notable Technical Benefits and Considerations + +### Benefits + +**Maximum Performance Potential**: By explicitly adjusting memory access patterns and tiling strategies that are difficult for automatic compilers to discover, very fast kernels can be created. + +**Hardware Abstraction Maintained**: You still work within the Python language and JAX ecosystem, and can use features like `jit` and `grad` as-is. + +**Multi-Backend Support**: Supports both GPUs and TPUs, providing flexibility to run the same kernel definition on multiple hardware platforms. + +### Considerations + +**Experimental API**: Pallas is still in an experimental stage with frequent changes. Breaking changes may occur with version updates, and some features may be incomplete or throw "unimplemented" errors. + +**Learning Curve**: You need to understand details like memory hierarchy, tiling, and data transfer. Poorly designed kernels can actually cause performance loss. + +**Need for Debugging and Optimization Tools**: Optimization is difficult without profiling, metrics, and system-level observation tools. It's not yet as mature as GPU's CUDA tools. + +--- + +## Conclusion + +As the saying goes, "Know your enemy, know yourself"—by understanding and responding to hardware (especially TPU) through Pallas, it can become a weapon in the performance war. + +Pallas is an important tool that bridges the gap between automated compilation stacks and research or production performance. The combination of Ironwood and the Mosaic compiler enables more effective processing of complex operations required by new generations of models on TPU. + +As more operators with non-standard memory access and high data dependencies—like attention, MoE, and sparse operations—emerge, the importance of such custom kernel capabilities will only grow. + +In the next article, we'll actually write a simple kernel using Pallas and compare performance, experiencing "know your enemy" in practice. + +--- + +## Reference + +- [Pallas Documentation](https://docs.jax.dev/en/latest/pallas/index.html) +- [Pallas TPU Details](https://docs.jax.dev/en/latest/pallas/tpu/details.html) +- [Inside the Ironwood TPU codesigned AI stack](https://cloud.google.com/blog/products/compute/inside-the-ironwood-tpu-codesigned-ai-stack) +- [PyTorch/XLA 2.4 improves Pallas and adds eager mode](https://cloud.google.com/blog/products/ai-machine-learning/pytorch-xla-2-4-improves-pallas-and-adds-eager-mode) +- [TPU Architecture (qsysarch.com)](https://qsysarch.com/posts/tpu-architecture/) +- [TPU v7 Documentation](https://docs.cloud.google.com/tpu/docs/tpu7x?hl=ko) + +--- + +## P.S. HyperAccel is Hiring! + +"Know your enemy, know yourself," but to win every battle, we need many talented people! + +If you're interested in the technologies we work with, please apply at [HyperAccel Career](https://hyperaccel.career.greetinghr.com/ko/guide)! + +HyperAccel has many excellent and brilliant engineers. We're waiting for your application. diff --git a/content/posts/pallas-programming-model/index.ko.md b/content/posts/pallas-programming-model/index.ko.md new file mode 100644 index 0000000..dad0e22 --- /dev/null +++ b/content/posts/pallas-programming-model/index.ko.md @@ -0,0 +1,375 @@ +--- +date: '2026-01-19T14:50:09+09:00' +draft: false +title: '지피지기면 백전불태 2.5편: Pallas Programming Model' +cover: + image: "images/00-main-image.png" + # can also paste direct link from external site + # ex. https://i.ibb.co/K0HVPBd/paper-mod-profilemode.png + alt: "Pallas Programming Model" + caption: "" + relative: true # To use relative path for cover image, used in hugo Page-bundles +authors: ["Donghyeon Choi"] # must match with content/authors +tags: ["Pallas", "TPU", "Google", "Ironwood", "커널", "커스텀 커널", "Programming Model"] +categories: ["AI Hardware", "Computer Architecture"] +series: ["지피지기면 백전불태"] +summary: "TPU에서 커스텀 커널을 작성할 수 있게 해주는 Pallas 프로그래밍 모델에 대해 알아봅니다." +comments: true +description: "지피지기면 백전불태 2편에서 살펴본 TPU 아키텍처를 바탕으로, TPU에서 커스텀 커널을 작성할 수 있게 해주는 Pallas 프로그래밍 모델에 대해 알아봅니다." +keywords: [ + "Pallas", "TPU", "Google", "Ironwood", "커스텀 커널", "커널 프로그래밍", + "AI Hardware", "Programming Model", "TPU 최적화" +] +--- + +# 지피지기면 백전불태 2.5편: Pallas Programming Model + +> **지피지기면 백전불태(知彼知己 百戰不殆)** +> 상대를 알고 나를 알면 백 번 싸워도 위태롭지 않다는 뜻입니다. +> 이 시리즈는 AI 가속기 설계를 위해 경쟁사들의 하드웨어를 깊이 이해하는 것을 목표로 합니다. +> 2.5편에서는 TPU에서 커스텀 커널을 작성할 수 있게 해주는 **Pallas** 프로그래밍 모델에 대해 다룹니다. + +2편에서 TPU의 하드웨어 아키텍처와 소프트웨어 스택에 대해 살펴봤습니다. 특히 XLA 컴파일러가 어떻게 연산을 최적화하는지, 그리고 그 한계가 무엇인지 다뤘습니다. 그리고 마지막에 다음 글(3편)에서 Groq의 LPU를 소개하겠다고 예고했습니다. + +그런데 TPU를 제대로 이해하려면 하드웨어와 소프트웨어 스택만으로는 부족합니다. 특히 최신 TPU 세대인 Ironwood에서는 **Pallas**라는 프로그래밍 모델이 성능 최적화의 핵심 역할을 하고 있습니다. Pallas는 TPU에서 CUDA나 Triton처럼 커스텀 커널을 작성할 수 있게 해주는 도구로, 하드웨어의 세부 구조를 직접 제어할 수 있으면서도 Python 환경에서 비교적 쉽게 사용할 수 있다는 점이 특징입니다. + +2편에서 XLA 컴파일러의 한계를 언급했듯이, 자동 컴파일러만으로는 최신 알고리즘들을 최적화하기 어려운 경우가 많습니다. Pallas는 바로 이런 한계를 넘어서기 위해 구글이 만든 커널 언어입니다. Ironwood와 함께 공개되면서 "Extreme performance: Custom kernels via Pallas"라는 슬로건으로 강조되고 있을 정도로, TPU의 성능을 극대화하는 데 중요한 역할을 합니다. + +따라서 3편에서 Groq의 LPU를 다루기 전에, 2.5편으로 Pallas 프로그래밍 모델을 한번 짚고 넘어가겠습니다. 이번 글에서는 Pallas가 무엇인지, 왜 필요한지, 그리고 어떻게 TPU의 성능을 끌어올리는지 알아봅니다. + +--- + +## Background: TPU 아키텍처와 Pallas의 필요성 + +Pallas를 이해하기 위해서는 먼저 TPU가 CPU나 GPU와 어떻게 다른지 알아야 합니다. + +--- + +### 전통적인 구조 vs TPU의 Systolic Array + +전통적인 CPU나 GPU 아키텍처는 연산을 위해 SRAM에서 레지스터 파일로 데이터를 매번 가져오고 저장하는 과정을 반복합니다. 이 방식은 행렬곱처럼 동일한 데이터를 반복해서 사용하는 연산에서 막대한 **메모리 대역폭 낭비**를 초래합니다. + +TPU는 이를 해결하기 위해 **한 번 로드된 데이터가 연산기 사이를 흐르며 연속적으로 계산되는 구조**를 채택했습니다. 이것이 바로 2편에서 설명한 Systolic Array 구조입니다. + +![TPU v1 아키텍처 다이어그램](images/01-tpuv1-arch.png) + +TPU 내에서의 연산은 다음과 같은 흐름을 가집니다: + +1. **Weight 공급**: DRAM에 저장된 가중치 데이터가 Weight FIFO를 거쳐 MXU(Matrix Multiply Unit)로 로드됩니다. +2. **Activation 공급**: Unified Buffer(UB)에 저장된 입력 데이터가 MXU로 전달됩니다. +3. **행렬 연산 수행**: MXU 내부의 Systolic Array 구조를 통해 곱셈과 누적(Multiply-Accumulate) 연산이 동시다발적으로 일어납니다. +4. **후처리 파이프라인**: MXU의 출력값은 Accumulator, Activation Unit(ReLU 등), Normalize/Pool Unit을 차례로 통과하며 AI 모델에 필요한 추가 연산을 수행합니다. +5. **결과 저장**: 모든 연산 플로우를 마친 최종 출력은 다시 Unified Buffer로 돌아와 저장됩니다. + +--- + +### GPU와의 핵심 차이점: 실행 단위의 변화 + +이러한 구조적 차이는 프로그래밍 모델의 관점에서 매우 큰 차이를 만듭니다. + +- **GPU (SIMT)**: 수많은 스레드가 독립적으로 계산을 수행하며, 32개의 스레드를 워프(Warp) 단위로 묶어 관리합니다. 데이터 하나하나에 집중하는 세밀한 병렬 처리가 특징입니다. +- **TPU (Tensor-centric, SPMD)**: 한 번의 데이터 로드로 전체 연산 시퀀스를 끝내버립니다. 즉, 개별 스레드가 아닌 **'텐서(또는 타일) 하나에 대한 전체 프로그램'을 하나의 최소 실행 단위**로 간주합니다. + +> **Why Pallas?** Pallas는 바로 이러한 TPU의 하드웨어적 특성(Unified Buffer 직접 제어, MXU 스케줄링 등)을 추상화하면서도, 개발자가 직접 이 '텐서 단위의 흐름'을 최적화할 수 있도록 설계된 언어입니다. + +![TPU v7 Ironwood 아키텍처 다이어그램](images/02-tpuv7-arch.png) + +TPU는 세대를 거듭하며 벡터 연산 유닛의 강화와 대규모 모델 대응을 위한 스케일업 기술을 발전시켜 왔습니다. 특히 최신 아키텍처인 Ironwood에 이르러서는 칩렛(Chiplet) 구조를 채택한 거대한 패키지 위에 **256x256 크기의 Systolic Array(MXU)를 총 4개**나 탑재한 "괴물 연산기"가 되었습니다. + +하지만 하드웨어의 체급이 비약적으로 커졌음에도 불구하고, **한 번의 데이터 로드를 통해 최대한 많은 연산을 수행한다**는 핵심 설계 기조는 변함없이 유지되고 있습니다. + +--- + +### 왜 Pallas가 필요한가? + +2편에서 XLA 컴파일러에 대해 다뤘을 때, XLA는 강력한 최적화 컴파일러지만 한계가 있다고 했습니다. 새로운 연산 알고리즘이 등장하면, 컴파일러가 이를 최적화할 수 있는 버전으로 업데이트되기 전까지는 수동으로 만든 커스텀 커널의 성능을 따라잡기 어렵습니다. + +예를 들어, Flash Attention이나 MoE(Mixture of Experts) 같은 최신 알고리즘들은 메모리 접근 패턴이 복잡하거나 데이터 의존성이 높아서, 자동 컴파일러가 최적화하기 어려운 경우가 많습니다. 이런 경우에는 개발자가 직접 메모리 계층 구조를 이해하고, 데이터를 어떻게 타일링할지, 언제 메모리 간 데이터 이동을 할지 등을 세밀하게 제어해야 합니다. + +GPU에서는 이런 문제를 CUDA나 Triton으로 해결했습니다. CUDA는 하드웨어를 직접 제어할 수 있지만 진입장벽이 높고, Triton은 CUDA보다 추상화 수준이 높아서 비교적 쉽게 사용할 수 있습니다. 하지만 둘 다 GPU에서만 작동합니다. + +TPU에서는 어떻게 할까요? 구글은 2023년 무렵부터 JAX의 실험적인 확장 기능으로 **Pallas**라는 커널 언어 API를 제공하기 시작했습니다. Pallas는 Triton과 비슷한 철학을 가지고 있지만, GPU와 TPU 양쪽을 모두 지원한다는 점이 큰 차이입니다. + +--- + +## Pallas란 무엇인가? + +Pallas는 JAX 생태계 내에서 커스텀 커널을 작성할 수 있게 해주는 커널 언어입니다. GPU와 TPU 양쪽에서 작동하며, 하드웨어의 메모리 계층 구조, 데이터 타일링, 블록 단위 병렬성 등을 직접 제어할 수 있게 해줍니다. + +Pallas의 핵심 아이디어는 간단합니다. 높은 수준의 자동화된 컴파일러가 처리하기 어려운 연산들을, 개발자가 직접 하드웨어에 가까운 수준에서 제어할 수 있게 해주는 것입니다. 하지만 CUDA처럼 완전히 낮은 수준으로 내려가는 것은 아니고, Python 환경에서 비교적 쉽게 사용할 수 있도록 추상화되어 있습니다. 이는 자동 컴파일러의 편의성과 수동 최적화의 제어력을 모두 제공하는 균형점을 찾은 것입니다. + +--- + +### 핵심 추상화: Grid, BlockSpec, Ref + +Pallas는 세 가지 핵심 추상화를 통해 하드웨어를 제어합니다: **Grid**, **BlockSpec**, 그리고 **Ref**입니다. + +> **Grid: 병렬 실행 추상화** + +![Grid와 Program ID: Grid는 병렬 실행의 반복 공간을 지정하며, 각 프로그램은 고유한 program_id를 가집니다](images/04-grid-and-block.png) + +Pallas는 **Grid**라는 추상화를 통해 실행 단위를 모델링합니다. Grid는 하드웨어별로 다른 의미를 가지지만, 통일된 인터페이스를 제공합니다. + +```python +def kernel(o_ref): + i = pl.program_id(0) # 현재 프로그램의 ID + o_ref[i] = i + +# Grid 크기 지정: (8,) = 8개의 병렬 프로그램 +result = pl.pallas_call( + kernel, + out_shape=jax.ShapeDtypeStruct((8,), jnp.int32), + grid=(8,), # 8개의 병렬 인스턴스 +)() +``` + +Pallas에서 **Grid**는 전체 작업을 잘게 쪼갠 실행 단위들의 집합이지만, 하드웨어가 이 Grid를 처리하는 방식은 아키텍처에 따라 근본적인 차이가 있습니다. + +  + +> **GPU에서의 Grid** + +GPU 백엔드(Triton/Mosaic)에서 Grid는 하드웨어 스케줄러에 의한 **완전 병렬 실행**을 전제로 합니다. 각 Grid 아이템은 하나의 **스레드 블록(Thread Block)** 에 매핑되어 개별 SM(Streaming Multiprocessor)에서 독립적으로 실행됩니다. + +하드웨어가 비결정적으로 스레드를 스케줄링하므로 Grid 간의 실행 순서가 보장되지 않습니다. 이는 개발 시 중요한 고려사항입니다. 서로 다른 프로그램이 동일한 HBM 위치에 쓰지 않도록 `BlockSpec`의 `index_map` 설계 시 경쟁 조건(Race condition)을 엄격히 관리해야 합니다. + +  + +> **TPU에서의 Grid** + +![TPU에서의 Grid와 Block: TPU 백엔드에서 Grid가 어떻게 병렬성과 파이프라이닝을 조합하는지 보여주는 다이어그램](images/05-grid-block-in-tpu.png) + +TPU 백엔드에서 Grid는 다중 코어 간의 병렬성과 단일 코어 내의 **순차적 파이프라이닝**을 조합한 모델입니다. + +TPU는 매우 넓은 SIMD 머신이지만, 소프트웨어적으로는 **단일 스레드 프로세서**처럼 코딩할 수 있습니다. 컨트롤유닛(TCS)이 루프를 돌며 전체 연산을 제어하는 직관적인 흐름을 제공합니다. + +여러 개의 TensorCore가 있을 경우 작업을 분산하여 **물리적으로 동시에** 실행하는 Parallel 차원과, 단일 TensorCore 내에서 **직렬 실행**을 보장하는 Sequential 차원이 있습니다. + +순차 차원은 단순히 느리게 실행하는 것이 목적이 아닙니다. **세마포어(Semaphore)** 를 통해 현재 연산과 다음 데이터의 로드를 중첩(Overlap)시켜 **메모리 레이턴시를 숨기기 위한 전략적 수단**으로 사용됩니다. +  + +> **BlockSpec: 메모리 레이아웃 추상화** + +Pallas는 **BlockSpec**을 통해 거대한 데이터를 하드웨어가 소화할 수 있는 크기(Chunk)로 쪼개는 과정을 추상화합니다. 이는 단순히 데이터를 자르는 것을 넘어, **HBM(Remote)과 SRAM(Local) 사이의 데이터 전송 방식 설정**입니다. + +BlockSpec은 세 가지 구성 요소로 이루어져 있습니다: + +- **`block_shape`**: 각 프로그램 인스턴스가 Local Memory(SRAM) 작업대 위에 올려둘 데이터의 크기입니다. 하드웨어의 SRAM 용량을 초과하지 않도록 설계하는 것이 성능 최적화의 핵심입니다. +- **`index_map` 함수**: `grid` 인덱스(i, j)를 입력받아 HBM 상의 블록 시작 위치를 반환합니다. 이 함수는 컴파일 시점에 분석되어 하드웨어의 DMA(Direct Memory Access) 주소 계산 로직으로 변환됩니다. +- **`memory_space`**: 조각난 데이터가 머물 물리적 그릇을 지정합니다. 지정하지 않을 경우 백엔드 설정에 따라 기본 공간(`pl.SRAM`)으로 할당됩니다. + +  + +> **Ref: 메모리 참조 추상화** + +Pallas는 `Ref` 객체를 통해 복잡한 하드웨어 메모리 주소 체계를 추상화합니다. 이는 단순히 데이터를 가리키는 포인터를 넘어, **SRAM(Local Memory) 상의 특정 데이터 블록에 대한 논리적 뷰**를 제공합니다. + +`Ref`의 핵심 특징은 다음과 같습니다: + +- **Local Memory 참조**: `Ref`는 HBM이 아니라, 하드웨어의 가장 빠른 작업대인 **SRAM(TPU의 VMEM/SMEM, GPU의 Shared Memory)**에 올라온 데이터를 가리킵니다. +- **Dereferencing (역참조)**: `x_ref[...]`와 같이 대괄호를 사용하는 시점에 비로소 **SRAM → Register File**로의 실제 데이터 로드가 발생하며, 값(`Value`)으로 변환됩니다. +- **하드웨어 추상화**: 동일한 `Ref` 인터페이스를 사용하더라도, 백엔드에 따라 TPU의 Vector/Scalar 메모리나 GPU의 Shared Memory 접근으로 자동 변환됩니다. 명확한 최적화를 위해 직접 설정할 수도 있습니다. + +--- + +### Pallas Hardware Modeling + +Pallas는 TPU와 GPU 각각에 대하여 큰 그림은 같으나 세부적으로는 다른 방향으로 하드웨어 모델을 제공합니다. 먼저 공통적으로 구조화되어 있는 전체적인 구조를 살펴봅시다. + +![Pallas 공통 하드웨어 모델 계층 구조: Remote Memory(HBM), Multiple Core structure, Local Memory(SRAM), Register Files](images/06-common-hw-model.png) + +Pallas가 추상화하는 공통 하드웨어 모델은 다음과 같은 계층 구조를 가집니다: + +- **Remote Memory (HBM)**: 고대역폭 메모리(High Bandwidth Memory)로, 가장 느리지만 용량이 큰 메모리 공간입니다. `pltpu.HBM` 등을 통해 명시적으로 지정하거나 `pl.ANY`를 통해 자동으로 지정할 때 주로 사용됩니다. +- **Multiple "Core" structure**: 그리드를 병렬적으로 처리하기 위한 독립적 연산 유닛들의 집합입니다. 각 코어는 독립적으로 연산을 수행할 수 있습니다. + - **Local Memory (SRAM 혹은 Cache)**: 코어 내부에 위치한 빠른 메모리입니다. Remote Memory보다 훨씬 빠르지만 용량은 제한적입니다. HBM상의 데이터와 일치성을 위해 READ용 변수와 WRITE용 변수를 분리하는 것을 권장합니다. + - **Register Files in Execution Units**: 연산 유닛과 직접적으로 데이터를 주고 받을 수 있는 가장 빠른 메모리입니다. 연산에 필요한 데이터를 즉시 제공합니다. + +일반적으로 Local Memory는 Register File의 10배의 레이턴시가 소요되고, Remote Memory는 다시 Local Memory의 10배의 레이턴시가 소요된다고 볼 수 있습니다. 그러나 Chiplet 구조를 채택한 TPU v7 Ironwood에서는 하나의 칩 안에 HBM까지 모두 가지고 있기 때문에 Remote Memory 통신이 기존 10배에서 2~5배로 빨라질 수 있습니다. + +  + +> **Pallas의 TPU 하드웨어 모델** + +![Pallas TPU 하드웨어 모델링: HBM → VMEM/SMEM → Register File로의 데이터 흐름과 파이프라이닝 과정](images/07-tpu-hw-model.png) + +TPU의 하드웨어 모델은 공통 모델을 기반으로 하되, TPU만의 특수한 구조를 가집니다: + +**메모리 계층 구조**: + +- **Local Memory**는 Vector Memory와 Scalar Memory로 분리되어 있습니다: + - **Vector Memory(VMEM)**: Vector 및 Matrix 관련 연산을 위한 데이터를 저장하는 메모리입니다. 동일 텐서 코어 내의 Vector/Matrix Unit (VPU/MXU/XLU)에서 접근 가능합니다. + - **Scalar Memory(SMEM)**: Logic flow(loop, condition 등)를 위한 scalar 연산 관련 데이터를 저장하는 메모리입니다. 동일 텐서 코어 내의 Scalar Unit (TCS)에서 접근 가능합니다. TCS에서 생성하는 고수준의 명령을 통해 SMEM 상에 저장된 데이터 또한 직접적으로 벡터 연산에 사용될 수 있습니다. + +- 각 VMEM과 SMEM에는 READ/WRITE를 위한 영역을 따로 배치하여 HBM과의 data sync를 확보합니다. READ로 불러와진 데이터를 수정 시 다시 write-back을 안 하기 때문에 HBM과의 데이터 불일치가 발생할 수 있어서입니다. + +**데이터 흐름과 파이프라이닝**: + +- 커널 호출 시 Remote Memory(HBM)에서 Local Memory(VMEM/SMEM)로 지정된 BlockSpec에 맞추어 데이터를 로드합니다. 이때 설정된 Grid 번호에 맞추어 해당하는 데이터를 자동으로 가져옵니다. + +- 커널 내에서 Ref 데이터 호출 시 Local Memory에서 Register File로 데이터를 가져가 연산을 진행합니다. 이때 Local Memory ↔ Register File 간 이동 및 연산 작업 수행 중 다음 Grid에 대한 Remote Memory → Local Memory 데이터 로드가 자동으로 겹쳐져서 진행됩니다. 이를 통해 메모리 레이턴시를 상당량 감출 수 있습니다. + +- 커널에서 설정된 출력 Ref 데이터에 결과 데이터를 쓰면 Local Memory로 write-back되며, 이후 해당 그리드에 대한 모든 프로그램 종료 시에 출력 Ref 데이터가 Remote Memory에 write-back됩니다. + +  + +> **Pallas의 GPU 하드웨어 모델** + +![Pallas GPU 하드웨어 모델링: HBM → SMEM → Register File로의 데이터 흐름과 파이프라이닝 과정](images/08-gpu-hw-model.png) + +GPU 또한 하드웨어 모델은 공통 모델을 기반으로 하되, GPU만의 구조를 가집니다: + +**메모리 계층 구조**: + +- **Local Memory는 Shared Memory(SMEM)**: 각 SM(Streaming Multiprocessor)별로 독립적인 Shared Memory/L1 Cache 공간을 가집니다. TPU와 달리 Scalar/Vector 구분이 없는 통합된 작업 공간이며, 수많은 스레드가 이 공간에 동시 접근하여 병렬 연산합니다. + +**데이터 흐름과 파이프라이닝**: + +- 단순히 하드웨어 스케줄링에만 의존하는 기존 방식과 달리, Pallas GPU는 **HBM(Remote) → SMEM(Local) 데이터 로드와 TensorCore 연산을 명시적으로 중첩(Overlap)** 합니다. `plgpu.emit_pipeline` 등을 통해 연산 유닛이 현재 데이터를 처리하는 동안, 비동기적으로 다음 그리드의 데이터를 미리 가져와 메모리 레이턴시를 감춥니다. +- 커널 내부의 Ref 변수는 HBM이 아닌, 이미 Local Memory(SMEM)에 올라왔거나 올라올 예정인 데이터 조각의 주소입니다. `memory_space=plgpu.GPUMemorySpace.GMEM` 설정을 통해 HBM(Remote) 공간을 명시하고, 이를 `copy_gmem_to_smem`으로 직접 제어할 수 있습니다. + +--- + +### 메모리 파이프라이닝과 세마포어 + +TPU 아키텍처에서 HBM(Remote Memory)과 VMEM(Local Memory) 간의 데이터 이동은 수백 사이클의 레이턴시를 발생시킵니다. MXU가 아무리 빠르게 행렬 연산을 수행해도, 데이터가 도착할 때까지 기다려야 한다면 전체 성능은 메모리 병목에 갇히게 됩니다. + +Pallas는 이 문제를 하드웨어 수준의 **세마포어(Semaphore)** 를 활용한 파이프라이닝으로 해결합니다. + +> **Pallas의 동기화 메커니즘: Semaphore 기반의 비동기 제어** + +Pallas에서 DMA와 연산 유닛은 서로 독립적으로 작동하며, 이들의 속도 차이를 조율하기 위해 **세마포어(Semaphore)** 를 사용합니다. 동작 원리는 다음과 같이 상호 신호 체계로 이루어집니다. + +- **Data Load (Producer)**: DMA가 외부 메모리(HBM)에서 로컬 메모리(VMEM/SMEM)로 데이터 로드를 완료하면, 세마포어 값을 증가(Signal)시켜 데이터가 준비되었음을 알립니다. +- **Compute (Consumer)**: 연산 유닛은 세마포어 값을 확인하며 데이터가 로드될 때까지 대기(Wait)합니다. 값이 충족되면 즉시 연산을 시작합니다. +- **Feedback Loop**: 연산이 완료되면 연산 유닛은 다시 세마포어를 Signal하여 DMA에게 버퍼가 비었음을 알리고, 다음 데이터를 로드할 수 있도록 허용합니다. + +이러한 Wait-Signal 구조를 통해 데이터 로드와 연산을 겹쳐서 실행(Overlapping)함으로써 하드웨어 가동률을 극대화합니다. + +  + +> **Double Buffering: 레이턴시를 숨기는 핵심 기법** + +세마포어를 활용한 가장 기본적인 파이프라이닝 기법이 Double Buffering입니다. Local Memory 내에 두 개의 버퍼(Buffer 0, Buffer 1)를 할당하고, 다음과 같이 동작합니다: + +1. **초기화**: DMA가 첫 번째 데이터를 Buffer 0에 로드 +2. **루프 시작**: VPU가 Buffer 0을 연산하는 **동시에**, DMA는 다음 데이터를 Buffer 1에 로드 +3. **버퍼 교체**: 연산과 로드가 모두 완료되면 역할 교대 - VPU는 Buffer 1을, DMA는 Buffer 0에 로드 +4. **반복**: 이 과정을 데이터가 끝날 때까지 반복 + +결과적으로 **메모리 로드 시간이 연산 시간 뒤로 완전히 숨겨집니다**. 연산이 데이터 로드보다 오래 걸리는 compute-bound 워크로드에서 특히 효과적입니다. + +Pallas에서는 `pltpu.semaphore`를 통해 이러한 동기화를 명시적으로 제어할 수 있습니다. 'Read-before-ready(데이터 도착 전 읽기)'와 'Write-over-active(연산 중 덮어쓰기)' 같은 위험한 상황을 하드웨어 수준에서 방지하며, 컴파일러가 자동으로 **Prefetch와 Overlap을 최적화**합니다. + +--- + +### Backend Lowering과 JAX 통합 + +![Pallas 프로그래밍 모델 전체 구조: Python 커널이 JAX Program Representation을 거쳐 TPU용 Mosaic 또는 GPU용 Triton/Mosaic GPU로 lowering되는 과정](images/09-pallas-lowering.png) + +Pallas로 작성된 커널은 최종적으로 하드웨어 코드로 변환됩니다. GPU에서는 Triton 또는 Mosaic GPU 백엔드를 통해, TPU에서는 Mosaic 컴파일러를 통해 MLIR 형태로 낮춰지고, 최종적으로 하드웨어 코드로 변환됩니다. 이 과정에서 연산자 fusion, 타일링 자동화, 데이터 전송과 계산의 오버랩(overlap) 등이 최적화됩니다. 개발자가 작성한 고수준 코드가 하드웨어에 최적화된 코드로 변환되는 것입니다. + +또한 Pallas 커널은 JAX의 `jit`, `vmap`, `grad` 같은 변환(Transform)과 호환됩니다. 따라서 고성능 커널을 작성하면서도 자동 미분, 매핑, 컴파일 등의 기능을 여전히 활용할 수 있다는 것이 큰 장점입니다. 이는 Pallas가 단순한 커널 언어를 넘어 JAX 생태계와 완전히 통합된 도구임을 보여줍니다. + +--- + +## CUDA vs Pallas: 프로그래밍 모델 비교 + +1편에서 살펴본 CUDA와 Pallas를 비교해봅니다. + +### 핵심 비교 + +| 구분 | CUDA | Pallas | +|:---:|:---|:---| +| **추상화** | 스레드 중심 (Thread → Warp → Block → Grid) | 데이터 중심 (Grid + BlockSpec + Ref) | +| **메모리 제어** | 명시적 (`__shared__`, `__global__`) | 선언적 (`memory_space`, 자동 매핑) | +| **동기화** | 수동 (`__syncthreads()`) | 자동 (세마포어 기반 파이프라이닝) | +| **하드웨어** | NVIDIA GPU 전용 | GPU + TPU 지원 | +| **생태계** | 성숙 (Nsight, cuBLAS, cuDNN) | 실험적 (JAX 통합) | +| **언어** | C/C++ | Python | + +### 코드 비교: 벡터 덧셈 + +**CUDA** - 스레드 ID 직접 계산, 경계 체크 필요: +```cuda +__global__ void vector_add(float *A, float *B, float *C, int N) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < N) C[idx] = A[idx] + B[idx]; +} +``` + +**Pallas** - 데이터 블록 단위 추상화: +```python +def vector_add_kernel(a_ref, b_ref, c_ref): + c_ref[...] = a_ref[...] + b_ref[...] + +result = pl.pallas_call(vector_add_kernel, out_shape=..., grid=(N,))(a, b) +``` + +### 선택 가이드 + +| CUDA 선택 | Pallas 선택 | +|:---|:---| +| NVIDIA GPU 최대 성능 필요 | GPU/TPU 이식성 필요 | +| 프로덕션 안정성 중요 | 빠른 프로토타이핑 | +| 기존 CUDA 코드베이스 활용 | JAX 생태계 (`jit`, `vmap`, `grad`) 활용 | +| 성숙한 디버깅/프로파일링 도구 필요 | TPU 커스텀 커널 작성 | + +--- + +## TPU에서 Pallas가 어떻게 활용되는가 + +최신 TPU 세대인 Ironwood에서는 Pallas를 중심에 두고 "Extreme performance: Custom kernels via Pallas"라는 슬로건으로 강조하고 있습니다. 특히 Ironwood 스택에서는 Pallas 정의를 통해 개발자가 Python 속에서 메모리 타일링, 데이터 이동, MXU 활용에 대한 전략을 직접 기술할 수 있게 하고, Mosaic 컴파일러가 이를 TPU 코드로 낮춰서 구현합니다. + +이 조합을 통해 데이터 고착(input stationary), 가중치 고착(weight stationary), 출력 고착(output stationary) 등의 타일링 전략과 스트라이드, 배치(batch) 단위의 분산 처리가 효율적으로 설계됩니다. 또한, 동시에 HBM ↔ 온칩 메모리 데이터 전송과 MXU 계산을 겹치게 하면서 전체 파이프라인의 스케줄링 병목을 최소화하는 설계가 가능해졌습니다. + +- **최대 성능 잠재력 실현**: 자동 컴파일러가 발견하기 힘든 메모리 액세스 패턴, 타일링 전략 등을 명시적으로 조정할 수 있기 때문에, 매우 빠른 커널을 만들 수 있습니다. + +- **하드웨어 추상화 유지**: 여전히 Python 언어, JAX 생태계 내에서 작업하며, `jit`, `grad` 등의 기능을 그대로 사용할 수 있습니다. + +- **다중 백엔드 대응**: GPU와 TPU 양쪽을 지원하며, 동일한 커널 정의로 여러 하드웨어에서 동작하게 할 수 있는 유연성이 있습니다. + +### 고려해야 할 사항 + +- **실험 단계**: Pallas는 아직 자주 변경되는 실험적 단계에 있습니다. 버전 업에 따라 breaking changes가 있을 수 있고, 일부 기능이 완전하지 않거나 "미구현(unimplemented)" 에러가 나올 수 있습니다. + +- **학습 곡선 및 진입 장벽**: 메모리 계층, 타일링, 데이터 전송 등의 세부사항을 이해해야 합니다. 잘못 설계된 커널은 오히려 성능 손실을 가져올 수 있습니다. + +- **디버깅 및 최적화 도구 필요성**: 프로파일링, 메트릭, 시스템 레벨 관측 도구들 없이 최적화를 진행하기 어렵습니다. 아직 GPU의 CUDA 도구들만큼 성숙하지는 않습니다. + +--- + +## 마치며 + +이번 글에서는 TPU에서 커스텀 커널을 작성할 수 있게 해주는 **Pallas** 프로그래밍 모델에 대해 살펴봤습니다. + +먼저, TPU의 Systolic Array 구조가 CPU/GPU와 어떻게 다른지, 그리고 이러한 하드웨어적 특성이 프로그래밍 모델에 어떤 영향을 미치는지 확인했습니다. TPU는 텐서 단위의 실행 모델을 채택하여, 한 번의 데이터 로드로 전체 연산 시퀀스를 수행하는 구조적 특징을 가지고 있습니다. + +XLA 컴파일러가 자동으로 최적화하기 어려운 영역을 위해 등장한 Pallas는 Grid, BlockSpec, Ref라는 핵심 추상화를 통해 하드웨어의 메모리 계층 구조와 데이터 타일링을 직접 제어할 수 있게 해줍니다. 특히 TPU와 GPU 각각의 특성에 맞춘 메모리 모델(TPU의 VMEM/SMEM, GPU의 Shared Memory)을 제공하면서도, Python 환경에서 비교적 쉽게 사용할 수 있는 추상화 수준을 유지합니다. + +CUDA와의 비교를 통해 Pallas가 더 높은 추상화 수준을 제공하면서도 하드웨어 제어력을 유지한다는 점을 확인했고, Ironwood 세대에서 Pallas가 "Extreme performance"를 위한 핵심 도구로 자리잡고 있음을 살펴봤습니다. + +Pallas는 자동화된 컴파일러와 수동 최적화 사이의 간극을 메우는 중요한 도구로, 최신 알고리즘들이 요구하는 복잡한 메모리 접근 패턴을 효율적으로 처리할 수 있게 해줍니다. 특히 Flash Attention, MoE, Sparse 연산처럼 비표준적인 연산 패턴이 중요해질수록, Pallas와 같은 커스텀 커널 도구의 역할은 더욱 커질 것입니다. + +이제 다시 본 시리즈로 돌아가, 3편에서는 Groq의 LPU(Language Processing Unit)를 분석하며, 또 다른 AI 가속기의 설계 철학과 프로그래밍 모델을 살펴보겠습니다. + +--- + +## Reference + +- [Pallas Documentation](https://docs.jax.dev/en/latest/pallas/index.html) +- [Pallas TPU Details](https://docs.jax.dev/en/latest/pallas/tpu/details.html) +- [Inside the Ironwood TPU codesigned AI stack](https://cloud.google.com/blog/products/compute/inside-the-ironwood-tpu-codesigned-ai-stack) +- [PyTorch/XLA 2.4 improves Pallas and adds eager mode](https://cloud.google.com/blog/products/ai-machine-learning/pytorch-xla-2-4-improves-pallas-and-adds-eager-mode) +- [TPU Architecture (qsysarch.com)](https://qsysarch.com/posts/tpu-architecture/) +- [TPU v7 Documentation](https://docs.cloud.google.com/tpu/docs/tpu7x?hl=ko) + +--- + +## 추신: HyperAccel은 채용 중입니다! + +지피지기면 백전불태라지만 백전백승을 위해서는 훌륭한 인재가 많이 필요합니다! + +저희가 다루는 기술들을 보시고, 관심이 있으시다면 [HyperAccel Career](https://hyperaccel.career.greetinghr.com/ko/guide)로 지원해 주세요! + +HyperAccel에는 정말 훌륭하고 똑똑한 엔지니어분들이 많습니다. 여러분의 지원을 기다립니다.