GPU 컴퓨팅 (GPGPU)

/* 오큐펀시 최적화 — 최적 블록 크기 자동 결정 */
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize, &blockSize,
    myKernel,         /* 대상 커널 */
    0,                /* 동적 공유메모리 크기 */
    0                 /* 블록 크기 제한 (0 = 제한 없음) */
);
/* blockSize = SM 리소스를 최대 활용하는 블록 크기 */

int gridSize = (N + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(data, N);

/* 커널 레지스터 제한 — 스필과 오큐펀시 트레이드오프 */
__global__ void
__launch_bounds__(256, 8)  /* 최대 256 스레드/블록, 최소 8 블록/SM */
myOptimizedKernel(float *data) {
    /* 컴파일러가 레지스터를 256*8 블록에 맞게 할당 */
    /* → 레지스터/스레드 = 65536 / (256*8) = 32 */
}

NVIDIA 커널 모듈 상세

NVIDIA 독점 드라이버는 5개의 커널 모듈로 구성됩니다. 각 모듈은 역할이 명확히 분리되어 있으며, sysfs와 procfs를 통해 런타임 상태를 조회하거나 파라미터를 변경할 수 있습니다.

# 주요 모듈 파라미터 확인
cat /proc/driver/nvidia/params
# NVreg_EnablePCIeGen3=1
# NVreg_MemoryPoolSize=256  (MB, nvidia-uvm 내부 풀)
# NVreg_PreserveVideoMemoryAllocations=0

# GPU 정보 조회
cat /proc/driver/nvidia/gpus/0000:01:00.0/information
# Model:           NVIDIA H100 80GB HBM3
# IRQ:             153
# GPU UUID:        GPU-xxxx-xxxx-xxxx-xxxx

# sysfs 전원 관리
cat /sys/bus/pci/devices/0000:01:00.0/power_state
# D0 (활성) / D3hot (절전)

# modprobe 옵션으로 파라미터 설정 (/etc/modprobe.d/nvidia.conf)
options nvidia NVreg_PreserveVideoMemoryAllocations=1
options nvidia NVreg_MemoryPoolSize=512

# DKMS 빌드 상태 확인
dkms status nvidia
# nvidia/550.127.05, 6.1.0-26-amd64, x86_64: installed

# nvidia-persistenced — GPU 컨텍스트 지속 (초기화 지연 제거)
sudo systemctl enable nvidia-persistenced
sudo systemctl start nvidia-persistenced
# GPU 초기화가 첫 CUDA 호출 시 ~0.5s 걸리는 문제 해소
# HPC/ML 서버에서 필수 — 잦은 CUDA 프로세스 시작/종료 시

# Fabricmanager — NVSwitch 기반 멀티 GPU 시스템 (DGX)
sudo systemctl enable nvidia-fabricmanager
# NVSwitch 토폴로지 관리, GPU 간 NVLink 풀 메시 구성

# /proc/driver/nvidia/ 전체 구조 확인
ls /proc/driver/nvidia/
# gpus/  params  patches  registry  version
cat /proc/driver/nvidia/version
# NVRM version: NVIDIA UNIX x86_64 Kernel Module  550.127.05

GPUDirect / RDMA

GPUDirect 기술은 GPU 메모리와 외부 디바이스(다른 GPU, NIC, NVMe) 간의 직접 DMA 경로를 제공하여 CPU 메모리 복사를 제거합니다. HPC와 대규모 AI 학습에서 노드 간 통신 병목을 해소하는 핵심 기술입니다.

# GPUDirect P2P 토폴로지 확인
nvidia-smi topo -m
#         GPU0  GPU1  GPU2  GPU3  NIC0  CPU
# GPU0      X   NV12  NV12  NV12  SYS   SYS
# GPU1    NV12    X   NV12  NV12  SYS   SYS
# NV12 = NVLink 12 hops, SYS = PCIe through CPU

# GPU 간 P2P 접근 가능 여부 확인
nvidia-smi topo -p2p r
# GPU0 GPU1: OK (NVLink P2P 가능)

# nvidia-peermem 로드 확인
lsmod | grep nvidia_peermem
# nvidia_peermem  16384  0
cat /sys/kernel/mm/memory_peers/nvidia-peermem/version
# 1.3

# NCCL 환경 변수로 GPUDirect 제어
export NCCL_NET_GDR_LEVEL=5       # GPUDirect RDMA 활성화 수준
export NCCL_IB_HCA=mlx5_0:1       # InfiniBand HCA 지정
export NCCL_P2P_LEVEL=NVL          # NVLink P2P 사용
export NCCL_DEBUG=INFO             # 디버그 로그

NCCL 집합 통신

NCCL(NVIDIA Collective Communications Library)은 다중 GPU 간 집합 통신을 최적화하는 라이브러리로, 분산 딥러닝 학습의 핵심 인프라입니다. NVLink, PCIe, InfiniBand 등 사용 가능한 모든 인터커넥트를 자동으로 감지하여 최적의 통신 토폴로지를 구성합니다.

/* NCCL AllReduce — 4-GPU 그래디언트 동기화 */
#include <nccl.h>

ncclComm_t comms[4];
ncclCommInitAll(comms, 4, devs);  /* 4 GPU 커뮤니케이터 생성 */

/* 각 GPU에서 비동기 AllReduce 실행 */
ncclGroupStart();
for (int i = 0; i < 4; i++) {
    cudaSetDevice(i);
    ncclAllReduce(
        sendbuff[i], recvbuff[i],
        count, ncclFloat, ncclSum,
        comms[i], streams[i]
    );
}
ncclGroupEnd();

/* 모든 스트림 동기화 */
for (int i = 0; i < 4; i++) {
    cudaSetDevice(i);
    cudaStreamSynchronize(streams[i]);
}

/* 정리 */
for (int i = 0; i < 4; i++)
    ncclCommDestroy(comms[i]);

NVIDIA 오픈 GPU 커널 모듈

2022년 5월, NVIDIA는 open-gpu-kernel-modules 저장소를 통해 GPU 커널 드라이버의 소스 코드를 공개했습니다(GPL-2.0 / MIT 듀얼 라이선스). Turing(GTX 16xx) 이후 아키텍처에서 사용할 수 있으며, 장기적으로 독점 모듈을 대체할 계획입니다.

NVIDIA 컨테이너 / 가상화(Virtualization)

클라우드 및 AI 인프라에서 GPU를 컨테이너·가상 머신에 안전하게 공유하려면 별도 런타임이 필요합니다. NVIDIA는 Container Toolkit, MIG, vGPU, K8s Device Plugin을 통해 GPU 격리와 스케줄링을 제공합니다.

# Docker에서 CUDA 컨테이너 실행
docker run --gpus all nvidia/cuda:12.4.0-runtime-ubuntu22.04 nvidia-smi

# MIG 인스턴스 생성 (A100/H100)
sudo nvidia-smi mig -cgi 19,19,19 -C  # 3x 3g.40gb 프로파일
nvidia-smi mig -lgi  # GPU 인스턴스 목록
nvidia-smi mig -lci  # 컴퓨트 인스턴스 목록

# Kubernetes GPU 리소스 요청 (Pod spec)
# resources:
#   limits:
#     nvidia.com/gpu: 1

# nvidia-container-toolkit 설치 (Ubuntu)
distribution=$(. /etc/os-release; echo $ID$VERSION_ID)
curl -fsSL https://nvidia.github.io/libnvidia-container/gpgkey | \
  sudo gpg --dearmor -o /usr/share/keyrings/nvidia-container-toolkit-keyring.gpg
sudo apt-get update && sudo apt-get install nvidia-container-toolkit

# Docker 런타임 설정
sudo nvidia-ctk runtime configure --runtime=docker
sudo systemctl restart docker

# MIG + 컨테이너 — 특정 MIG 인스턴스에서 컨테이너 실행
# MIG 디바이스 UUID 확인
nvidia-smi -L
# GPU 0: A100-SXM4-80GB (UUID: GPU-xxxxx)
#   MIG 1g.10gb Device 0: (UUID: MIG-yyyyy)

# MIG UUID로 컨테이너 실행
docker run --gpus '"device=MIG-yyyyy"' nvidia/cuda:12.4.0-base-ubuntu22.04 nvidia-smi

# Kubernetes MIG 리소스 — Device Plugin 설정
# helm install --set migStrategy=single nvidia-device-plugin ...
# Pod spec: nvidia.com/mig-1g.10gb: 1

CUDA 컴파일 파이프라인

CUDA 소스(.cu)는 nvcc 컴파일러 드라이버를 통해 여러 단계를 거칩니다. 호스트 코드는 gcc/clang에 위임되고, 디바이스 코드는 PTX(가상 ISA) → SASS(실제 기계어(Machine Code))로 변환됩니다.

# 기본 컴파일 (sm_80 = Ampere, sm_90 = Hopper)
nvcc -arch=sm_80 -o matmul matmul.cu

# Fatbinary: 여러 아키텍처 동시 타겟
nvcc -gencode arch=compute_80,code=sm_80 \
     -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     -o matmul matmul.cu
# compute_90,code=compute_90 → PTX 포함 (미래 GPU JIT 호환)

# PTX 어셈블리 확인
nvcc -arch=sm_80 --ptx -o matmul.ptx matmul.cu
# SASS 디스어셈블
cuobjdump -sass matmul

# NVRTC — 런타임 CUDA 컴파일 (Driver API)
# PTX를 프로그램 실행 중에 동적 생성하는 시나리오
# AI 프레임워크(PyTorch, JAX)의 커널 퓨전에 사용

# Compute Capability별 아키텍처 코드 매핑
# sm_70 = Volta (V100)
# sm_75 = Turing (T4, RTX 20xx)
# sm_80 = Ampere (A100, A10)
# sm_86 = Ampere (RTX 30xx, A40)
# sm_89 = Ada Lovelace (RTX 40xx, L4, L40)
# sm_90 = Hopper (H100, H200)
# sm_90a = Hopper (SM 전용 기능, GH200)
# sm_100 = Blackwell (B200, GB200)

# 현재 GPU의 Compute Capability 확인
nvidia-smi --query-gpu=compute_cap --format=csv,noheader
# 9.0

# NVCC 상세 컴파일 과정 보기
nvcc -v -arch=sm_80 -o test test.cu 2>&1 | head -50
# cudafe++ → cicc → ptxas → fatbinary → gcc 순서 확인 가능

CUDA 디버깅 및 프로파일링

CUDA 프로그램의 성능 분석과 디버깅에는 NVIDIA가 제공하는 전용 도구체인과 Linux 커널의 /proc/driver/nvidia/ 인터페이스를 활용합니다.

# GPU 상태 모니터링
nvidia-smi
# +-------------------------+------+------+
# | GPU  Name        Temp   | Util | MIG  |
# | Fan  Perf  Pwr:Usage/Cap|  GPU | Mode |
# |=========================+======+======|
# |   0  NVIDIA H100    38C |  85% |  On  |
# |  N/A  P0    310W / 350W |      |      |

# 시스템 프로파일링 (Nsight Systems)
nsys profile --stats=true ./my_cuda_app
# → report.nsys-rep 생성 (GUI에서 타임라인 분석)

# 커널 단위 프로파일링 (Nsight Compute)
ncu --set full --target-processes all ./my_cuda_app
# → 오큐펀시, SM 활용률, 메모리 throughput 상세 리포트

# 메모리 오류 탐지
compute-sanitizer --tool memcheck ./my_cuda_app
# → 범위 초과 접근, 초기화되지 않은 읽기 탐지

# 레이스 컨디션 탐지
compute-sanitizer --tool racecheck ./my_cuda_app
# → 공유 메모리 WAR/WAW/RAW 레이스 탐지

# 동기화 오류 탐지
compute-sanitizer --tool synccheck ./my_cuda_app
# → __syncthreads() 누락, 불균형 배리어 탐지

# cuda-gdb 디버깅 세션
cuda-gdb ./my_cuda_app
# (cuda-gdb) set cuda break_on_launch all
# (cuda-gdb) run
# (cuda-gdb) cuda thread (0,0,0)   ← 특정 스레드 선택
# (cuda-gdb) info cuda threads     ← 워프/블록 상태
# (cuda-gdb) print threadIdx       ← 현재 스레드 인덱스
# (cuda-gdb) cuda kernel            ← 활성 커널 정보

멀티 GPU 프로그래밍

단일 노드에 여러 GPU가 장착된 환경에서 CUDA는 cudaSetDevice()로 활성 GPU를 전환하고, P2P(Peer-to-Peer) 접근과 비동기 전송으로 GPU 간 데이터를 교환합니다. 대규모 AI 학습에서는 NCCL과 결합하여 데이터/모델/파이프라인 병렬처리를 구현합니다.

/* 멀티 GPU — P2P 메모리 접근 + 비동기 전송 */
int deviceCount;
cudaGetDeviceCount(&deviceCount);

/* P2P 접근 활성화 (NVLink/PCIe P2P) */
for (int i = 0; i < deviceCount; i++) {
    cudaSetDevice(i);
    for (int j = 0; j < deviceCount; j++) {
        if (i != j) {
            int canAccess;
            cudaDeviceCanAccessPeer(&canAccess, i, j);
            if (canAccess)
                cudaDeviceEnablePeerAccess(j, 0);
        }
    }
}

/* GPU 0의 메모리를 GPU 1에서 직접 접근 (UVA) */
cudaSetDevice(0);
float *d_gpu0;
cudaMalloc(&d_gpu0, size);

cudaSetDevice(1);
/* P2P 활성화 시 GPU1 커널에서 d_gpu0 직접 읽기 가능 */
readFromGpu0Kernel<<<grid, block>>>(d_gpu0);

/* 비동기 GPU간 복사 (Copy Engine 사용) */
cudaMemcpyPeerAsync(d_gpu1, 1,   /* dst GPU 1 */
                     d_gpu0, 0,   /* src GPU 0 */
                     size, stream);

CUDA 에러 처리 패턴

CUDA API 호출은 cudaError_t를 반환하며, 커널 실행 오류는 비동기적으로 발생합니다. 프로덕션 코드에서는 모든 CUDA 호출을 검사하는 매크로(Macro)를 사용하고, 커널 실행 후 cudaGetLastError()와 cudaDeviceSynchronize()로 에러를 포착합니다.

/* CUDA 에러 검사 매크로 — 프로덕션 필수 패턴 */
#define CUDA_CHECK(call) do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

/* 사용 예 */
CUDA_CHECK(cudaMalloc(&d_ptr, size));
CUDA_CHECK(cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice));

/* 커널 실행 에러 검사 — 2단계 필수 */
myKernel<<<grid, block>>>(d_ptr);
CUDA_CHECK(cudaGetLastError());          /* 실행 설정 오류 (즉시) */
CUDA_CHECK(cudaDeviceSynchronize());     /* 실행 중 오류 (비동기) */

/* 스트림 콜백으로 비동기 에러 처리 */
cudaLaunchHostFunc(stream, [](void* data) {
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess)
        handle_gpu_error(err);
}, nullptr);

CUDA 라이브러리 에코시스템

CUDA 생태계의 진정한 강점은 수십 년간 최적화된 도메인별 라이브러리에 있습니다. 대부분의 AI/HPC 워크로드는 커널을 직접 작성하지 않고 이 라이브러리들을 조합합니다.

/* cuBLAS GEMM — 행렬 곱셈 C = α·A·B + β·C */
cublasHandle_t handle;
cublasCreate(&handle);

float alpha = 1.0f, beta = 0.0f;
cublasSgemm(handle,
    CUBLAS_OP_N, CUBLAS_OP_N,
    M, N, K,
    &alpha,
    d_A, M,     /* A: M×K */
    d_B, K,     /* B: K×N */
    &beta,
    d_C, M);    /* C: M×N */

cublasDestroy(handle);
/* Tensor Core 자동 활용 (FP16/TF32/FP8 입력 시) */

/* cuDNN — 합성곱(Convolution) 연산 예제 */
cudnnHandle_t cudnn;
cudnnCreate(&cudnn);

cudnnTensorDescriptor_t inputDesc, outputDesc;
cudnnFilterDescriptor_t filterDesc;
cudnnConvolutionDescriptor_t convDesc;

/* 텐서/필터 디스크립터 설정 (NCHW 포맷) */
cudnnCreateTensorDescriptor(&inputDesc);
cudnnSetTensor4dDescriptor(inputDesc, CUDNN_TENSOR_NCHW,
    CUDNN_DATA_FLOAT, batch, channels, height, width);

/* 최적 합성곱 알고리즘 자동 선택 */
cudnnConvolutionFwdAlgo_t algo;
cudnnGetConvolutionForwardAlgorithm_v7(cudnn,
    inputDesc, filterDesc, convDesc, outputDesc,
    1, &returnedCount, &perfResults);
algo = perfResults.algo;  /* 가장 빠른 알고리즘 선택 */

/* 합성곱 실행 — Tensor Core 자동 활용 */
cudnnSetConvolutionMathType(convDesc, CUDNN_TENSOR_OP_MATH);
float alpha = 1.0f, beta = 0.0f;
cudnnConvolutionForward(cudnn, &alpha,
    inputDesc, d_input, filterDesc, d_filter,
    convDesc, algo, d_workspace, workspaceSize,
    &beta, outputDesc, d_output);

cudnnDestroy(cudnn);

딥러닝 프레임워크 통합

PyTorch, TensorFlow, JAX 등 주요 딥러닝 프레임워크는 CUDA를 통해 GPU 가속을 구현합니다. 프레임워크는 내부적으로 cuBLAS, cuDNN, NCCL, cuFFT 등 CUDA 라이브러리를 호출하며, 사용자는 Python API만으로 Tensor Core, 멀티 GPU, 혼합 정밀도 학습을 활용할 수 있습니다.

# PyTorch CUDA 사용 확인
python3 -c "import torch; print(torch.cuda.is_available())"
# True
python3 -c "import torch; print(torch.cuda.get_device_name(0))"
# NVIDIA H100 80GB HBM3

# PyTorch 혼합 정밀도 학습 (AMP)
# scaler = torch.amp.GradScaler()
# with torch.amp.autocast(device_type='cuda', dtype=torch.bfloat16):
#     output = model(input)
#     loss = criterion(output, target)
# scaler.scale(loss).backward()
# scaler.step(optimizer)

# TensorRT 모델 최적화 (FP16 추론)
trtexec --onnx=model.onnx \
  --fp16 \
  --workspace=4096 \
  --saveEngine=model_fp16.trt \
  --verbose
# → 레이어 퓨전, 텐서 레이아웃 최적화, Tensor Core 활용

# torch.compile — Triton 커널 자동 생성 (PyTorch 2.0+)
# model = torch.compile(model, mode='max-autotune')
# → 커널 퓨전, 메모리 접근 패턴 최적화, Triton→PTX 컴파일

OpenCL — 크로스 플랫폼 GPU 컴퓨트

OpenCL(Open Computing Language)은 Khronos Group이 표준화한 범용 병렬 컴퓨팅 API입니다. NVIDIA·AMD·Intel GPU, ARM Mali, Qualcomm Adreno 등 다양한 가속기를 동일한 코드로 활용할 수 있습니다. Linux에서는 Mesa Rusticl(신규 Rust 구현)과 Clover(레거시 C++ 구현)가 오픈소스 드라이버를 제공합니다.

실행 모델: Work-item / Work-group / NDRange

/* OpenCL 메모리 계층 */
/*
 * Global Memory  (~수 GB VRAM): 모든 work-item 공유, 높은 지연
 * Local Memory   (~16~64 KB):   한 work-group 내 공유, 빠름
 * Private Memory (레지스터):    각 work-item 전용, 최고속
 * Constant Memory (캐시):       읽기 전용, GPU가 캐싱 최적화
 */

/* GEMM OpenCL 커널 예제 (타일링 최적화) */
__kernel void gemm_tiled(
    __global const float *A,     /* Global memory 입력 */
    __global const float *B,
    __global       float *C,
    const int N)
{
    __local float tileA[16][16];   /* Local memory 타일 */
    __local float tileB[16][16];

    int row = get_global_id(0);    /* 전체 좌표 */
    int col = get_global_id(1);
    int lrow = get_local_id(0);   /* work-group 내 좌표 */
    int lcol = get_local_id(1);

    float sum = 0.0f;

    for (int t = 0; t < N / 16; t++) {
        /* 타일을 Local memory로 협력 로딩 */
        tileA[lrow][lcol] = A[row * N + (t * 16 + lcol)];
        tileB[lrow][lcol] = B[(t * 16 + lrow) * N + col];
        barrier(CLK_LOCAL_MEM_FENCE);  /* work-group 동기화 */

        for (int k = 0; k < 16; k++)
            sum += tileA[lrow][k] * tileB[k][lcol];
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    C[row * N + col] = sum;
}

Mesa Rusticl vs Clover

# OpenCL 컴파일 파이프라인
# OpenCL C → LLVM IR → SPIR-V → GPU ISA

# clinfo로 OpenCL 드라이버 확인
clinfo | head -40
# Platform: Mesa/Rusticl or AMD ROCm or Intel OpenCL

# Rusticl 강제 활성화 (환경 변수)
RUSTICL_ENABLE=radeonsi clinfo

# 오프라인 OpenCL 커널 컴파일 (SPIR-V 생성)
clang -cl-std=CL3.0 -target spir64 -emit-llvm -c gemm.cl -o gemm.bc
llvm-spirv gemm.bc -o gemm.spv

OpenCL 호스트 워크플로우

OpenCL 호스트 프로그램은 플랫폼 탐색부터 결과 수집까지 정해진 API 순서를 따릅니다. 아래 다이어그램은 완전한 호스트 측 워크플로우를 보여 줍니다.

다음은 위 워크플로우를 완전하게 구현한 C 예제입니다. 오류 처리와 리소스 해제까지 포함한 실제 사용 가능한 코드입니다.

#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define CHECK_CL(err, msg) \
    do { if ((err) != CL_SUCCESS) { \
        fprintf(stderr, "%s: %d\n", (msg), (err)); exit(1); } } while(0)

static const char *kernel_src =
    "__kernel void vadd(__global const float *a,\n"
    "                   __global const float *b,\n"
    "                   __global       float *c) {\n"
    "    int i = get_global_id(0);\n"
    "    c[i] = a[i] + b[i];\n"
    "}\n";

int main(void) {
    const int N = 1024;
    const size_t sz = N * sizeof(float);
    cl_int err;

    /* ① 플랫폼 탐색 */
    cl_platform_id platform;
    err = clGetPlatformIDs(1, &platform, NULL);
    CHECK_CL(err, "clGetPlatformIDs");

    /* ② 디바이스 선택 (GPU 우선) */
    cl_device_id device;
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    if (err != CL_SUCCESS)
        err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &device, NULL);
    CHECK_CL(err, "clGetDeviceIDs");

    /* ③ 컨텍스트 생성 */
    cl_context ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    CHECK_CL(err, "clCreateContext");

    /* ④ 커맨드 큐 생성 (OpenCL 2.0+ 스타일) */
    cl_command_queue queue = clCreateCommandQueueWithProperties(
        ctx, device, NULL, &err);
    CHECK_CL(err, "clCreateCommandQueueWithProperties");

    /* ⑤ 프로그램 생성 & 빌드 */
    cl_program prog = clCreateProgramWithSource(
        ctx, 1, &kernel_src, NULL, &err);
    CHECK_CL(err, "clCreateProgramWithSource");
    err = clBuildProgram(prog, 1, &device,
        "-cl-std=CL3.0 -cl-fast-relaxed-math", NULL, NULL);
    if (err != CL_SUCCESS) {
        char log[4096];
        clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG,
            sizeof(log), log, NULL);
        fprintf(stderr, "Build error:\n%s\n", log);
        exit(1);
    }

    /* ⑥ 커널 생성 */
    cl_kernel kernel = clCreateKernel(prog, "vadd", &err);
    CHECK_CL(err, "clCreateKernel");

    /* ⑦ 버퍼 할당 */
    cl_mem buf_a = clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sz, NULL, &err);
    cl_mem buf_b = clCreateBuffer(ctx, CL_MEM_READ_ONLY,  sz, NULL, &err);
    cl_mem buf_c = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, sz, NULL, &err);

    /* 호스트 데이터 초기화 */
    float *ha = malloc(sz), *hb = malloc(sz), *hc = malloc(sz);
    for (int i = 0; i < N; i++) { ha[i] = (float)i; hb[i] = (float)(N - i); }

    /* 버퍼 쓰기 (호스트 → 디바이스) */
    err = clEnqueueWriteBuffer(queue, buf_a, CL_TRUE, 0, sz, ha, 0, NULL, NULL);
    err = clEnqueueWriteBuffer(queue, buf_b, CL_TRUE, 0, sz, hb, 0, NULL, NULL);

    /* ⑧ 커널 인수 설정 */
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_a);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_b);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_c);

    /* ⑨ 커널 실행 */
    size_t global = N, local = 64;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
        &global, &local, 0, NULL, NULL);
    CHECK_CL(err, "clEnqueueNDRangeKernel");

    /* ⑩ 결과 읽기 (디바이스 → 호스트) */
    err = clEnqueueReadBuffer(queue, buf_c, CL_TRUE, 0, sz, hc, 0, NULL, NULL);
    clFinish(queue);
    printf("c[0]=%.1f c[N-1]=%.1f\n", hc[0], hc[N - 1]);

    /* ⑪ 리소스 해제 */
    free(ha); free(hb); free(hc);
    clReleaseMemObject(buf_a); clReleaseMemObject(buf_b); clReleaseMemObject(buf_c);
    clReleaseKernel(kernel);
    clReleaseProgram(prog);
    clReleaseCommandQueue(queue);
    clReleaseContext(ctx);
    return 0;
}

OpenCL 메모리 모델

OpenCL은 4단계 메모리 계층을 명시적으로 관리합니다. 글로벌 메모리(Global Memory)와 로컬 메모리(Local Memory) 간의 데이터 이동을 수동으로 제어하여 메모리 대역폭을 최적화합니다.

다음 예제는 명시적 버퍼 관리와 서브 버퍼(Sub-buffer) 활용법을 보여 줍니다.

/* 명시적 버퍼 관리 및 서브 버퍼 예제 */

/* 큰 버퍼 한 번 할당 후 분할하여 재사용 */
const size_t TOTAL = 64 * 1024 * 1024;  /* 64 MB */
cl_mem big_buf = clCreateBuffer(ctx,
    CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
    TOTAL, NULL, &err);

/* 서브 버퍼: 오프셋+크기로 큰 버퍼의 일부를 참조 */
cl_buffer_region region = {
    .origin = 0,
    .size   = TOTAL / 2
};
cl_mem sub_a = clCreateSubBuffer(big_buf,
    CL_MEM_READ_ONLY, CL_BUFFER_CREATE_TYPE_REGION,
    &region, &err);

region.origin = TOTAL / 2;
cl_mem sub_b = clCreateSubBuffer(big_buf,
    CL_MEM_READ_ONLY, CL_BUFFER_CREATE_TYPE_REGION,
    &region, &err);

/* 비동기 전송: 이벤트로 완료 대기 */
cl_event write_ev;
clEnqueueWriteBuffer(queue, sub_a, CL_FALSE, 0,
    TOTAL / 2, host_data, 0, NULL, &write_ev);

/* 커널 실행을 write_ev 완료 후로 의존성 설정 */
clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
    &global, &local, 1, &write_ev, NULL);

/* 맵(Map)을 통한 제로 카피: PCIe 전송 없이 호스트에서 직접 접근 */
void *ptr = clEnqueueMapBuffer(queue, big_buf,
    CL_TRUE, CL_MAP_READ,
    0, TOTAL, 0, NULL, NULL, &err);
memcpy(host_result, ptr, TOTAL);
clEnqueueUnmapMemObject(queue, big_buf, ptr, 0, NULL, NULL);

clReleaseEvent(write_ev);
clReleaseMemObject(sub_a);
clReleaseMemObject(sub_b);
clReleaseMemObject(big_buf);

OpenCL 3.0 주요 변경사항

OpenCL 3.0(2020년 공개)은 기존 2.x 기능을 모두 선택적(Optional)으로 변경하고, OpenCL 1.2 기능만 필수(Mandatory)로 지정하여 임베디드·모바일 구현의 진입 장벽을 낮추었습니다. 런타임에서 기능 쿼리를 통해 지원 여부를 확인한 뒤 코드 경로를 선택하는 방식이 필요합니다.

런타임에서 OpenCL 3.0 기능 지원 여부를 쿼리하는 방법은 다음과 같습니다.

/* OpenCL 3.0 기능 쿼리 예제 */
#include <CL/cl.h>
#include <stdio.h>

void query_opencl30_features(cl_device_id device) {
    /* OpenCL C 버전 확인 */
    char version[128];
    clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION,
        sizeof(version), version, NULL);
    printf("OpenCL C Version: %s\n", version);

    /* SVM 지원 여부 */
    cl_device_svm_capabilities svm_caps;
    clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES,
        sizeof(svm_caps), &svm_caps, NULL);
    if (svm_caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER)
        printf("  SVM: Coarse-Grain Buffer 지원\n");
    if (svm_caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER)
        printf("  SVM: Fine-Grain Buffer 지원\n");
    if (svm_caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM)
        printf("  SVM: Fine-Grain System 지원\n");

    /* 원자 연산 메모리 모델 */
    cl_device_atomic_capabilities atomic_caps;
    clGetDeviceInfo(device, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
        sizeof(atomic_caps), &atomic_caps, NULL);
    printf("Atomic caps: 0x%lx\n", (unsigned long)atomic_caps);

    /* Work-group collective 함수 (reduce, scan) */
    cl_bool wg_collective;
    clGetDeviceInfo(device,
        CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT,
        sizeof(wg_collective), &wg_collective, NULL);
    printf("Work-group collective: %s\n",
        wg_collective ? "지원" : "미지원");

    /* 디바이스 측 enqueue */
    cl_device_device_enqueue_capabilities enqueue_caps;
    clGetDeviceInfo(device,
        CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES,
        sizeof(enqueue_caps), &enqueue_caps, NULL);
    printf("Device-side enqueue: %s\n",
        enqueue_caps ? "지원" : "미지원");

    /* OpenCL C 3.0 세부 기능 목록 */
    size_t feat_size;
    clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_FEATURES,
        0, NULL, &feat_size);
    size_t n = feat_size / sizeof(cl_name_version);
    cl_name_version *feats = malloc(feat_size);
    clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_FEATURES,
        feat_size, feats, NULL);
    printf("OpenCL C 3.0 기능 (%zu개):\n", n);
    for (size_t i = 0; i < n; i++)
        printf("  %s (v%u.%u)\n", feats[i].name,
            CL_VERSION_MAJOR(feats[i].version),
            CL_VERSION_MINOR(feats[i].version));
    free(feats);
}

# clinfo로 OpenCL 3.0 기능 일괄 확인
clinfo --human | grep -E "Version|SVM|Atomic|Enqueue|collective"

# Mesa Rusticl OpenCL 3.0 지원 확인
RUSTICL_ENABLE=radeonsi clinfo | grep "Device Version"

# OpenCL C 3.0으로 컴파일 (선택적 기능 활성화)
clang -cl-std=CL3.0 \
      -cl-ext=+__opencl_c_int64,+__opencl_c_atomic_order_acq_rel \
      -target spir64 -emit-llvm -c kernel.cl -o kernel.bc
llvm-spirv kernel.bc -o kernel.spv

Vulkan Compute 파이프라인

Vulkan은 Khronos Group이 표준화한 저수준 크로스 플랫폼 그래픽/컴퓨트 API입니다. Compute Pipeline을 통해 Graphics Pipeline 없이 GPGPU 연산을 수행할 수 있으며, 래스터라이저·프래그먼트 셰이더 대신 Compute Shader만으로 구성됩니다. Vulkan 컴퓨트는 CUDA와 달리 벤더 중립적이어서 NVIDIA, AMD, Intel, Qualcomm, ARM Mali 등 모든 Vulkan 호환 GPU에서 동일한 SPIR-V 셰이더를 실행할 수 있습니다.

Linux에서 Vulkan 드라이버는 크게 두 계열로 나뉩니다: Mesa 오픈소스 드라이버(RadV, ANV/Hasvk, PanVK, Turnip 등)와 벤더 독점 드라이버(NVIDIA, AMDGPU-PRO). 모든 Vulkan 드라이버는 DRM 서브시스템의 /dev/dri/renderD* 노드를 통해 GPU에 접근하므로, root 권한 없이도 컴퓨트 작업을 수행할 수 있습니다(render node 그룹 소속 필요).

Vulkan Linux 드라이버 스택

Vulkan 애플리케이션이 GPU에 접근하는 전체 소프트웨어 스택은 다음과 같습니다. Vulkan 로더(libvulkan.so)가 ICD(Installable Client Driver)를 검색하여 적절한 드라이버를 동적 로드하고, 드라이버는 DRM 렌더 노드 ioctl로 GPU와 통신합니다.

# Vulkan 드라이버 정보 확인
vulkaninfo --summary
# GPU0: AMD Radeon RX 7900 XTX (RADV NAVI31)
#   apiVersion    = 1.3.274
#   driverVersion = 24.0.99
#   driverID      = DRIVER_ID_MESA_RADV

# ICD 파일 확인 (로더가 검색하는 JSON 매니페스트)
ls /usr/share/vulkan/icd.d/
# radeon_icd.x86_64.json  intel_icd.x86_64.json  nvidia_icd.json

# 특정 드라이버 강제 선택 (다중 GPU 시스템)
VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/radeon_icd.x86_64.json ./my_app

# Vulkan 레이어 확인
vulkaninfo --show-layer-list 2>&1 | head -20

# Mesa 디버그: NIR→ISA 컴파일 덤프
RADV_DEBUG=preoptir,shaders ./my_vulkan_app 2>shader_dump.txt

Compute Shader와 파이프라인 구성

Vulkan Compute Pipeline의 구성 과정은 크게 6단계로 나뉩니다: ① GLSL/HLSL 작성 → ② SPIR-V 컴파일 → ③ VkDevice/VkQueue 생성 → ④ 리소스 할당(버퍼, 디스크립터) → ⑤ 파이프라인 생성 → ⑥ 커맨드 기록 및 제출.

/* GLSL Compute Shader — 행렬 곱셈 (공유 메모리 타일링) */
/* matmul.comp → glslc matmul.comp -o matmul.spv */

#version 450

#define TILE_SIZE 16
layout(local_size_x = TILE_SIZE, local_size_y = TILE_SIZE) in;

layout(set = 0, binding = 0) readonly buffer MatA { float A[]; };
layout(set = 0, binding = 1) readonly buffer MatB { float B[]; };
layout(set = 0, binding = 2)          buffer MatC { float C[]; };

layout(push_constant) uniform PushConst { uint N; } pc;

/* 공유 메모리 — 워크그룹 내 invocation 간 공유 */
shared float tileA[TILE_SIZE][TILE_SIZE];
shared float tileB[TILE_SIZE][TILE_SIZE];

void main() {
    uint row = gl_GlobalInvocationID.y;
    uint col = gl_GlobalInvocationID.x;
    uint lr  = gl_LocalInvocationID.y;
    uint lc  = gl_LocalInvocationID.x;

    float sum = 0.0;
    uint numTiles = (pc.N + TILE_SIZE - 1) / TILE_SIZE;

    for (uint t = 0; t < numTiles; t++) {
        tileA[lr][lc] = A[row * pc.N + t * TILE_SIZE + lc];
        tileB[lr][lc] = B[(t * TILE_SIZE + lr) * pc.N + col];
        barrier();  /* 워크그룹 동기화 (CUDA __syncthreads 대응) */

        for (uint k = 0; k < TILE_SIZE; k++)
            sum += tileA[lr][k] * tileB[k][lc];
        barrier();
    }
    C[row * pc.N + col] = sum;
}

/* vkCmdDispatch(N/TILE_SIZE, N/TILE_SIZE, 1) 으로 실행
 * → work-group 수: (N/16)² 개
 * → work-group 내 invocation: 16×16 = 256개 */

SPIR-V 셰이더 컴파일

SPIR-V(Standard Portable Intermediate Representation)는 Vulkan의 셰이더 바이트코드 형식입니다. GLSL·HLSL·Slang 등 고급 셰이딩 언어에서 SPIR-V로 컴파일한 뒤 Vulkan 드라이버에 전달하면, 드라이버가 대상 GPU의 네이티브 ISA(GCN/RDNA, Xe, SASS 등)로 최종 번역합니다.

# GLSL → SPIR-V 컴파일
glslc -fshader-stage=compute shader.comp -o shader.spv

# 최적화 옵션
glslc -O shader.comp -o shader.spv                    # 기본 최적화
glslc --target-env=vulkan1.3 shader.comp -o shader.spv # Vulkan 1.3 타겟

# SPIR-V 디스어셈블 (spirv-tools)
spirv-dis shader.spv
# OpCapability Shader
# OpMemoryModel Logical GLSL450
# OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
# OpExecutionMode %main LocalSize 256 1 1

# SPIR-V 유효성 검사
spirv-val shader.spv

# SPIR-V 최적화 (size/performance)
spirv-opt -O shader.spv -o shader_opt.spv

# HLSL → SPIR-V (DirectX Shader Compiler)
dxc -spirv -T cs_6_0 -E CSMain shader.hlsl -Fo shader.spv

# spirv-cross: SPIR-V → GLSL/HLSL/MSL 역변환
spirv-cross --output shader_back.glsl shader.spv

Vulkan 메모리 관리

Vulkan은 명시적 메모리 관리를 요구합니다. OpenGL이 드라이버에 위임하던 메모리 할당/바인딩을 애플리케이션이 직접 제어하며, 이를 통해 메모리 레이아웃과 전송을 최적화할 수 있습니다. GPU 메모리는 VkPhysicalDeviceMemoryProperties로 조회한 메모리 타입과 힙(Heap) 정보를 기반으로 할당합니다.

Vulkan 동기화 모델

Vulkan의 동기화는 완전히 명시적입니다. 드라이버가 자동으로 동기화하는 OpenGL과 달리, Vulkan에서는 모든 리소스 의존성을 개발자가 직접 선언해야 합니다. 잘못된 동기화는 데이터 레이스, 렌더링 아티팩트, 크래시의 주요 원인입니다.

/* Vulkan Pipeline Barrier — Compute 셰이더 결과 읽기 전 동기화 */
VkBufferMemoryBarrier barrier = {
    .sType               = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER,
    .srcAccessMask       = VK_ACCESS_SHADER_WRITE_BIT,   /* Compute 셰이더 쓰기 */
    .dstAccessMask       = VK_ACCESS_HOST_READ_BIT,      /* CPU 읽기 */
    .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
    .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
    .buffer              = outputBuffer,
    .offset              = 0,
    .size                = VK_WHOLE_SIZE,
};

vkCmdPipelineBarrier(cmdBuf,
    VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT,  /* src: Compute 완료 후 */
    VK_PIPELINE_STAGE_HOST_BIT,            /* dst: CPU 접근 전 */
    0,                                     /* 플래그 */
    0, NULL,                               /* 메모리 배리어 */
    1, &barrier,                            /* 버퍼 배리어 */
    0, NULL                                /* 이미지 배리어 */
);

/* Vulkan 1.3 Synchronization2 — 더 직관적인 API */
VkBufferMemoryBarrier2 barrier2 = {
    .sType         = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
    .srcStageMask  = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
    .srcAccessMask = VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT,
    .dstStageMask  = VK_PIPELINE_STAGE_2_HOST_BIT,
    .dstAccessMask = VK_ACCESS_2_HOST_READ_BIT,
    .buffer        = outputBuffer,
    .size          = VK_WHOLE_SIZE,
};
VkDependencyInfo depInfo = {
    .sType                    = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
    .bufferMemoryBarrierCount = 1,
    .pBufferMemoryBarriers    = &barrier2,
};
vkCmdPipelineBarrier2(cmdBuf, &depInfo);

Vulkan Compute 호스트 코드 작성

Vulkan Compute의 전체 호스트 코드(C)는 초기화 → 리소스 할당 → 파이프라인 생성 → 디스패치 → 정리 순서로 구성됩니다. 아래는 GLSL Compute Shader로 벡터 덧셈을 수행하는 최소 완전 예제입니다.

/* Vulkan Compute 최소 예제 — 벡터 덧셈 (핵심 부분만 발췌) */
/* 전체 코드는 ~300줄이지만, 핵심 흐름만 표시 */

/* 1. Instance + Physical Device + Logical Device */
VkInstance instance;
vkCreateInstance(&instInfo, NULL, &instance);
VkPhysicalDevice physDev;
vkEnumeratePhysicalDevices(instance, &count, &physDev);

/* Compute 큐 패밀리 검색 */
uint32_t computeQueueFamily = UINT32_MAX;
for (uint32_t i = 0; i < queueFamilyCount; i++) {
    if (queueFamilies[i].queueFlags & VK_QUEUE_COMPUTE_BIT)
        computeQueueFamily = i;
}

VkDevice device;
vkCreateDevice(physDev, &devInfo, NULL, &device);
VkQueue computeQueue;
vkGetDeviceQueue(device, computeQueueFamily, 0, &computeQueue);

/* 2. 버퍼 생성 + 메모리 할당 (입력 A, B / 출력 C) */
VkBuffer bufA, bufB, bufC;
VkDeviceMemory memA, memB, memC;
createBuffer(device, size, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, &bufA, &memA);
createBuffer(device, size, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, &bufB, &memB);
createBuffer(device, size, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, &bufC, &memC);

/* 호스트 데이터 업로드 (HOST_VISIBLE 메모리의 경우) */
float *mapped;
vkMapMemory(device, memA, 0, size, 0, (void**)&mapped);
memcpy(mapped, hostDataA, size);
vkUnmapMemory(device, memA);

/* 3. Descriptor Set Layout + Pipeline Layout */
VkDescriptorSetLayoutBinding bindings[3] = {
    { 0, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 1, VK_SHADER_STAGE_COMPUTE_BIT },
    { 1, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 1, VK_SHADER_STAGE_COMPUTE_BIT },
    { 2, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 1, VK_SHADER_STAGE_COMPUTE_BIT },
};
/* ... DescriptorSetLayout, PipelineLayout 생성 생략 ... */

/* 4. Compute Pipeline 생성 */
VkShaderModule shaderModule;
/* shader.spv 파일 로드 → vkCreateShaderModule() */
VkComputePipelineCreateInfo pipelineInfo = {
    .sType  = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
    .stage  = {
        .sType  = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
        .stage  = VK_SHADER_STAGE_COMPUTE_BIT,
        .module = shaderModule,
        .pName  = "main",
    },
    .layout = pipelineLayout,
};
VkPipeline pipeline;
vkCreateComputePipelines(device, NULL, 1, &pipelineInfo, NULL, &pipeline);

/* 5. Command Buffer 기록 */
VkCommandBuffer cmdBuf;
vkBeginCommandBuffer(cmdBuf, &beginInfo);
vkCmdBindPipeline(cmdBuf, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
vkCmdBindDescriptorSets(cmdBuf, VK_PIPELINE_BIND_POINT_COMPUTE,
    pipelineLayout, 0, 1, &descriptorSet, 0, NULL);
vkCmdDispatch(cmdBuf, (N + 255) / 256, 1, 1);  /* 디스패치! */
vkEndCommandBuffer(cmdBuf);

/* 6. 제출 + 완료 대기 */
VkFence fence;
vkCreateFence(device, &fenceInfo, NULL, &fence);
VkSubmitInfo submitInfo = { .commandBufferCount = 1, .pCommandBuffers = &cmdBuf };
vkQueueSubmit(computeQueue, 1, &submitInfo, fence);
vkWaitForFences(device, 1, &fence, VK_TRUE, UINT64_MAX);

/* 7. 결과 읽기 */
vkMapMemory(device, memC, 0, size, 0, (void**)&mapped);
memcpy(hostResult, mapped, size);
vkUnmapMemory(device, memC);

Subgroup 연산과 고급 Compute 기능

Vulkan 1.1에서 도입된 Subgroup은 CUDA의 Warp에 대응하는 개념으로, GPU 하드웨어가 동시에 SIMD 실행하는 invocation 그룹입니다. Subgroup 크기는 GPU마다 다르며(NVIDIA: 32, AMD: 64/32, Intel: 8~32), gl_SubgroupSize와 subgroupBallot() 등 내장 함수로 워프 수준 최적화를 수행합니다.

/* Subgroup 리덕션 — 공유 메모리 없이 워크그룹 합 계산 */
#version 450
#extension GL_KHR_shader_subgroup_arithmetic : enable

layout(local_size_x = 256) in;

layout(set = 0, binding = 0) readonly buffer Input { float data[]; };
layout(set = 0, binding = 1)          buffer Output { float result[]; };

shared float partialSums[8];  /* 256/32 = 8 subgroups */

void main() {
    uint idx = gl_GlobalInvocationID.x;
    float val = data[idx];

    /* 1단계: Subgroup 내 합 (Warp 리덕션, 레지스터 수준) */
    float subgroupSum = subgroupAdd(val);

    /* 2단계: 각 Subgroup의 lane 0이 부분합 저장 */
    if (subgroupElect())
        partialSums[gl_SubgroupID] = subgroupSum;

    barrier();

    /* 3단계: 첫 Subgroup이 최종 합 계산 */
    if (gl_SubgroupID == 0) {
        float v = (gl_SubgroupInvocationID < gl_NumSubgroups)
                  ? partialSums[gl_SubgroupInvocationID] : 0.0;
        float total = subgroupAdd(v);
        if (subgroupElect())
            result[gl_WorkGroupID.x] = total;
    }
}

/* Specialization Constants — 파이프라인 생성 시 상수 주입 */
#version 450

/* 컴파일 시 값이 결정되지 않고, VkSpecializationInfo로 주입 */
layout(constant_id = 0) const uint BLOCK_SIZE = 256;
layout(constant_id = 1) const uint ALGORITHM  = 0;

layout(local_size_x_id = 0) in;  /* local_size를 상수로 지정 */

void main() {
    if (ALGORITHM == 0) {
        /* 알고리즘 A — 상수 접기로 데드 코드 제거됨 */
    } else {
        /* 알고리즘 B */
    }
}

Vulkan Compute 디버깅

Vulkan은 드라이버가 에러 검증을 하지 않는 저수준 API이므로, Validation Layers를 활성화하여 API 오용을 검출하는 것이 필수적입니다.

# Validation Layer 활성화 (환경 변수)
VK_INSTANCE_LAYERS=VK_LAYER_KHRONOS_validation ./my_vulkan_app

# GPU 선택 (다중 GPU 시스템)
VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/radeon_icd.x86_64.json ./my_app

# Mesa 드라이버 디버그 (RadV)
RADV_DEBUG=info,preoptir ./my_vulkan_app  # 드라이버 정보 + 셰이더 IR 덤프
RADV_PERFTEST=nosam ./my_vulkan_app       # 성능 실험 플래그

# Mesa 드라이버 디버그 (ANV — Intel)
INTEL_DEBUG=cs ./my_vulkan_app            # Compute Shader 컴파일 로그

# RenderDoc CLI 캡처
renderdoccmd capture --wait-for-exit ./my_vulkan_app
# → .rdc 파일 생성 → GUI에서 Compute Dispatch 분석

# vulkaninfo — 디바이스 Compute 제한 확인
vulkaninfo 2>/dev/null | grep -A5 "maxComputeWorkGroupSize"
# maxComputeWorkGroupSize[0]       = 1024
# maxComputeWorkGroupSize[1]       = 1024
# maxComputeWorkGroupSize[2]       = 64
# maxComputeWorkGroupInvocations   = 1024
# maxComputeSharedMemorySize       = 49152  (48 KB)

Vulkan Compute 실전 사용 사례

Vulkan Compute는 크로스 플랫폼 GPU 가속이 필요하면서 벤더 종속성을 피하고 싶은 시나리오에서 사용됩니다. 특히 모바일/임베디드 ML 추론, 영상 처리, 게임 엔진 물리/파티클, 그래픽 후처리 분야에서 활발합니다.

ROCm / HIP — AMD GPU 컴퓨트

ROCm(Radeon Open Compute)은 AMD의 오픈소스 GPU 컴퓨트 플랫폼입니다. HIP(Heterogeneous Interface for Portability)은 CUDA와 호환되는 API를 제공해 CUDA 코드를 AMD GPU용으로 이식하기 용이합니다. 커널 레벨에서는 /dev/kfd(KFD — Kernel Fusion Driver)를 통해 GPU와 통신하며, amdgpu DRM 드라이버가 GFX/SDMA/VCN 등 IP 블록을 관리합니다.

ROCm 소프트웨어 스택 개요

ROCm 스택은 사용자 애플리케이션부터 GPU 하드웨어까지 여러 계층으로 구성됩니다. HIP 경로와 OpenCL 경로가 공통 런타임(ROCr/HSA) 위에서 수렴하며, 최하단에서 /dev/kfd ioctl을 통해 KFD 커널 드라이버와 통신합니다.

각 계층의 역할은 다음과 같습니다. HIP Runtime은 amdhip64 라이브러리로 제공되며, CUDA와 유사한 API를 AMD GPU에서 실행할 수 있도록 변환합니다. ROCr(HSA Runtime)은 Heterogeneous System Architecture(HSA) 표준을 구현하여 큐 관리, 시그널, 에이전트 추상화를 담당합니다. Thunk(libhsakmt)는 /dev/kfd ioctl을 래핑하는 사용자 공간 라이브러리로, 커널 드라이버와의 바이너리 인터페이스를 제공합니다.

HIP 프로그래밍 — CUDA 이식성

HIP(Heterogeneous Interface for Portability)의 핵심 설계 목표는 CUDA 코드베이스를 최소한의 수정으로 AMD GPU에서 실행하는 것입니다. 대부분의 cuda* 함수는 hip* 함수로 1:1 대응되며, hipify-perl 또는 hipify-clang 도구로 자동 변환이 가능합니다.

CUDA API vs HIP API 대응표

HIP 벡터 덧셈 예제

#include <hip/hip_runtime.h>
#include <stdio.h>

// GPU 커널: 두 벡터를 더해 결과를 c에 저장합니다
__global__ void vectorAdd(const float* a, const float* b,
                          float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n)
        c[idx] = a[idx] + b[idx];
}

int main() {
    const int N = 1 << 20;  // 1M 요소
    size_t size = N * sizeof(float);

    // 호스트 메모리 할당 및 초기화
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    for (int i = 0; i < N; i++) { h_a[i] = i; h_b[i] = i * 2; }

    // 디바이스 메모리 할당
    float *d_a, *d_b, *d_c;
    hipMalloc(&d_a, size);
    hipMalloc(&d_b, size);
    hipMalloc(&d_c, size);

    // 호스트 → 디바이스 복사
    hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice);

    // 커널 실행: 그리드 크기와 블록 크기 설정
    dim3 block(256);
    dim3 grid((N + block.x - 1) / block.x);
    vectorAdd<<<grid, block>>>(d_a, d_b, d_c, N);

    // 커널 완료 대기 및 결과 복사
    hipDeviceSynchronize();
    hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost);

    printf("h_c[0]=%.0f, h_c[1]=%.0f\n", h_c[0], h_c[1]);

    // 메모리 해제
    hipFree(d_a); hipFree(d_b); hipFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

hipify 도구로 CUDA 코드 변환

hipify-perl은 CUDA 소스 파일을 파싱하여 HIP 등가 API로 텍스트 치환합니다. 더 정확한 AST 기반 변환이 필요할 경우 hipify-clang을 사용합니다.

# hipify-perl: 정규식 기반 빠른 변환 (간단한 프로젝트에 적합)
hipify-perl cuda_vector_add.cu > hip_vector_add.hip

# hipify-clang: Clang AST 기반 정밀 변환 (복잡한 템플릿/매크로에 권장)
hipify-clang --cuda-path=/usr/local/cuda \
    cuda_vector_add.cu -o hip_vector_add.hip

# 변환 통계 확인: 변환 성공/실패 항목 목록 출력
hipify-perl --print-stats cuda_project.cu

# 대규모 프로젝트: 디렉터리 일괄 변환
find . -name "*.cu" | xargs -I{} hipify-perl {} -o {}.hip

# AMD GPU 타겟 (CDNA2 — MI250X)
hipcc --offload-arch=gfx90a -O3 -o vadd hip_vector_add.hip

# AMD GPU 타겟 (RDNA3 — RX 7900 XT)
hipcc --offload-arch=gfx1100 -O3 -o vadd hip_vector_add.hip

# NVIDIA 백엔드 (CUDA 경유): HIP_PLATFORM=nvidia 환경 변수 필요
HIP_PLATFORM=nvidia hipcc --offload-arch=sm_80 -O3 -o vadd hip_vector_add.hip

# 복수 아키텍처 동시 빌드 (fat binary)
hipcc --offload-arch=gfx90a --offload-arch=gfx1100 -O3 -o vadd vadd.hip

RDNA vs CDNA 아키텍처 비교

AMD GPU 아키텍처는 크게 소비자용 게이밍 GPU인 RDNA 계열과 데이터센터·AI 컴퓨트 전용 CDNA 계열로 분류됩니다. 두 아키텍처 모두 64스레드 웨이브프론트(Wavefront)를 기본 실행 단위로 사용하지만, CDNA는 디스플레이 엔진을 제거하고 행렬 연산 가속기(Matrix Core)를 추가하였습니다.

AMD Compute Unit 내부 구조

AMD의 CU(Compute Unit)는 NVIDIA의 SM(Streaming Multiprocessor)에 대응하는 기본 컴퓨트 블록입니다. 각 CU는 4개의 SIMD 유닛(SIMD16)을 포함하며, 이들이 함께 64스레드 웨이브프론트를 처리합니다. CU 내부에는 LDS(Local Data Share — 공유 메모리), 스칼라 유닛, 스케줄러가 포함됩니다.

ROCm 개발 도구

ROCm은 GPU 프로파일링, 디버깅, 시스템 관리, 코드 변환을 위한 다양한 도구를 제공합니다. 이 도구들은 /opt/rocm/bin/에 설치됩니다.

ROCm 진단 및 운영 명령

# ── 시스템 정보 확인 ──────────────────────────────────────

# ROCm 버전 및 설치 경로 확인
cat /opt/rocm/.info/version
rocminfo | head -40

# GPU 에이전트 목록 및 속성 (VRAM, CU 수, 클럭 등)
rocminfo

# GPU 상태 모니터링: 클럭·온도·전력·VRAM 사용량
rocm-smi
rocm-smi --showclocks       # 현재 클럭 속도
rocm-smi --showmeminfo vram # VRAM 사용량
rocm-smi --showtemp         # 온도
rocm-smi --showpower        # 전력 소비

# ── 컴파일 및 실행 ───────────────────────────────────────

# hipcc 버전 및 타겟 아키텍처 확인
hipcc --version
rocminfo | grep -A 5 "gfx"

# HIP 프로그램 컴파일 (MI250X 타겟)
hipcc --offload-arch=gfx90a -O3 -o vadd vector_add.hip

# ── 성능 프로파일링 ──────────────────────────────────────

# rocprof: 커널별 실행 시간 측정
rocprof --stats ./vadd

# rocprof: HIP/HSA API 추적
rocprof --hsa-trace ./vadd
rocprof --hip-trace ./vadd

# 특정 HW 카운터 수집 (metrics.xml 파일로 지정)
rocprof -i metrics.xml -o output.csv ./vadd

# ── 커널 디버깅 ─────────────────────────────────────────

# rocgdb: GPU 커널 디버깅 (디버그 빌드 필요: -g -O0)
hipcc -g -O0 --offload-arch=gfx90a -o vadd_dbg vector_add.hip
rocgdb ./vadd_dbg
# (gdb 프롬프트 내 명령 예시)
# break vectorAdd     -- GPU 커널 중단점 설정
# info threads        -- 웨이브프론트 목록
# thread 2            -- 특정 웨이브프론트로 전환
# print idx           -- 변수 출력

# ── KFD 커널 드라이버 확인 ───────────────────────────────

# KFD 디바이스 노드 확인
ls -la /dev/kfd /dev/dri/renderD*

# amdgpu 드라이버 로드 상태
lsmod | grep amdgpu
dmesg | grep -i amdgpu | tail -20

# sysfs를 통한 GPU 정보 확인
cat /sys/class/drm/card0/device/gpu_busy_percent
cat /sys/class/drm/card0/device/mem_info_vram_used

Intel oneAPI / Level Zero

Intel oneAPI는 CPU·GPU·FPGA를 통합하는 오픈 표준 프로그래밍 플랫폼입니다. Level Zero는 그 중 GPU와 직접 통신하는 저수준 API로, DRM render node(xe/i915 드라이버) 위에서 동작합니다. SYCL/DPC++ 컴파일러가 Level Zero를 백엔드로 사용합니다.

Level Zero 아키텍처

Level Zero API 핵심 객체

/* Level Zero 핵심 API 예제 (행렬 곱) */
#include <level_zero/ze_api.h>

/* 1. 초기화 및 디바이스 선택 */
zeInit(ZE_INIT_FLAG_GPU_ONLY);

ze_driver_handle_t  hDriver;
ze_device_handle_t  hDevice;
ze_context_handle_t hContext;
zeDriverGet(&driverCount, &hDriver);
zeDeviceGet(hDriver, &deviceCount, &hDevice);
zeContextCreate(hDriver, &ctxDesc, &hContext);

/* 2. GPU 메모리 할당 */
ze_device_mem_alloc_desc_t memDesc = {
    .stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC,
    .flags = ZE_DEVICE_MEM_ALLOC_FLAG_BIAS_CACHED,
};
void *d_A, *d_B, *d_C;
zeMemAllocDevice(hContext, &memDesc, N*N*sizeof(float), 64, hDevice, &d_A);

/* 3. 커맨드 리스트 생성 및 커널 제출 */
ze_command_list_handle_t hCmdList;
zeCommandListCreate(hContext, hDevice, &cmdListDesc, &hCmdList);

zeCommandListAppendLaunchKernel(
    hCmdList, hKernel,
    &groupCount,      /* dispatch 크기 */
    hSignalEvent,     /* 완료 시그널 (ze_event) */
    0, NULL           /* 대기 이벤트 없음 */
);
zeCommandListClose(hCmdList);
zeCommandQueueExecuteCommandLists(hCmdQueue, 1, &hCmdList, hFence);
zeFenceHostSynchronize(hFence, UINT64_MAX); /* CPU 대기 */

Intel XMX — AI 행렬 가속 유닛

Intel Arc GPU(Xe-HPG 마이크로아키텍처)부터 탑재된 XMX(Xe Matrix eXtensions)는 행렬 곱을 하드웨어에서 가속하는 전용 유닛입니다. 딥러닝 추론 성능을 크게 향상시킵니다.

# Intel GPU 컴퓨트 환경 확인
# xe 드라이버 (Linux 6.2+, Intel Arc/Meteor Lake/Lunar Lake)
ls /dev/dri/renderD*
dmesg | grep -i "xe\|i915"

# Level Zero 디바이스 정보
zello_world     # Level Zero 기본 테스트 도구

# Intel GPU top (GPU 사용률 모니터링)
intel_gpu_top

# DPC++ 컴파일 (SYCL → SPIR-V → Intel GPU ISA)
icpx -fsycl -o matmul matmul.cpp

# OpenCL + Rusticl 또는 Intel NEO 드라이버로 Intel GPU 사용
OCL_ICD_FILENAMES=/usr/lib/intel-opencl/libigdrcl.so clinfo

SYCL / DPC++ 프로그래밍

SYCL은 Khronos Group이 표준화한 단일 소스 C++ 이기종 프로그래밍 모델입니다. 같은 코드가 CPU와 GPU에서 모두 실행되며, Intel의 DPC++(Data Parallel C++)는 SYCL 표준을 구현하면서 Level Zero 또는 OpenCL을 백엔드로 사용합니다.

// SYCL 벡터 덧셈 예제 (DPC++ 컴파일러)
#include <sycl/sycl.hpp>
using namespace sycl;

int main() {
    constexpr size_t N = 1024;

    // 1. 큐(Queue) 생성 — GPU 선택기 사용
    queue q{ gpu_selector_v };

    // 2. 호스트 데이터 초기화
    std::vector<float> a(N, 1.0f), b(N, 2.0f), c(N);

    {
        // 3. 버퍼(Buffer) 생성 — 호스트↔장치 데이터 소유권 관리
        buffer<float> buf_a(a.data(), range<1>{N});
        buffer<float> buf_b(b.data(), range<1>{N});
        buffer<float> buf_c(c.data(), range<1>{N});

        // 4. 커맨드 그룹(CGH) 제출
        q.submit([&](handler& cgh) {
            // 5. 접근자(Accessor) 획득 — 읽기/쓰기 권한 선언
            auto acc_a = buf_a.get_access<access::mode::read>(cgh);
            auto acc_b = buf_b.get_access<access::mode::read>(cgh);
            auto acc_c = buf_c.get_access<access::mode::write>(cgh);

            // 6. parallel_for — GPU에서 N개 워크아이템 병렬 실행
            cgh.parallel_for(range<1>{N}, [=](id<1> idx) {
                acc_c[idx] = acc_a[idx] + acc_b[idx];
            });
        });
    } // 버퍼 소멸 시 GPU→호스트 자동 전송

    // 7. 결과 검증 (큐 완료 보장 후)
    for (size_t i = 0; i < N; ++i)
        assert(c[i] == 3.0f);
    return 0;
}

xe vs i915 커널 드라이버 비교

Intel GPU는 리눅스 커널에서 두 가지 DRM 드라이버로 지원됩니다. i915는 2004년부터 유지된 레거시 드라이버이고, xe는 Linux 6.2부터 병합된 차세대 드라이버로 Xe 마이크로아키텍처 이후 GPU를 주 대상으로 합니다.

Intel GPU 컴퓨트 세대별 스펙

Intel GPU는 Gen(Generation) 번호에서 Xe(크세) 브랜드 체계로 전환하였으며, 실행 유닛(EU, Execution Unit)의 명칭도 XVE(Xe Vector Engine)로 변경되었습니다. 아래 표는 컴퓨트 관련 세대별 주요 스펙을 비교합니다.

참고 사항

관련 문서

이 주제와 관련된 다른 문서를 더 깊이 이해하고 싶다면 다음을 참고하세요.

• GPU 서브시스템 (DRM/KMS) 개요 — GPU 서브시스템 전체 구조와 핵심 요약
• DRM 코어 및 디스플레이 (KMS) — DRM 코어, KMS, atomic modesetting, 디스플레이 프로토콜
• GPU 메모리 관리 및 스케줄러 — GEM/TTM, DMA-BUF, GPU 스케줄러, GPUVM, 전원 관리
• ROCm / HIP — AMD GPU 컴퓨트 플랫폼 상세
• NPU (Neural Processing Unit) — AI 가속 디바이스
• 디바이스 드라이버 — 디바이스 드라이버 모델
• PCI / PCIe — PCI/PCIe 버스 아키텍처