CUDA C++ සහ Triton තුළ GEMM, torch.compile සහ ROCm අතේ ගෙන යා හැකි බව
එම පෙරමුණු ක්ෂේත්‍ර තුන—CUDA C++ එදිරිව Triton තුළ GEMM ක්‍රියාත්මක කිරීම, torch.compile (TorchInductor) ස්වයංක්‍රීය කේත ජනනය සහ AMD ROCm බහු-විකුණුම්කරු අතේ ගෙන යා හැකි බව—පිළිබඳ විශ්ලේෂණාත්මක ගැඹුරු ගමනක් නවීන AI පද්ධති පිටුපස ඇති ප්‍රායෝගික ඉංජිනේරු තරඟකාරී තීරණ හෙළි කරයි.
1. පැති-පැති විශ්ලේෂණය: CUDA C++ එදිරිව Triton තුළ Matrix Multiplication (GEMM)
Matrix multiplication (C = A × B) GPU ක්‍රමලේඛනයේ අවසාන මිණුම් සලකුණ වන්නේ එය compute-සීමිත වන නිසාත් මතකයේ නවතිනු නොලැබ Tensor Cores වලට අඛණ්ඩව සැපයිය යුතු නිසාත්ය.
- නවීන CUDA C++ (CUTLASS / CuTe ප්‍රවේශය):
- අතින් දෘඩාංග සම්බන්ධීකරණය: CUDA C++ තුළ උපරිම GEMM වේගය ලබා ගැනීමට කටයුතු දැඩි ධරාසූතිකයකට බෙදිය යුතුය: Grid → Thread Block (CTA) → Warpgroup (4 warps) → Warp → Tensor Core උපදෙස (wgmma).
- පැහැදිලි Pipelining & Swizzling: ඉංජිනේරුවන් අතින් double- හෝ triple-buffered Shared Memory අදියර (cp.async හෝ TMA) වෙන් කළ යුත්තේ Tensor Cores Tile N ගුණ කරන අතරතුර අසමමුහුර්ත දෘඩාංග එන්ජින් Tile N+1 පූරණය කිරීමටය. ඔවුන් Shared Memory bank ගැටුම් වළක්වා ගැනීමට XOR memory swizzling ද යෙදිය යුතුය.
- විශ්ලේෂණාත්මක තීන්දුව: සසඳා බැලිය නොහැකි උපරිම කාර්යසාධනය (න්‍යායාත්මක Tensor Core FLOPS වලින් 90%–98%), නමුත් පේළි සිය ගණනක් බිඳී යා හැකි, ගෘහ නිර්මාණයට-විශේෂිත C++ template කේත අවශ්‍ය වේ.
- OpenAI Triton ප්‍රවේශය:
- ප්‍රකාශනීය Tile ගණිතය: Triton තුළ, ඉහළ කාර්යසාධන tiled GEMM එකක් Python පේළි ආසන්න 40ක් ගනී. ඉංජිනේරුවා අභ්‍යන්තර මානය K මත loop එකක් ලියමින්, 2D blocks පූරණය කර (tl.load), අර්ධ එකතු එකතු කිරීමට tl.dot(a, b) අමතයි.
- Compiler-මත පදනම් වූ Pipelining: Triton ගේ MLIR backend එක software pipelining අදියර (num_stages=3) ස්වයංක්‍රීයව අනුමාන කර, Shared Memory swizzling ඇතුළත් කර, Tensor Core උපදෙස් නිකුත් කරයි. මීට අමතරව, @triton.autotune ඉලක්ක GPU මත විවිධ tile ප්‍රමාණ (BLOCK_M, BLOCK_N, BLOCK_K) පරීක්ෂා කර වේගවත්ම වින්‍යාසය ස්වයංක්‍රීයව තෝරා ගනියි.
- විශ්ලේෂණාත්මක තීන්දුව: 10× ක් අඩු කේතයකින් CUTLASS කාර්යසාධනයෙන් 85%–95% ක් ලබා දෙන බැවින්, සම්මත පුස්තකාල ප්‍රමාණවත් නොවන අභිරුචි fused GEMMs (උදා: GEMM + Bias + GeLU + Quantization) සඳහා එය උසස් වේ.
2. torch.compile (TorchInductor) විශ්ලේෂණය: PyTorch ස්වයංක්‍රීයව Triton ලියන ආකාරය
PyTorch 2.0 ගේ torch.compile, TorchInductor භාවිතයෙන් ස්වයංක්‍රීයව Triton කේත ජනනය කිරීම මඟින් eager Python නම්‍යශීලීබවය සහ compiled kernel වේගය අතර හිඩැස පුරවයි.
[Eager Python Code]
│ 1. TorchDynamo (Bytecode Capture)
▼
[FX Graph (ATen Ops)]
│ 2. AOTAutograd (Forward + Backward Decomposition)
▼
[TorchInductor Loop IR]
│ 3. Operator Fusion & Memory Layout Analysis
▼
[Generated Triton Python Kernel] ──► [PTX / SASS Binary]- ස්වයංක්‍රීය Fusion ක්‍රියා කරන ආකාරය:
- ඔබ ආකෘතියක් @torch.compile සමඟ අලංකාර කරන විට, TorchDynamo Python bytecode උදුරා ගෙන Python interpreter පොදු කාර්යයක් නොමැතිව ස්ථිතික ගණනය කිරීමේ ග්‍රාෆ් එකක් ග්‍රහණය කර ගනියි.
- TorchInductor ඉහළ-මට්ටමේ PyTorch operators ~2,000 ක් මූලික primitives ~250 ක් දක්වා පහත හෙළමින්, memory යැපීම් විශ්ලේෂණය කර ආසන්න pointwise (උදා: add, mul, silu) සහ reduction (උදා: sum, mean, softmax) මෙහෙයුම් එකට කාණ්ඩ කරයි.
- වෙනම පෙර-සංයුක්ත C++ kernels අමතනවා වෙනුවට, TorchInductor තැටියේ අලුත් Triton Python script එකක් නිකුත් කර (TORCH_LOGS="output_code" මඟින් දැකිය හැකි), එය Triton ගේ MLIR backend එකෙන් සංයුක්ත කර, ලැබෙන GPU binary එක cache කරයි.
- විශ්ලේෂණාත්මක තරඟකාරී තීරණ (ශක්තීන් එදිරිව සීමාවන්):
- ශක්තිය: සම්මත ආකෘති බාධකවලින් 80%–90% ක් සඳහා අතින් kernel ලිවීම ඉවත් කරමින්, පේළියක කේතයකින් පුහුණුව සහ inference නිතිපතා 1.5× සිට 2.5× දක්වා වේගවත් කරයි.
- සීමාව: සහාය නොදක්වන ගතික Python control flow හමුවන විට graph breaks වලට ලක් වන අතර, සංකීර්ණ ඇල්ගොරිතමික නැවත සැකසුම් (FlashAttention ගේ online Softmax tiling වැනි) ස්වයංක්‍රීයව සොයා ගැනීමට තවමත් පොරබදන අතර, ඒවාට තවමත් අතින් ලියන ලද Triton හෝ CUDA kernels අවශ්‍ය වේ.
3. බහු-විකුණුම්කරු අතේ ගෙන යා හැකි බව පිළිබඳ විශ්ලේෂණය: NVIDIA CUDA එදිරිව AMD ROCm
වසර ගණනාවක් තිස්සේ, NVIDIA ගේ හිමිකාර CUDA පරිසර පද්ධතිය AI compute මත ආසන්න-ඒකාධිකාරයක් නිර්මාණය කළේය. නවීන compiler stacks ඉහළ-මට්ටමේ AI කේතය විකුණුම්කරු-විශේෂිත දෘඩාංග backends වලින් වෙන් කිරීමෙන් මෙම අගුළු වැට ක්‍රමානුකූලව කඩා දමා ඇත.
| අතේ ගෙන යා හැකි බවේ ස්තරය | NVIDIA පරිසර පද්ධතිය | AMD ROCm පරිසර පද්ධතිය | විශ්ලේෂණාත්මක සංසන්දනය |
|---|---|---|---|
| පහළ-මට්ටමේ C++ Kernels | CUDA C++ (nvcc → PTX → SASS) | HIP (hipcc → AMDGPU LLVM → GCN/CDNA) | ඉහළ ඝර්ෂණය. AMD ගේ hipify මෙවලම CUDA syntax HIP බවට පරිවර්තනය කරයි, නමුත් warp ප්‍රමාණ වෙනස් වේ (NVIDIA මත threads 32 එදිරිව AMD CDNA මත threads 64), hardcoded warp-shuffle ගණිතය බිඳ දමමිනි. |
| Block-Level Kernels | OpenAI Triton (NVPTX Backend) | OpenAI Triton (AMDGPU Backend) | ආසන්න-බාධාරහිත. Triton warp ප්‍රමාණ 2D blocks පිටුපස වියුක්ත කරන බැවින්, තනි Triton kernel එකක් NVIDIA H100/B200 සහ AMD Instinct MI300X/MI325X දෙක මතම ස්වදේශීයව ක්‍රියාත්මක වේ. |
| Framework මට්ටම | PyTorch + torch.compile + vLLM | PyTorch + torch.compile + vLLM | පරිපක්ව සමානාත්මතාවය. ප්‍රධාන inference එන්ජින් (vLLM, SGLang) සහ PyTorch, custom ops Triton සහ AMD ගේ rocBLAS/CK (Composable Kernel) හරහා යොමු කිරීමෙන් ROCm මත එළඹිමෙන් ක්‍රියාත්මක වේ. |
- උපායාත්මක නිගමනය: අතින් ලියන ලද CUDA C++ NVIDIA සිලිකන් වලට සීමා වී පවතින අතර AMD HIP වෙත අතින් ගෙන යෑම අවශ්‍ය වන නමුත්, OpenAI Triton අතේ ගෙන යා හැකි දෘඩාංග වියුක්ත ස්තරයක් ලෙස ක්‍රියා කරයි. custom kernels Triton මත සම්මත කරන AI කණ්ඩායම් වෙනම C++ codebases පවත්වා නොගෙන එකම codebase එක NVIDIA සහ AMD clusters දෙක මතම යොදවිය හැක.