Belajar LLVM - GPU & Heterogeneous Computing

#Pendahuluan

Setelah di episode 22 kita memahami MLIR untuk domain-specific compilation, pada episode ini kita membahas target GPU di LLVM — bagaimana LLVM IR diturunkan ke code untuk GPU vendor berbeda.

Heterogeneous computing (CPU + GPU + TPU + custom accelerators) adalah tren dominan 2025-2026. LLVM mendukung berbagai target GPU melalui backend dedicated, memungkinkan satu framework kompilasi untuk seluruh spektrum hardware.

#Target GPU

#NVPTX (NVIDIA)

Backend untuk NVIDIA GPU — menghasilkan PTX (Parallel Thread Execution) yang dijalankan oleh CUDA driver:

Compile ke PTX NVIDIA

clang -O2 --target=nvptx64-nvidia-cuda -S kernel.cu -o kernel.ptx

PTX adalah assembly-level ISA untuk NVIDIA GPU. Dari PTX, CUDA driver mengonversi ke SASS (shader assembly) yang dieksekusi langsung oleh hardware.

#AMDGPU (ROCm)

Backend untuk AMD GPU — menghasilkan AMDGPU ISA yang dijalankan oleh ROCm runtime:

Compile ke AMDGPU

clang -O2 --target=amdgcn-amd-amdhsa -march=gfx900 kernel.cl -o kernel.bc

AMDGPU mendukung COV6 (Compiler Volume 6) sebagai default ISA di ROCm 6.3+, dengan perbaikan register pressure dan scheduling.

#SPIR-V

SPIR-V adalah intermediate format portable untuk GPU — didukung oleh Vulkan, OpenCL, dan OneAPI:

Compile ke SPIR-V

clang -O2 --target=spirv64-vulkan-library kernel.cl -o kernel.spv

SPIR-V mendukung BFloat16 sejak 2026 melalui kontribusi AMD — format 16-bit yang lebih cocok untuk ML workload dibanding FP16 standar.

#DirectX (DXIL)

Backend DirectX memungkinkan kompilasi shader ke format yang bisa dijalankan di Windows GPU.

#SYCL & OneAPI

SYCL adalah standar C++ untuk heterogeneous computing. libsycl adalah runtime SYCL yang di-upstream ke LLVM 22 (2026) oleh Intel:

Contoh SYCL kernel

#include <sycl/sycl.hpp>
 
int main() {
    sycl::queue q{sycl::gpu_selector_v};
    q.parallel_for(sycl::range<1>(1024), [](sycl::id<1> i) {
        // Kernel GPU berjalan di sini
        a[i] = b[i] + c[i];
    }).wait();
}

libsycl di LLVM 22 memungkinkan satu codebase SYCL dijalankan di NVIDIA, AMD, dan Intel GPU tanpa perubahan — model oneAPI.

#Praktik: Compile Kernel ke PTX dan SPIR-V

Compile kernel ke berbagai target GPU

# Kernel OpenCL sederhana
cat > kernel.cl << 'EOF'
__kernel void vecadd(__global const float *a,
                     __global const float *b,
                     __global float *c) {
    int i = get_global_id(0);
    c[i] = a[i] + b[i];
}
EOF
 
# Compile ke PTX (NVIDIA)
clang -O2 --target=nvptx64-nvidia-cuda --cuda-path=/usr/local/cuda \
    kernel.cl -o kernel.ptx
 
# Compile ke AMDGPU
clang -O2 --target=amdgcn-amd-amdhsa -march=gfx900 \
    kernel.cl -o kernel-amd.bc
 
# Compile ke SPIR-V
clang -O2 --target=spirv64-vulkan-library kernel.cl -o kernel.spv
 
# Inspeksi hasil
llvm-dis kernel-amd.bc -o - | head -50

Note

Untuk NVIDIA GPU, kalian juga bisa menggunakan clang --cuda-gpu-arch=sm_80 untuk menghasilkan PTX + SASS dalam satu compilation. Ini adalah workflow standar untuk CUDA development di LLVM — Clang mendukung CUDA compilation sejak LLVM 3.8.

#Penutup

Inti yang harus dibawa pulang:

NVPTX untuk NVIDIA GPU (PTX output); AMDGPU untuk AMD GPU (ROCm); SPIR-V untuk Vulkan/OpenCL.
AMDGPU COV6 (ROCm 6.3+) adalah default ISA baru dengan perbaikan scheduling.
SYCL (libsycl di LLVM 22) memungkinkan satu codebase berjalan di GPU vendor berbeda.
BFloat16 di SPIR-V membuka dukungan ML workload di GPU portable.

Di episode 24 selanjutnya kita akan membahas cross-compilation & embedded toolchain — triple target (aarch64, riscv32, arm-none-eabi), sysroot, multilib, bare-metal ARM/RISC-V, dan praktik build firmware blink LED untuk QEMU RISC-V. Sampai jumpa di episode 24!