ROCm / HIP — AMD GPU 컴퓨트 플랫폼

AMD ROCm(Radeon Open Compute)은 GPU 범용 컴퓨팅을 위한 완전 오픈소스 소프트웨어 스택입니다. Linux 커널의 KFD(Kernel Fusion Driver)와 amdgpu 드라이버 위에서 동작하며, HIP(Heterogeneous-compute Interface for Portability) 언어 런타임, rocBLAS/MIOpen 등의 수학 라이브러리, rocProfiler/rocTracer 프로파일링 인프라, RCCL 집합 통신 라이브러리를 포함합니다. RDNA(그래픽+컴퓨트 겸용)와 CDNA(순수 HPC 컴퓨트 전용) 아키텍처를 지원하며, MI300X 같은 최신 APU 칩렛 구조까지 다룹니다. PyTorch, TensorFlow, JAX 등 주요 AI 프레임워크가 ROCm 백엔드를 공식 지원합니다.

ROCm 발전 역사와 로드맵

ROCm(Radeon Open Compute)의 역사는 AMD가 2016년 오픈소스 GPU 컴퓨팅 생태계를 구축하겠다는 선언과 함께 시작됩니다. 당시 NVIDIA CUDA가 GPU 컴퓨팅 시장을 독점하는 상황에서, AMD는 Linux Foundation과의 협력을 통해 HSA(Heterogeneous System Architecture) 표준을 기반으로 한 완전 오픈소스 대안을 제시했습니다. 초기에는 Fiji(R9 Fury) 아키텍처를 지원하는 단일 드라이버로 출발했지만, 이후 7년간 CDNA 서버 GPU, AI 프레임워크 통합, WMMA 연산, MI300X APU 지원까지 급속히 발전했습니다.

ROCm의 근간에는 HSA(Heterogeneous System Architecture) 철학이 있습니다. HSA Consortium(AMD, ARM, Qualcomm, TI 등이 참여)이 정의한 이 표준은 CPU와 GPU가 동일한 가상 주소 공간에서 효율적으로 협력하는 모델을 규정합니다. AQL(Architected Queuing Language) 패킷 형식, 신호(Signal) 객체, 통합 메모리(SVM) 모두 HSA 1.1 사양에서 유래했습니다. ROCm 1.0이 이 사양을 최초로 오픈소스로 구현한 프로덕션 스택입니다.

컴파일러 측면에서의 주요 전환점은 HCC(Heterogeneous Compute Compiler)에서 HIP/hipcc로의 이행입니다. ROCm 3.x까지 AMD는 HCC를 주 GPU 컴파일러로 사용했으나, HCC는 CUDA 코드 이식성이 부족하고 LLVM 업스트림과 괴리가 있었습니다. ROCm 3.3부터 hipcc(LLVM Clang 기반)가 기본 컴파일러가 되면서 CUDA 코드의 직접 이식이 크게 쉬워졌고, LLVM 업스트림 AMDGPU 백엔드와의 통합도 강화되었습니다. ROCm 5.0에서 HCC는 공식 deprecated되었습니다.

버전별 주요 이정표

# ROCm 설치 버전 및 컴포넌트 확인
cat /opt/rocm/.info/version
# 출력: 6.1.0-1

# 설치된 모든 ROCm 패키지 나열
dpkg -l | grep rocm
# 또는 RPM 기반 시스템:
rpm -qa | grep rocm

# 컴파일러 버전 확인
/opt/rocm/bin/hipcc --version
# 출력: HIP version: 6.1.40091-5e38b8620
#        clang version 17.0.0 (amd-mainline-open ...)

# 주요 라이브러리 버전 확인
apt show rocblas 2>/dev/null | grep Version
apt show miopen-hip 2>/dev/null | grep Version
apt show rccl 2>/dev/null | grep Version

# rocm-smi로 현재 드라이버 버전 확인
rocm-smi --showdriverversion
# Driver Version: 6.7.0
# ROCm Version: 6.1.0

# KFD ioctl 버전 (ROCm ABI 호환성 확인)
cat /sys/class/kfd/kfd/version
# 출력: 1.14

ROCm vs CUDA 생태계 비교

ROCm은 CUDA와의 호환성을 명시적인 설계 목표로 삼습니다. HIP API는 CUDA Runtime API와 90% 이상 일치하며, hipify-perl과 hipify-clang 도구를 사용하면 대부분의 CUDA 소스 코드를 자동 변환할 수 있습니다. 단, CUDA device-specific intrinsics나 PTX 인라인 어셈블리는 수동 변환이 필요하고, CUDA 고유 기능(cooperative groups 일부, CUDA graphs 완전 지원 등)은 ROCm에서 구현 범위가 다를 수 있습니다.

향후 로드맵과 전략 방향

2024년 이후 ROCm의 핵심 전략은 크게 세 방향으로 정리됩니다. 첫째, AI/LLM 추론 최적화입니다. MI300X의 192 GB HBM3 용량을 활용한 대형 언어 모델 단일 노드 추론 시나리오에서 vLLM, SGLang, llama.cpp 등이 ROCm 백엔드를 공식 지원하기 시작했습니다. hipBLASLt의 Epilogue Fusion과 Custom Kernel 지원이 Transformer attention 계층 최적화를 크게 개선했습니다.

둘째, LLVM 업스트림 통합 강화입니다. ROCm 팀은 자체 LLVM 포크에 의존하는 대신 AMDGPU 백엔드 코드를 지속적으로 LLVM 메인라인에 기여하고 있습니다. 이를 통해 Clang/LLVM 최신 버전을 사용하는 일반 개발자도 ROCm 없이 AMDGPU 코드를 컴파일할 수 있게 됩니다.

셋째, PyTorch/JAX 우선 지원입니다. AMD는 Meta와의 협력을 통해 PyTorch ROCm 빌드의 성능 격차를 CUDA 대비 5% 이내로 좁히는 것을 목표로 합니다. FlashAttention-ROCm, Triton-ROCm 등의 핵심 커널 라이브러리가 공식 AMD 팀에 의해 유지 관리되고 있습니다.

ROCm 소프트웨어 스택

ROCm은 단일 라이브러리가 아니라 여러 레이어로 구성된 소프트웨어 스택입니다. 맨 아래의 Linux 커널 드라이버부터 최상위 AI 프레임워크까지 각 레이어가 명확한 인터페이스 계약을 통해 연결됩니다. 이 계층 구조를 이해하면 성능 문제나 API 오류가 발생했을 때 어느 레이어를 조사해야 하는지 즉시 판단할 수 있습니다.

스택의 핵심 철학은 레이어 간 명확한 계약입니다. HIP API는 AMD GPU와 NVIDIA GPU 양쪽에서 동작하는 이식성 레이어이며, ROCr(HSA Runtime)은 AQL(Architected Queuing Language) 패킷을 통해 GPU에 커맨드를 제출합니다. 커널 레이어에서는 KFD가 사용자 공간 런타임의 ioctl 요청을 받아 amdgpu 드라이버의 컴퓨트 큐를 설정합니다.

# ROCm 버전 및 GPU 상태 확인
rocm-smi --showdriverversion
# 출력 예: ROCm version: 6.1.0

rocm-smi --showproductname
# 출력 예: GPU[0]: Instinct MI300X

rocm-smi --showmeminfo vram
# VRAM 사용량 조회 (Used / Total in MB)

# ROCm 전체 컴포넌트 버전 확인
cat /opt/rocm/.info/version
# 출력: 6.1.0-1

# amdgpu 드라이버 로드 확인
lsmod | grep amdgpu
dmesg | grep -i kfd | head -20

HSA/AQL 패킷 구조 심층 분석

AQL(Architected Queuing Language)은 HSA 사양이 정의하는 GPU 커맨드 제출 형식입니다. CUDA의 경우 CPU가 GPU 커맨드 링 버퍼에 드라이버 내부 포맷으로 패킷을 써야 했지만, AQL은 이 포맷을 공개 사양으로 표준화하여 사용자 공간에서 직접 GPU 큐에 패킷을 쓸 수 있게 했습니다. 이 설계 덕분에 커널 진입 없이 GPU 작업을 제출할 수 있어 레이턴시가 크게 줄어듭니다.

AQL 큐는 mmap된 메모리 링 버퍼입니다. 사용자 프로세스가 패킷을 채운 후 도어벨(doorbell) MMIO 레지스터에 새 쓰기 포인터(wptr)를 기록하면, GPU의 Command Processor(CP)가 이를 감지하여 패킷을 소비합니다. 이 과정에서 시스템 콜이나 드라이버 개입이 전혀 없습니다. ROCr HSA 런타임은 이 AQL 큐 관리를 캡슐화하고, HIP 런타임은 ROCr 위에서 hipLaunchKernel을 구현합니다.

AQL 패킷은 64바이트 고정 크기이며, 4가지 타입이 존재합니다. Kernel Dispatch Packet은 GPU 커널을 실행할 때 사용하고, Agent Dispatch Packet은 특수 에이전트 작업 요청에, Barrier-AND/OR Packet은 의존성 동기화에 사용됩니다. 커널 디스패치 패킷 하나가 hipLaunchKernel 호출 하나에 해당합니다.

AQL 패킷 구조 상세

도어벨 메커니즘

/* AQL Kernel Dispatch Packet 직접 구성 (C 로우레벨 예제) */
#include <hsa/hsa.h>
#include <stdint.h>
#include <string.h>

/* HSA 1.1 사양 4.2절: Kernel Dispatch Packet */
typedef struct __attribute__((packed)) {
    uint16_t     header;               /* [7:0]=type [8]=barrier [12:9]=fence */
    uint16_t     setup;
    uint16_t     workgroup_size_x;
    uint16_t     workgroup_size_y;
    uint16_t     workgroup_size_z;
    uint16_t     reserved0;
    uint32_t     grid_size_x;
    uint32_t     grid_size_y;
    uint32_t     grid_size_z;
    uint32_t     private_segment_size;
    uint32_t     group_segment_size;
    uint64_t     kernel_object;
    uint64_t     kernarg_address;
    uint64_t     reserved2;
    hsa_signal_t completion_signal;
} aql_kdp_t;
_Static_assert(sizeof(aql_kdp_t) == 64, "must be 64 bytes");

void submit_kernel(hsa_queue_t *q, uint64_t code, void *args,
                   uint32_t gx, uint32_t bx, hsa_signal_t done)
{
    /* 1. 원자적으로 wptr 인덱스 획득 */
    uint64_t wi = hsa_queue_add_write_index_relaxed(q, 1);
    while (wi - hsa_queue_load_read_index_acquire(q) >= q->size)
        ;

    /* 2. 슬롯 포인터 계산 */
    aql_kdp_t *p = (aql_kdp_t *)q->base_address + (wi & (q->size - 1));

    /* 3. 패킷 필드 작성 (header는 반드시 마지막에!) */
    memset(p, 0, 64);
    p->setup             = 1;
    p->workgroup_size_x  = (uint16_t)bx;
    p->workgroup_size_y  = p->workgroup_size_z = 1;
    p->grid_size_x       = gx;
    p->grid_size_y       = p->grid_size_z = 1;
    p->kernel_object     = code;
    p->kernarg_address   = (uint64_t)args;
    p->completion_signal = done;

    /* 4. header 원자적 쓰기 */
    uint16_t hdr =
        (uint16_t)HSA_PACKET_TYPE_KERNEL_DISPATCH
        | (1u << HSA_PACKET_HEADER_BARRIER)
        | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE)
        | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
    __atomic_store_n(&p->header, hdr, __ATOMIC_RELEASE);

    /* 5. 도어벨 MMIO 쓰기 (syscall 없음!) */
    hsa_signal_store_relaxed(q->doorbell_signal, (hsa_signal_value_t)wi);
}

/* 완료 대기 */
hsa_signal_wait_scacquire(done, HSA_SIGNAL_CONDITION_LT, 1,
                          UINT64_MAX, HSA_WAIT_STATE_BLOCKED);

/* ROCr 런타임 큐 생성 및 완료 대기 패턴 */
hsa_init();

hsa_agent_t gpu;
hsa_iterate_agents(pick_gpu_callback, &gpu);

/* AQL 큐 생성: KFD_IOC_CREATE_QUEUE -> MQD 할당 -> 도어벨 mmap */
hsa_queue_t *q;
hsa_queue_create(gpu, 256, HSA_QUEUE_TYPE_SINGLE,
                 NULL, NULL, UINT32_MAX, UINT32_MAX, &q);

/* 완료 신호 (초기값=1, GPU 완료 후 0이 됨) */
hsa_signal_t done;
hsa_signal_create(1, 0, NULL, &done);

/* 커널 인자 버퍼: GPU 접근 가능 메모리에 할당 */
struct { float *a, *b, *c; int n; } kargs = {d_a, d_b, d_c, N};
void *kernarg;
hsa_memory_allocate(kernarg_region, sizeof(kargs), &kernarg);
memcpy(kernarg, &kargs, sizeof(kargs));

submit_kernel(q, kernel_code_ptr, kernarg, N, 256, done);
hsa_signal_wait_scacquire(done, HSA_SIGNAL_CONDITION_LT, 1,
                          UINT64_MAX, HSA_WAIT_STATE_BLOCKED);

hsa_signal_destroy(done);
hsa_queue_destroy(q);
hsa_shut_down();

KFD (Kernel Fusion Driver)

KFD(Kernel Fusion Driver)는 Linux 커널의 amdgpu 드라이버 내부에 통합된 컴퓨트 전용 커널 모듈입니다. drivers/gpu/drm/amd/amdkfd/에 위치하며, /dev/kfd 문자 디바이스를 통해 사용자 공간 ROCm 런타임에 GPU 컴퓨트 자원을 노출합니다. KFD의 핵심 설계 목표는 다중 프로세스 안전한 GPU 접근과 HSA 표준 준수입니다.

KFD는 프로세스별로 독립적인 GPU 주소 공간을 관리합니다. IOMMU를 활용한 PASID(Process Address Space ID) 메커니즘으로 한 프로세스의 GPU 메모리 접근이 다른 프로세스의 메모리를 침범할 수 없도록 격리합니다. 또한 Compute Queue(MQD/HQD)를 프로세스별로 할당하여 GPU 스케줄러가 공정하게 작업을 처리하도록 합니다.

# /dev/kfd 존재 및 권한 확인
ls -la /dev/kfd
# crw-rw---- 1 root render 234, 0 /dev/kfd
# render 그룹 멤버십 필요: sudo usermod -aG render $USER

# KFD topology: 연결된 GPU 정보 조회
ls /sys/bus/event_source/devices/ | grep kfd
cat /sys/class/kfd/kfd/topology/nodes/1/name
# 출력 예: gfx942 (MI300X의 GFX IP 버전)

cat /sys/class/kfd/kfd/topology/nodes/1/mem_banks/0/size_in_bytes
# HBM3 총 바이트 크기

# KFD ioctl 버전 확인 (ROCr이 호출하는 방식)
# GET_VERSION ioctl → struct kfd_ioctl_get_version_args
python3 -c "
import fcntl, struct, os
fd = os.open('/dev/kfd', os.O_RDWR | os.O_CLOEXEC)
KFD_IOC_GET_VERSION = 0x4008AE01  # _IOWR(0xAE, 0x01, ...)
buf = struct.pack('II', 0, 0)
res = fcntl.ioctl(fd, KFD_IOC_GET_VERSION, bytearray(buf))
maj, minor = struct.unpack('II', res)
print(f'KFD version: {maj}.{minor}')
os.close(fd)
"

# KFD topology sysfs: GPU 노드 상세 정보
for node in /sys/class/kfd/kfd/topology/nodes/*/; do
    echo "=== Node: $node ==="
    cat "$node/name" 2>/dev/null
    cat "$node/properties" 2>/dev/null | grep -E "simd_count|max_waves|mem_banks"
done

# 출력 예 (MI300X, 8 XCD):
# simd_count 1792      (CU 수 × 4 SIMD)
# max_waves_per_simd 10
# max_flat_address_watch 4

KFD ioctl 내부 동작 상세

KFD의 핵심 인터페이스는 /dev/kfd에 대한 ioctl 호출입니다. 유저스페이스 HSA 런타임은 이 ioctl을 통해 큐 생성, 메모리 할당, 이벤트 등록 등 모든 GPU 자원 관리 작업을 수행합니다. 커널 내부에서는 kfd_ioctl.c가 ioctl을 수신하고 각 명령에 해당하는 핸들러를 호출하며, 핸들러는 결국 amdgpu 백엔드를 통해 실제 하드웨어 자원을 조작합니다.

/* KFD CREATE_QUEUE ioctl 처리 흐름 (커널 내부 단순화) */
/* drivers/gpu/drm/amd/amdkfd/kfd_ioctl.c */

static long kfd_ioctl(struct file *filep, unsigned int cmd, unsigned long arg)
{
    const struct kfd_ioctl_desc *ioctl = NULL;
    kfd_ioctl_t *func;
    void *kdata = NULL;
    unsigned int nr = _IOC_NR(cmd);

    /* ioctl 번호로 핸들러 테이블에서 함수 포인터 조회 */
    if (nr < ARRAY_SIZE(kfd_ioctl_funcs))
        ioctl = &kfd_ioctl_funcs[nr];

    func = ioctl->func;
    /* 유저스페이스 인자를 커널 스택으로 복사 */
    copy_from_user(kdata, (void __user *)arg, ioctl->usize);

    return func(filep, kdata);   /* 예: kfd_ioctl_create_queue() */
}

/* CREATE_QUEUE 핸들러 - 큐 생성 주요 단계 */
static int kfd_ioctl_create_queue(struct file *filep, void *data)
{
    struct kfd_ioctl_create_queue_args *args = data;
    struct kfd_process *p;
    struct kfd_dev *dev;
    int err, queue_id;

    /* filep에서 현재 kfd_process 조회 (프로세스 격리 단위) */
    p = kfd_get_process(current());
    /* gpu_id로 대상 GPU 장치 조회 */
    dev = kfd_device_by_id(args->gpu_id);

    /* MQD 메모리 할당 및 큐 링 버퍼 매핑 */
    err = pqm_create_queue(&p->pqm, dev, filep, &args->queue_properties,
                            &queue_id, NULL, NULL, NULL, NULL);

    args->queue_id = queue_id;
    /* HQD에 MQD 등록 → CP가 AQL 패킷 소비 시작 */
    return err;
}

/* KFD GPU 메모리 할당 경로 */
/* ALLOC_MEMORY_OF_GPU ioctl → amdgpu_amdkfd_gpuvm_alloc_memory */

static int kfd_ioctl_alloc_memory_of_gpu(struct file *filep, void *data)
{
    struct kfd_ioctl_alloc_memory_of_gpu_args *args = data;
    struct kfd_process *p = kfd_get_process(current());
    struct kfd_dev *dev = kfd_device_by_id(args->gpu_id);
    void *mem;
    uint64_t offset;

    /* amdgpu TTM을 통한 VRAM/GTT 버퍼 할당 */
    int err = amdgpu_amdkfd_gpuvm_alloc_memory_of_gpu(
        dev->adev,        /* amdgpu_device 포인터 */
        args->va_addr,    /* 가상 주소 힌트 */
        args->size,       /* 할당 크기 (바이트) */
        p->mm,            /* 프로세스 VM */
        &mem,             /* 출력: BO 핸들 */
        &offset,          /* 출력: mmap offset */
        args->flags       /* ALLOC_MEM_FLAGS_VRAM 등 */
    );
    if (err) return err;

    /* kfd_process_device의 BO 목록에 등록 (eviction 추적용) */
    err = kfd_process_device_create_obj_handle(
        kfd_get_process_device_data(dev, p), mem);

    args->handle = err;
    args->mmap_offset = offset;
    return 0;
}

/* 할당 플래그 비트 필드 */
#define KFD_IOC_ALLOC_MEM_FLAGS_VRAM        (1 << 0)  /* GPU VRAM에 할당 */
#define KFD_IOC_ALLOC_MEM_FLAGS_GTT         (1 << 1)  /* CPU 접근 가능 GTT */
#define KFD_IOC_ALLOC_MEM_FLAGS_USERPTR     (1 << 2)  /* 기존 CPU 메모리 핀닝 */
#define KFD_IOC_ALLOC_MEM_FLAGS_DOORBELL    (1 << 3)  /* doorbell 페이지 */
#define KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE    (1 << 31) /* GPU 쓰기 가능 */
#define KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE  (1 << 30) /* 셰이더 실행 가능 */

KFD 프로세스 Eviction과 복구

KFD 프로세스 eviction은 시스템 메모리 압박 시 GPU VRAM을 회수하기 위해 실행 중인 GPU 프로세스를 일시적으로 중단하고 그 메모리를 시스템 RAM으로 이동하는 메커니즘입니다. ROCm 런타임은 이 과정에서 애플리케이션이 크래시하지 않도록 투명하게 처리합니다. amdgpu TTM의 메모리 회수 콜백이 KFD에 eviction을 요청하며, KFD는 해당 프로세스의 모든 큐를 중단하고 BO(Buffer Object)를 GTT(CPU 접근 가능 영역)로 이동시킵니다.

/* KFD Eviction 흐름 (간략화) */
/* drivers/gpu/drm/amd/amdkfd/kfd_process.c */

/* TTM이 메모리 회수 필요 시 호출하는 진입점 */
int kgd2kfd_schedule_evict_and_restore_process(
    struct mm_struct *mm, struct dma_fence *fence)
{
    struct kfd_process *p = kfd_lookup_process_by_mm(mm);
    if (!p) return -ESRCH;

    if (fence)
        dma_fence_add_callback(fence, &p->eviction_fence_cb, evict_callback);
    else
        queue_delayed_work(kfd_wq, &p->eviction_work, 0);
    return 0;
}

/* 실제 eviction 워크큐 함수 */
static void evict_process_worker(struct work_struct *work)
{
    struct kfd_process *p =
        container_of(work, struct kfd_process, eviction_work.work);

    /* 1단계: 모든 큐를 HQD에서 unmap → GPU 실행 중단 */
    kfd_process_evict_queues(p, KFD_QUEUE_EVICTION_TRIGGER_TTM);

    /* 2단계: 모든 BO를 VRAM → GTT로 이동 */
    amdgpu_amdkfd_evict_userptr(p->mm, p->priv);

    /* 3단계: restore 워크 스케줄 (메모리 여유 생기면 복구) */
    queue_delayed_work(kfd_restore_wq, &p->restore_work, KFD_RESTORE_DELAY_MS);
    p->process_status = KFD_PROC_STATUS_EVICTED;
}

/* 복구 워크큐 함수 */
static void restore_process_worker(struct work_struct *work)
{
    struct kfd_process *p =
        container_of(work, struct kfd_process, restore_work.work);
    int ret;

    p->process_status = KFD_PROC_STATUS_RESTORING;

    /* 1단계: GTT → VRAM으로 BO 다시 이동 */
    ret = amdgpu_amdkfd_restore_userptr(p->mm, p->priv);
    if (ret) {
        /* VRAM 여전히 부족: 재시도 */
        queue_delayed_work(kfd_restore_wq, &p->restore_work, KFD_RESTORE_RETRY_MS);
        return;
    }

    /* 2단계: GPU 페이지 테이블 재구성 및 큐 복구 */
    kfd_process_restore_queues(p);
    p->process_status = KFD_PROC_STATUS_RUNNING;
}

amdgpu 커널 드라이버

amdgpu는 Linux mainline에 포함된 AMD GPU의 통합 DRM 드라이버입니다. drivers/gpu/drm/amd/ 하위에 위치하며, IP 블록 기반 설계로 GFX, SDMA, VCN, DCN, PSP, SMU, NBIO 등 각 하드웨어 IP를 독립적으로 초기화하고 관리합니다. 각 IP 블록은 amdgpu_ip_block_version 구조체를 통해 드라이버에 등록되며, GPU 세대(GFX9/GFX10/GFX11/GFX12)에 따라 다른 IP 구현을 자동으로 선택합니다.

# amdgpu 모듈 파라미터 확인
modinfo amdgpu | grep -E "^parm:"

# 주요 파라미터 설정 예 (/etc/modprobe.d/amdgpu.conf)
# options amdgpu ip_block_mask=0xff  # 특정 IP 블록 비트마스크로 활성화
# options amdgpu gpu_recovery=1     # GPU hang 시 자동 복구
# options amdgpu gttsize=8192       # GTT 최대 크기 (MB)
# options amdgpu exp_hw_support=1   # 실험적 GPU 지원 활성화
# options amdgpu ppfeaturemask=0xffffffff  # 전원 관리 기능 비트마스크
# options amdgpu runpm=1            # runtime PM 활성화

# 현재 로드된 파라미터 확인
cat /sys/module/amdgpu/parameters/gpu_recovery
cat /sys/module/amdgpu/parameters/gttsize

# dmesg에서 amdgpu 초기화 로그 확인
dmesg | grep -E "amdgpu|kfd" | head -50

# 정상 초기화 예시 출력:
# [    5.234] amdgpu 0000:03:00.0: amdgpu: Trusted Memory Zone (TMZ) enabled
# [    5.310] amdgpu 0000:03:00.0: amdgpu: VRAM: 49152M 0x0000008000000000 - 0x000000BFFFFFFFFF (49152M used)
# [    5.311] amdgpu 0000:03:00.0: amdgpu: GART: 512M 0x0000000000000000 - 0x000000001FFFFFFF
# [    5.312] amdgpu 0000:03:00.0: amdgpu: PSP is initialized, GFX version: 9.4.2
# [    5.890] kfd kfd: Initialized module
# [    5.891] kfd kfd: Added apertures for GPU 0000:03:00.0

# amdgpu debugfs 활용
ls /sys/kernel/debug/dri/0/
# amdgpu_fence_info  amdgpu_vm_info  amdgpu_gtt_mm  amdgpu_eviction_stats ...
cat /sys/kernel/debug/dri/0/amdgpu_fence_info

AMDGCN ISA 개요

AMDGCN(AMD Graphics Core Next)은 AMD GPU의 실제 하드웨어 ISA(Instruction Set Architecture)입니다. CUDA의 PTX가 가상 ISA인 것과 달리, AMDGCN은 진짜 하드웨어 ISA로 hipcc가 생성하는 최종 바이너리가 바로 이 명령어 집합으로 구성됩니다. 커널 개발자가 AMDGCN ISA를 이해하면 컴파일러가 생성한 코드를 직접 분석하고, 레지스터 압박이나 메모리 접근 패턴을 최적화하거나, 인라인 어셈블리로 특수 명령어를 활용할 수 있습니다.

AMDGCN ISA의 핵심 특징은 스칼라(SALU)와 벡터(VALU) 파이프라인의 분리입니다. 스칼라 연산은 Wavefront 내 모든 스레드가 동일하게 수행하는 제어 흐름과 상수 계산에 사용되고, 벡터 연산은 각 스레드가 독립적으로 다른 데이터에 대해 수행하는 SIMD 연산입니다. SGPR(스칼라 레지스터)은 Wavefront당 공유되며, VGPR(벡터 레지스터)은 스레드별 독립 레인을 가집니다.

또 하나의 중요한 특징은 Wave32 vs Wave64 모드입니다. GFX9(Vega/CDNA) 계열은 64개 스레드 Wavefront(Wave64)만 지원하지만, GFX10(RDNA)/GFX11(RDNA3)은 32개 스레드 Wavefront(Wave32)를 기본으로 사용합니다. Wave32는 분기 비용이 낮고 점유율 관리가 유리하지만, 64개 float 연산당 레지스터 오버헤드가 증가합니다.

명령어 범주

레지스터 파일과 점유율

AMDGCN에서 성능 최적화의 핵심은 레지스터 압박(register pressure)과 점유율(occupancy) 간의 균형입니다. 각 CU는 VGPR 풀을 여러 Wavefront가 공유합니다. CDNA2(MI200) 기준으로 CU당 512개의 AGPR(Accumulation VGPR)을 포함한 총 512 VGPR을 보유합니다. Wavefront당 사용하는 VGPR 수가 적을수록 더 많은 Wavefront가 CU에 동시에 상주할 수 있어 메모리 레이턴시를 숨길 수 있습니다.

# hipcc로 컴파일 시 임시 어셈블리 파일 저장
hipcc -save-temps -O3 --offload-arch=gfx942 vectoradd.hip -o vectoradd

# 생성된 .s 파일 (ISA 어셈블리) 확인
ls /tmp/vectoradd-*.s
cat /tmp/vectoradd-hip-amdgcn-amd-amdhsa-gfx942.s

# llvm-objdump으로 오브젝트 파일 역어셈블
/opt/rocm/llvm/bin/llvm-objdump -d vectoradd | grep -A 50 "vectorAdd:"

# 커널의 VGPR/SGPR 사용량 확인 (metadata에서)
/opt/rocm/llvm/bin/llvm-objdump --amdgpu-kernel-metadata vectoradd | grep -E "vgpr|sgpr|lds"
# 출력 예:
#   .vgpr_count:     12    (Wavefront당 VGPR 12개 사용)
#   .sgpr_count:     20    (SGPR 20개 사용)
#   .lds_size:       0     (LDS 미사용)

# 점유율 계산 (gfx942, 512 VGPR pool)
# VGPR_per_wf=12 → 할당 단위 4로 올림 → 12 → 512/12=42 wavefront/CU 가능
python3 -c "
vgpr_used = 12
vgpr_granule = 4  # GFX9 할당 단위
vgpr_alloc = ((vgpr_used + vgpr_granule - 1) // vgpr_granule) * vgpr_granule
max_wf_per_cu = 512 // vgpr_alloc
print(f'VGPR 사용: {vgpr_used} -> 할당: {vgpr_alloc} -> 최대 WF/CU: {max_wf_per_cu}')
"

인라인 어셈블리 활용

# AMDGCN 어셈블리 예: 벡터 덧셈 커널 (gfx942)
# 다음은 hipcc -save-temps 로 생성된 .s 파일 구조 예시
#
# .amdgpu_metadata 섹션: 커널 메타데이터
#   .group_segment_fixed_size: 0   (LDS 미사용)
#   .private_segment_fixed_size: 0  (scratch 미사용)
#   .vgpr_count: 10
#   .sgpr_count: 22
#
# 커널 진입점 (v_add_f32 예시):
#   s_load_dwordx4  s[0:3], s[4:5], 0x0    ; 커널 인자 로드 (SGPR)
#   s_load_dwordx2  s[4:5], s[6:7], 0x18   ; N 값 로드
#   v_lshlrev_b32   v0, 2, v0              ; threadIdx * 4 (byte offset)
#   s_waitcnt       lgkmcnt(0)             ; 스칼라 메모리 로드 완료 대기
#   global_load_dword v1, v[0:1], s[0:1]  ; A[idx] 로드
#   global_load_dword v2, v[0:1], s[2:3]  ; B[idx] 로드
#   s_waitcnt       vmcnt(0)              ; 벡터 메모리 로드 완료 대기
#   v_add_f32       v1, v1, v2            ; C = A + B
#   global_store_dword v[0:1], v1, s[4:5] ; C[idx] 저장
#   s_endpgm                              ; 커널 종료

# s_waitcnt 의미: 특정 메모리 연산 완료까지 대기
#   vmcnt(N): N개의 벡터 메모리 연산이 완료될 때까지 대기
#   lgkmcnt(N): N개의 LDS/GDS/스칼라 메모리 연산이 완료될 때까지 대기
#   expcnt(N): N개의 export 연산이 완료될 때까지 대기
#   s_waitcnt 0: 모든 종류의 메모리 연산 완료 대기 (완전 동기화)

/* HIP 커널 내 AMDGCN 인라인 어셈블리 활용 예제 */
#include <hip/hip_runtime.h>

/* 예제 1: EXEC 마스크 읽기 (활성화된 레인 수 확인) */
__device__ uint64_t get_exec_mask()
{
    uint64_t exec;
    __asm__ volatile("s_mov_b64 %0, exec" : "=s"(exec));
    return exec;
}

/* 예제 2: s_barrier 명령어로 Wavefront 간 동기화 */
__device__ void gpu_barrier()
{
    __asm__ volatile("s_barrier" ::: "memory");
    /* __syncthreads()와 동일 — LDS 데이터 일관성 보장 */
}

/* 예제 3: v_cvt_f16_f32 — FP32를 FP16으로 변환 (WMMA 입력용) */
__device__ __half fp32_to_fp16(float x)
{
    uint32_t result;
    __asm__ volatile(
        "v_cvt_f16_f32 %0, %1"
        : "=v"(result)
        : "v"(x)
    );
    return (__half)(result & 0xFFFF);
}

/* 예제 4: ds_bpermute — LDS 배열을 통한 레인 간 데이터 교환 */
__device__ float lane_shuffle(float val, int src_lane)
{
    float result;
    __asm__ volatile(
        "ds_bpermute_b32 %0, %1, %2"
        : "=v"(result)
        : "v"(src_lane * 4), "v"(val)
    );
    return result;
}

/* 예제 5: v_mad_f32 — FMA(Fused Multiply-Add) 강제 사용 */
__device__ float fma_explicit(float a, float b, float c)
{
    float result;
    __asm__ volatile(
        "v_fma_f32 %0, %1, %2, %3"
        : "=v"(result)
        : "v"(a), "v"(b), "v"(c)
    );
    return result;
}

/* 예제 6: s_waitcnt vmcnt(0) — 메모리 펜스 */
__device__ void vmem_fence()
{
    __asm__ volatile("s_waitcnt vmcnt(0)" ::: "memory");
    /* 모든 벡터 메모리 연산 완료 보장 — __threadfence()의 저수준 구현 */
}

# AMDGCN 커널 분석 실무 워크플로우

# 1. 커널 컴파일 및 임시 파일 저장
hipcc -save-temps -O3 --offload-arch=gfx942 -o kernel kernel.hip

# 2. GFX 코드 오브젝트 추출
/opt/rocm/llvm/bin/clang-offload-bundler --unbundle     --type=o --targets=hip-amdgcn-amd-amdhsa-gfx942     --inputs=kernel --outputs=kernel_gfx942.o

# 3. 역어셈블 (상세 주석 포함)
/opt/rocm/llvm/bin/llvm-objdump -d --disassemble-all kernel_gfx942.o | head -100

# 4. 커널 메타데이터 확인 (VGPR/SGPR/LDS 사용량)
/opt/rocm/llvm/bin/llvm-objdump --amdgpu-kernel-metadata kernel_gfx942.o

# 5. rocprof로 실제 레지스터 사용 통계 수집
rocprof --stats -i rocprof_input.txt ./kernel
# rocprof_input.txt 내용 예:
# pmc: VFetchInsts VWriteInsts LDSBankConflict

# 6. omniperf 전체 분석 (레지스터 압박 + 점유율 + 메모리 접근 패턴)
omniperf profile -n vectoradd -- ./vectoradd
omniperf analyze -p workloads/vectoradd/ --roof-only

HIP 프로그래밍 모델

HIP(Heterogeneous-compute Interface for Portability)는 AMD가 개발한 GPU 프로그래밍 모델로, CUDA와 거의 동일한 문법과 API를 제공합니다. hipcc 컴파일러로 AMD GPU에서는 amdgpu 백엔드를, NVIDIA GPU에서는 CUDA 백엔드를 투명하게 사용합니다. 이로써 단일 소스 코드베이스로 두 플랫폼을 동시에 지원할 수 있습니다.

HIP의 실행 모델은 CUDA와 동일한 계층 구조를 따릅니다. 최상위는 Grid(전체 커널 실행 단위)이며, Grid는 여러 Block으로, Block은 여러 Thread로 구성됩니다. AMD GPU에서는 32개 또는 64개의 Thread가 하나의 Wavefront(NVIDIA의 Warp에 해당)를 이루어 SIMD 방식으로 동시에 실행됩니다.

/* HIP 벡터 덧셈 예제: CUDA 문법과 99% 동일 */
#include <hip/hip_runtime.h>
#include <cstdio>

/* GPU 커널 함수: __global__ 키워드로 GPU에서 실행됨을 표시 */
__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 = 1024 * 1024;
    const size_t bytes = N * sizeof(float);

    /* GPU 메모리 할당 */
    float *d_A, *d_B, *d_C;
    hipMalloc(&d_A, bytes);
    hipMalloc(&d_B, bytes);
    hipMalloc(&d_C, bytes);

    /* 호스트→GPU 데이터 복사 (초기화 코드 생략) */
    /* hipMemcpy(d_A, h_A, bytes, hipMemcpyHostToDevice); */

    /* 커널 실행: <<>> 문법 */
    int blockSize = 256;
    int gridSize  = (N + blockSize - 1) / blockSize;
    vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

    /* 커널 완료 대기 */
    hipDeviceSynchronize();

    /* 메모리 해제 */
    hipFree(d_A); hipFree(d_B); hipFree(d_C);
    return 0;
}

/* 공유 메모리(LDS) 활용 예: 타일 행렬 곱셈 */
#define TILE 16

__global__ void tiledMatMul(const float *A, const float *B,
                              float *C, int N)
{
    /* __shared__: LDS(Local Data Share) 할당 — CU 내 모든 스레드 공유 */
    __shared__ float sA[TILE][TILE];
    __shared__ float sB[TILE][TILE];

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;
    float sum = 0.0f;

    for (int t = 0; t < N / TILE; ++t) {
        /* 타일 로드 */
        sA[ty][tx] = A[row * N + t * TILE + tx];
        sB[ty][tx] = B[(t * TILE + ty) * N + col];
        /* 블록 내 스레드 동기화 (LDS 쓰기 완료 보장) */
        __syncthreads();

        for (int k = 0; k < TILE; ++k)
            sum += sA[ty][k] * sB[k][tx];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

/* 스트림(비동기 실행)과 이벤트 사용 예 */
hipStream_t stream0, stream1;
hipStreamCreate(&stream0);
hipStreamCreate(&stream1);

hipEvent_t startEvent, stopEvent;
hipEventCreate(&startEvent);
hipEventCreate(&stopEvent);

/* 스트림 0: 비동기 복사 후 커널 실행 */
hipMemcpyAsync(d_A, h_A, bytes, hipMemcpyHostToDevice, stream0);
hipEventRecord(startEvent, stream0);
myKernel<<<grid, block, 0, stream0>>>(d_A, d_C, N);
hipEventRecord(stopEvent, stream0);

/* 스트림 1: 독립적 작업 동시 실행 */
hipMemcpyAsync(d_B, h_B, bytes, hipMemcpyHostToDevice, stream1);
myKernel<<<grid, block, 0, stream1>>>(d_B, d_D, N);

/* 두 스트림 모두 완료 대기 */
hipStreamSynchronize(stream0);
hipStreamSynchronize(stream1);

/* 실행 시간 측정 */
float ms;
hipEventElapsedTime(&ms, startEvent, stopEvent);
printf("Kernel time: %.3f ms\n", ms);

hipStreamDestroy(stream0);
hipStreamDestroy(stream1);

/* HIP 에러 처리 — 모든 HIP API 호출에 적용 권장 */
#define HIP_CHECK(call)                                              \
    do {                                                              \
        hipError_t err = (call);                                      \
        if (err != hipSuccess) {                                      \
            fprintf(stderr, "HIP error %s:%d: %s\n",               \
                    __FILE__, __LINE__, hipGetErrorString(err));      \
            exit(1);                                                  \
        }                                                             \
    } while(0)

/* 디바이스 속성 조회 */
hipDeviceProp_t prop;
HIP_CHECK(hipGetDeviceProperties(&prop, 0));
printf("GPU: %s\n", prop.name);
printf("VRAM: %zu MB\n", prop.totalGlobalMem / (1024*1024));
printf("CU count: %d\n", prop.multiProcessorCount);
printf("Max threads/block: %d\n", prop.maxThreadsPerBlock);
printf("Warp size: %d\n", prop.warpSize);  /* AMD: 64 */
printf("L2 cache: %d KB\n", prop.l2CacheSize / 1024);
printf("Memory BW: %.1f GB/s\n",
       2.0 * prop.memoryClockRate * (prop.memoryBusWidth / 8) / 1e6);

hipGraph API

CUDA Graph의 AMD 동등 구현인 hipGraph는 GPU 작업(커널 실행, 메모리 복사, 이벤트 등)을 유향 비순환 그래프(DAG)로 정의하여 반복 실행 비용을 대폭 낮추는 API입니다. 일반 스트림 실행에서는 매 반복마다 CPU가 커널 실행·동기화·메모리 복사 명령을 각각 드라이버에 제출하고, 드라이버는 이를 AQL 패킷으로 변환하여 GPU 큐에 기록합니다. 이 경로의 CPU 오버헤드는 커널 하나당 수 마이크로초 수준이며, 짧은 커널이 수백 개 이어지는 딥러닝 레이어별 실행 구간에서 병목이 됩니다.

hipGraph를 사용하면 첫 번째 반복에서 캡처(capture) 단계를 통해 작업 그래프를 기록하고, 인스턴스화(instantiate)로 실행 가능한 그래프 오브젝트를 생성한 뒤, 이후 반복에서는 단일 hipGraphLaunch() 호출만으로 전체 그래프를 GPU에 제출합니다. CPU 제출 오버헤드가 그래프당 O(1)로 고정되어, 반복 횟수가 많을수록 이득이 커집니다.

스트림 캡처 방식

스트림 캡처는 hipStreamBeginCapture()와 hipStreamEndCapture() 사이에 발생하는 모든 GPU 작업을 자동으로 기록하는 편리한 방법입니다. 기존 스트림 기반 코드를 최소한의 수정으로 그래프로 전환할 수 있습니다. 캡처 모드는 hipStreamCaptureModeGlobal(기본, 가장 안전), hipStreamCaptureModeThreadLocal(스레드별 격리), hipStreamCaptureModeRelaxed(이벤트 혼용 허용) 세 가지가 있으며, 멀티 스트림 캡처 시 hipEventRecord()로 스트림 간 의존 관계를 표현합니다.

// hipGraph 스트림 캡처 예제
#include <hip/hip_runtime.h>

__global__ void kernelA(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) data[idx] *= 2.0f;
}

__global__ void kernelB(float* a, float* b, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) b[idx] = a[idx] + b[idx];
}

int main() {
    const int N = 1 << 20;
    float *dA, *dB, *hA;
    hipMalloc(&dA, N * sizeof(float));
    hipMalloc(&dB, N * sizeof(float));
    hA = new float[N];
    for (int i = 0; i < N; i++) hA[i] = (float)i;

    hipStream_t stream;
    hipStreamCreate(&stream);

    // --- 캡처 단계 (최초 1회) ---
    hipGraph_t graph;
    hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal);

    hipMemcpyAsync(dA, hA, N * sizeof(float),
                   hipMemcpyHostToDevice, stream);
    hipMemsetAsync(dB, 0, N * sizeof(float), stream);

    dim3 block(256), grid((N + 255) / 256);
    kernelA<<<grid, block, 0, stream>>>(dA, N);
    kernelB<<<grid, block, 0, stream>>>(dA, dB, N);

    hipStreamEndCapture(stream, &graph);

    // --- 인스턴스화 ---
    hipGraphExec_t graphExec;
    hipGraphInstantiate(&graphExec, graph,
                        nullptr, nullptr, 0);

    // --- 반복 실행 (각 반복마다 O(1) CPU 오버헤드) ---
    for (int iter = 0; iter < 1000; iter++) {
        hipGraphLaunch(graphExec, stream);
    }
    hipStreamSynchronize(stream);

    hipGraphExecDestroy(graphExec);
    hipGraphDestroy(graph);
    hipStreamDestroy(stream);
    hipFree(dA); hipFree(dB);
    delete[] hA;
}

수동 그래프 구성

스트림 캡처가 불가능한 경우(조건부 분기, 루프 종속 구조 등) 노드와 엣지를 직접 추가하는 수동 구성 방식을 사용합니다. hipGraphCreate()로 빈 그래프를 생성한 뒤 각 노드 타입별 Add 함수로 노드를 추가하고, hipGraphAddDependencies()로 의존 관계를 선언합니다. 수동 구성의 장점은 노드 간 의존 관계를 세밀하게 제어할 수 있고, 스트림 캡처로는 표현하기 어려운 복잡한 DAG 위상을 직접 정의할 수 있다는 점입니다.

// 수동 그래프 구성: GEMM → Bias Add → Activation 파이프라인
hipGraph_t graph;
hipGraphCreate(&graph, 0);

// GEMM 커널 노드 파라미터
hipKernelNodeParams gemmParams = {};
gemmParams.func = (void*)gemmKernel;
gemmParams.gridDim  = dim3(gridX, gridY, 1);
gemmParams.blockDim = dim3(16, 16, 1);
gemmParams.sharedMemBytes = 0;
void* gemmArgs[] = { &dA, &dB, &dC, &M, &N, &K };
gemmParams.kernelParams = gemmArgs;
gemmParams.extra = nullptr;

hipGraphNode_t gemmNode, biasNode, actNode, cpNode;

// 루트 노드 (의존 없음)
hipGraphAddKernelNode(&gemmNode, graph,
                      nullptr, 0, &gemmParams);

// Bias: gemmNode에 의존
hipKernelNodeParams biasParams = {};
biasParams.func = (void*)biasKernel;
biasParams.gridDim = dim3(gridX, 1, 1);
biasParams.blockDim = dim3(256, 1, 1);
void* biasArgs[] = { &dC, &dBias, &N };
biasParams.kernelParams = biasArgs;
hipGraphAddKernelNode(&biasNode, graph,
                      &gemmNode, 1, &biasParams);

// Activation: biasNode에 의존
hipKernelNodeParams actParams = {};
actParams.func = (void*)reluKernel;
actParams.gridDim = dim3(gridX, 1, 1);
actParams.blockDim = dim3(256, 1, 1);
void* actArgs[] = { &dC, &N };
actParams.kernelParams = actArgs;
hipGraphAddKernelNode(&actNode, graph,
                      &biasNode, 1, &actParams);

// Memcpy D→H: actNode 이후
hipMemcpy3DParms cpParams = {};
cpParams.srcPtr = make_hipPitchedPtr(
    dC, M*sizeof(float), M, N);
cpParams.dstPtr = make_hipPitchedPtr(
    hC, M*sizeof(float), M, N);
cpParams.extent = make_hipExtent(
    M*sizeof(float), N, 1);
cpParams.kind = hipMemcpyDeviceToHost;
hipGraphAddMemcpyNode(&cpNode, graph,
                     &actNode, 1, &cpParams);

// 인스턴스화 및 실행
hipGraphExec_t exec;
hipGraphInstantiate(&exec, graph, nullptr, nullptr, 0);
hipGraphLaunch(exec, stream);
hipStreamSynchronize(stream);

그래프 업데이트 및 재사용

ROCm 5.x 이후 hipGraphExecKernelNodeSetParams()와 hipGraphExecUpdate()를 통해 이미 인스턴스화된 그래프의 파라미터(커널 인자, grid/block 크기, memcpy 포인터 등)를 재인스턴스화 없이 수정할 수 있습니다. 학습 루프처럼 입력 데이터 포인터만 바뀌는 경우에 특히 유용하며, 그래프의 위상(노드 수·의존 관계)이 동일하면 재인스턴스화 비용(수 밀리초)을 절약할 수 있습니다.

// 그래프 파라미터 업데이트: 재인스턴스화 없이 커널 인자 변경

// 1) 특정 노드 파라미터만 직접 변경
hipKernelNodeParams newParams = gemmParams;
void* newArgs[] = {
    &dA2, &dB2, &dC2, &M, &N, &K
}; // 새 배치 포인터
newParams.kernelParams = newArgs;
hipGraphExecKernelNodeSetParams(exec, gemmNode, &newParams);

// 2) 그래프 오브젝트 변경 후 Exec에 반영
//    파라미터만 변경이면 SUCCESS, 위상 변경이면 FAIL → 재인스턴스화
hipGraphExecUpdateResult updateResult;
hipGraphNode_t errNode;
hipError_t ret = hipGraphExecUpdate(
    exec, graph, &errNode, &updateResult);
if (updateResult != hipGraphExecUpdateSuccess) {
    hipGraphExecDestroy(exec);
    hipGraphInstantiate(&exec, graph,
                        nullptr, nullptr, 0);
}

// 3) 학습 루프에서의 전형적 사용 패턴
for (int epoch = 0; epoch < epochs; epoch++) {
    for (int step = 0; step < stepsPerEpoch; step++) {
        // 현재 배치로 입력 포인터 업데이트
        updateInputPointers(exec, batchPtr[step]);
        hipGraphLaunch(exec, stream);
        hipStreamSynchronize(stream);
    }
}

// 4) 디버깅: 그래프 DOT 파일로 시각화
hipGraphDebugDotPrint(graph, "graph.dot",
    hipGraphDebugDotFlagsVerbose);

HIP Cooperative Groups

Cooperative Groups는 HIP에서 스레드 그룹을 유연하게 정의하고 동기화하는 API입니다. 기존 __syncthreads()가 블록 내 모든 스레드를 동기화하는 것과 달리, Cooperative Groups는 서브셋 타일, 블록 전체, 그리드 전체 등 다양한 범위의 동기화를 지원합니다. 그리드 레벨 동기화를 위해서는 hipLaunchCooperativeKernel로 커널을 시작해야 하며, GPU가 cooperative launch를 지원해야 합니다.

/* HIP Cooperative Groups: 그리드 레벨 협동 실행 */
#include <hip/hip_cooperative_groups.h>
namespace cg = cooperative_groups;

/* 그리드 전체 동기화를 활용한 reduce 커널 */
__global__ void grid_reduce_kernel(float *data, float *result, int n)
{
    /* 그리드 전체를 하나의 그룹으로 묶음 */
    cg::grid_group grid = cg::this_grid();

    int tid = grid.thread_rank();
    int stride = grid.size();

    /* 1단계: 각 스레드가 담당 구간을 합산 */
    float local_sum = 0.0f;
    for (int i = tid; i < n; i += stride)
        local_sum += data[i];

    /* 2단계: 블록 내 shared memory로 reduce */
    __shared__ float smem[256];
    cg::thread_block block = cg::this_thread_block();
    smem[block.thread_rank()] = local_sum;
    block.sync();   /* __syncthreads()와 동일하지만 명시적 그룹 기반 */

    for (int s = block.size() / 2; s > 0; s >>= 1) {
        if (block.thread_rank() < s)
            smem[block.thread_rank()] += smem[block.thread_rank() + s];
        block.sync();
    }

    /* 3단계: 그리드 전체 동기화 후 블록 결과를 원자적으로 합산 */
    if (block.thread_rank() == 0)
        atomicAdd(result, smem[0]);
    grid.sync();   /* 그리드 레벨 동기화 (모든 블록 완료 대기) */
}

/* Cooperative 커널 시작 방법 */
void launch_cooperative(float *d_data, float *d_result, int n)
{
    int block_size = 256;

    /* 동시 실행 가능한 최대 블록 수 조회 */
    int num_blocks;
    hipOccupancyMaxActiveBlocksPerMultiprocessor(
        &num_blocks, grid_reduce_kernel, block_size, 0);

    int device;
    hipGetDevice(&device);
    hipDeviceProp_t props;
    hipGetDeviceProperties(&props, device);
    int grid_size = num_blocks * props.multiProcessorCount;

    /* Cooperative 커널 시작: 일반 <<<...>>> 구문 대신 */
    void *args[] = { &d_data, &d_result, &n };
    hipLaunchCooperativeKernel(
        (void *)grid_reduce_kernel,
        dim3(grid_size), dim3(block_size),
        args, 0, 0
    );
}

/* Thread Block Tile: wavefront 내 sub-group 동기화 */

__global__ void warp_reduce_kernel(float *data, float *out)
{
    cg::thread_block block = cg::this_thread_block();

    /* 32개 스레드 타일 생성 (warp/wavefront 서브셋) */
    auto tile32 = cg::tiled_partition<32>(block);

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[tid];

    /* 타일 내 butterfly reduce (lane XOR shuffle) */
    for (int offset = tile32.size() / 2; offset > 0; offset >>= 1)
        val += tile32.shfl_down(val, offset);

    /* tile의 lane 0이 블록 결과를 출력 */
    if (tile32.thread_rank() == 0)
        atomicAdd(out, val);

    /* 16개 타일도 지원: tiled_partition<16> */
    auto tile16 = cg::tiled_partition<16>(block);
    /* AMD에서는 wavefront 64개 스레드를 16개 x 4 타일로 분할 */
}

/* coalesced_threads: 조건부 활성 스레드 그룹 */
__global__ void conditional_sync_kernel(int *data, int *result, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    cg::thread_block block = cg::this_thread_block();

    if (tid < n && data[tid] > 0) {
        /* 조건을 만족하는 스레드만으로 그룹 구성 */
        auto active = cg::coalesced_threads();
        int sum = reduce(active, data[tid], cg::plus<int>());
        if (active.thread_rank() == 0)
            atomicAdd(result, sum);
    }
}

HIP Virtual Memory Management (VMM)

HIP VMM API는 CUDA Virtual Memory Management와 유사하게, 물리적 메모리 할당과 가상 주소 공간 매핑을 분리하는 세밀한 제어 기능을 제공합니다. 이를 통해 과도한 예약(overcommit), 동적 크기 조정, 여러 GPU에 동일한 물리 메모리 매핑 등 고급 메모리 관리 기법을 구현할 수 있습니다.

/* HIP VMM: 가상 주소 공간 예약 → 물리 메모리 할당 → 매핑 → 접근 권한 설정 */

#include <hip/hip_runtime.h>

void hip_vmm_example(size_t size)
{
    hipDeviceptr_t va_ptr;
    hipMemGenericAllocationHandle_t handle;

    /* 1단계: 가상 주소 공간 예약 (물리 메모리 소비 없음) */
    hipMemAddressReserve(&va_ptr, size,
                         2 * 1024 * 1024,  /* 2 MB 정렬 */
                         0,                 /* 주소 힌트 (0=자동) */
                         0);                /* flags */

    /* 2단계: 물리 메모리 할당 속성 설정 */
    hipMemAllocationProp prop = {};
    prop.type = hipMemAllocationTypePinned;
    prop.location.type = hipMemLocationTypeDevice;
    prop.location.id = 0;   /* GPU 0의 VRAM */

    /* 3단계: 물리 메모리 할당 (실제 VRAM 소비) */
    hipMemCreate(&handle, size, &prop, 0);

    /* 4단계: 가상 주소에 물리 메모리 매핑 */
    hipMemMap(va_ptr, size,
              0,      /* offset within physical allocation */
              handle,
              0);     /* flags */

    /* 5단계: 접근 권한 설정 (읽기/쓰기 활성화) */
    hipMemAccessDesc access_desc = {};
    access_desc.location.type = hipMemLocationTypeDevice;
    access_desc.location.id = 0;
    access_desc.flags = hipMemAccessFlagsProtReadWrite;
    hipMemSetAccess(va_ptr, size, &access_desc, 1);

    /* 이제 va_ptr를 GPU 커널에서 일반 포인터처럼 사용 가능 */
    /* ... GPU 커널 실행 ... */

    /* 해제 순서: 매핑 해제 → 가상 주소 해제 → 물리 메모리 해제 */
    hipMemUnmap(va_ptr, size);
    hipMemAddressFree(va_ptr, size);
    hipMemRelease(handle);
}

/* 동적 크기 조정 예: realloc 패턴 */
void vmm_dynamic_resize(hipDeviceptr_t *va_ptr, size_t old_size, size_t new_size)
{
    hipDeviceptr_t new_va;
    hipMemAllocationProp prop = {};
    prop.type = hipMemAllocationTypePinned;
    prop.location.type = hipMemLocationTypeDevice;
    prop.location.id = 0;

    /* 새 크기의 가상 주소 공간 예약 */
    hipMemAddressReserve(&new_va, new_size, 0, 0, 0);

    /* 기존 데이터를 새 영역에 매핑 (데이터 복사 없이 재매핑) */
    hipMemGenericAllocationHandle_t new_handle;
    hipMemCreate(&new_handle, new_size, &prop, 0);
    hipMemMap(new_va, new_size, 0, new_handle, 0);

    hipMemAccessDesc desc = {};
    desc.location.type = hipMemLocationTypeDevice;
    desc.location.id = 0;
    desc.flags = hipMemAccessFlagsProtReadWrite;
    hipMemSetAccess(new_va, new_size, &desc, 1);

    *va_ptr = new_va;
}

hipify 마이그레이션

기존 CUDA 코드를 HIP로 전환할 때는 hipify-perl과 hipify-clang 두 가지 도구를 사용합니다. hipify-perl은 단순 텍스트 치환 방식으로 빠르고 의존성이 없지만 복잡한 코드에서 오변환이 생길 수 있습니다. hipify-clang은 Clang AST를 기반으로 의미론적으로 정확하게 변환하며 통계 리포트도 제공합니다.

# hipify-perl: 단순 텍스트 치환
hipify-perl cuda_vector_add.cu > hip_vector_add.hip

# 디렉터리 전체 변환
hipify-perl --inplace src/*.cu src/*.cuh

# 변환 통계 확인
hipify-perl --print-stats myapp.cu

# 출력 예:
# [HIPIFY] info: converted 42 CUDA API calls
# [HIPIFY] info: 3 CUDA API call(s) not converted (manual action required)

# hipify-clang: AST 기반 변환 (더 정확)
# compilation database (compile_commands.json) 필요
cmake -DCMAKE_EXPORT_COMPILE_COMMANDS=ON ..

hipify-clang cuda_app.cu \
    -p ./build \
    --cuda-path=/usr/local/cuda \
    -o hip_app.cpp

# 변환 통계 리포트 생성
hipify-clang --print-stats=1 cuda_app.cu -p ./build

# 여러 파일 일괄 변환
find src -name "*.cu" | xargs -I{} hipify-clang {} -p ./build

# CMakeLists.txt: HIP 통합 빌드 설정
cmake_minimum_required(VERSION 3.21)
project(MyROCmApp LANGUAGES CXX HIP)

# ROCm 경로 설정
find_package(HIP REQUIRED)
find_package(rocblas REQUIRED)

set(CMAKE_HIP_ARCHITECTURES "gfx942;gfx1100")  # 타겟 아키텍처

add_executable(myapp main.hip vector_add.hip)

target_link_libraries(myapp
    hip::host          # HIP 런타임
    roc::rocblas       # rocBLAS
)

set_source_files_properties(main.hip vector_add.hip
    PROPERTIES LANGUAGE HIP)

# 디버그 빌드 시 sanitizer 활성화
target_compile_options(myapp PRIVATE
    "$<$<CONFIG:Debug>:-fsanitize=address>")

RDNA vs CDNA 아키텍처

AMD GPU는 용도에 따라 두 가지 주요 아키텍처 계열로 나뉩니다. RDNA(Radeon DNA)는 게임과 컨슈머 그래픽을 위한 아키텍처로, 디스플레이 엔진(DCN), 레이 트레이싱 유닛(BVH), 미디어 인코더(VCN)를 포함합니다. CDNA(Compute DNA)는 HPC와 AI 연산을 위한 서버 전용 아키텍처로, 그래픽 관련 유닛을 제거하고 더 많은 CU(Compute Unit), 대용량 HBM 메모리, XGMI 인터커넥트, 행렬 연산 가속기(Matrix Core)에 집중합니다.

# rocm-smi로 GPU 토폴로지 및 아키텍처 확인
rocm-smi --showtopo
# GPU[0]  GPU[1]  ... GPU[7]
# GPU[0]:  0       XGMI    XGMI  ...
# GPU[1]:  XGMI    0       XGMI  ...

# 개별 GPU 정보 조회
rocm-smi -d 0 --showproductname --showmeminfo vram
rocm-smi -d 0 --showclkfreq       # 현재 클럭 주파수
rocm-smi -d 0 --showtemp           # GPU 온도
rocm-smi -d 0 --showpower          # 전력 소비

# HIP으로 아키텍처 버전 조회
rocminfo | grep -E "Name|gfx_target_triple"
# 출력 예: gfx942 (MI300X의 GFX IP)

/* HIP으로 아키텍처 정보 프로그래밍 방식 조회 */
#include <hip/hip_runtime.h>
#include <hip/hip_ext.h>

hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, 0);

printf("Name: %s\n", prop.name);
printf("gcnArchName: %s\n", prop.gcnArchName);     /* gfx942 등 */
printf("CU count: %d\n", prop.multiProcessorCount);
printf("VRAM: %zu GB\n", prop.totalGlobalMem >> 30);
printf("L2 cache: %d MB\n", prop.l2CacheSize >> 20);
printf("Max shmem/block: %zu KB\n", prop.sharedMemPerBlock >> 10);
printf("XNACK: %d\n", prop.xnackEnabled);           /* SVM 폴트 지원 */
printf("Unified addr: %d\n", prop.unifiedAddressing);

Matrix Core / MFMA 프로그래밍 심화

Matrix Core(MFMA, Matrix Fused Multiply-Add)는 CDNA 아키텍처(MI100 이후)의 핵심 AI 가속 엔진입니다. 단일 명령어로 소형 행렬 곱셈-덧셈(D = A × B + C)을 실행하며, FP64/FP32/FP16/BF16/INT8/FP8 등 다양한 정밀도를 지원합니다. RDNA3(GFX11)에서는 WMMA(Wavefront Matrix Multiply-Accumulate)라는 유사하지만 범위가 더 제한된 형태로 제공됩니다. 프로그래밍 방법은 rocWMMA 라이브러리를 통한 고수준 API와, LLVM 내장 함수(intrinsic)를 직접 호출하는 저수준 방법이 있습니다.

/* rocWMMA를 사용한 MFMA 행렬 곱: C = A × B + C */
/* rocWMMA는 CUDA wmma API와 유사한 고수준 인터페이스 제공 */
#include <rocwmma/rocwmma.hpp>

using namespace rocwmma;

constexpr int WMMA_M = 16;
constexpr int WMMA_N = 16;
constexpr int WMMA_K = 16;

__global__ void rocwmma_gemm_kernel(
    half *A, half *B, float *C, float *D,
    int M, int N, int K)
{
    /* 블록 내 wavefront 위치 계산 */
    int wave_row = (blockIdx.y * blockDim.y + threadIdx.y) / warpSize;
    int wave_col = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize;

    /* fragment 선언: 16×16 행렬 타일 */
    fragment<matrix_a, WMMA_M, WMMA_N, WMMA_K, half, row_major> a_frag;
    fragment<matrix_b, WMMA_M, WMMA_N, WMMA_K, half, col_major> b_frag;
    fragment<accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
    fragment<accumulator, WMMA_M, WMMA_N, WMMA_K, float> d_frag;

    /* accumulator 초기화 */
    fill_fragment(d_frag, 0.0f);

    /* K 차원을 따라 WMMA_K 크기 타일로 반복 */
    for (int k = 0; k < K; k += WMMA_K) {
        int a_row = wave_row * WMMA_M;
        int a_col = k;
        int b_row = k;
        int b_col = wave_col * WMMA_N;

        /* 글로벌 메모리에서 fragment 로드 */
        load_matrix_sync(a_frag, A + a_row * K + a_col, K);
        load_matrix_sync(b_frag, B + b_row * N + b_col, N);

        /* MFMA 실행: d_frag += a_frag × b_frag */
        mma_sync(d_frag, a_frag, b_frag, d_frag);
    }

    /* C 행렬 로드 및 누산 */
    load_matrix_sync(c_frag, C + wave_row * WMMA_M * N + wave_col * WMMA_N, N,
                     mem_row_major);
    for (int i = 0; i < d_frag.num_elements; i++)
        d_frag.x[i] += c_frag.x[i];

    /* 결과를 글로벌 메모리에 저장 */
    store_matrix_sync(D + wave_row * WMMA_M * N + wave_col * WMMA_N,
                      d_frag, N, mem_row_major);
}

/* 저수준 MFMA intrinsic 직접 사용 */
/* __builtin_amdgcn_mfma_f32_16x16x16f16 */

__global__ void mfma_intrinsic_kernel(half *A, half *B, float *C)
{
    /* MFMA f32 16x16x16 f16 연산 */
    /* 각 wavefront(64 스레드)가 16x16 타일을 공동으로 처리 */

    /* 입력: A는 half×4 벡터 (각 스레드가 4개의 half 값 보유) */
    typedef float __attribute__((ext_vector_type(4))) float4v;
    typedef __fp16 __attribute__((ext_vector_type(4))) half4v;

    int lane = threadIdx.x % 64;  /* wavefront 내 lane 번호 */
    int row = lane / 4;           /* 담당 행 인덱스 */
    int col_base = (lane % 4) * 4; /* 담당 열 시작 인덱스 */

    /* A 행렬의 해당 스레드 담당 원소 로드 */
    half4v a_vec;
    a_vec[0] = A[row * 16 + col_base];
    a_vec[1] = A[row * 16 + col_base + 1];
    a_vec[2] = A[row * 16 + col_base + 2];
    a_vec[3] = A[row * 16 + col_base + 3];

    half4v b_vec;  /* B 행렬 원소 (유사하게 로드) */

    /* accumulator 초기화 (4개 float, 각 스레드가 accumulator 일부 소유) */
    float4v c_vec = {0.0f, 0.0f, 0.0f, 0.0f};

    /* MFMA 내장 함수 호출: D[16x16][float] = A[16x16][f16] × B[16x16][f16] + C */
    float4v result = __builtin_amdgcn_mfma_f32_16x16x16f16(
        a_vec,   /* A 타일 (각 스레드 담당 부분) */
        b_vec,   /* B 타일 */
        c_vec,   /* C accumulator (초기값) */
        0,       /* cbsz: broadcast 제어 */
        0,       /* abid: A 행렬 broadcast ID */
        0        /* blgp: B 행렬 layout 그룹 */
    );

    /* 결과 저장: 각 스레드가 accumulator 일부를 글로벌 메모리에 기록 */
    C[row * 16 + (lane % 4) * 4 + 0] = result[0];
    C[row * 16 + (lane % 4) * 4 + 1] = result[1];
    C[row * 16 + (lane % 4) * 4 + 2] = result[2];
    C[row * 16 + (lane % 4) * 4 + 3] = result[3];
}

메모리 관리

ROCm GPU 프로그래밍에서 메모리는 성능의 핵심 결정 요소입니다. GPU는 여러 계층의 메모리를 제공하며 각 계층은 용량, 대역폭, 레이턴시, 가시성 범위가 서로 다릅니다. 최적의 성능을 얻으려면 알고리즘의 데이터 접근 패턴에 맞는 메모리 계층을 선택하고, 데이터 지역성을 최대화하며, 메모리 계층 간 전송을 최소화해야 합니다.

/* 기본 GPU 메모리 할당 및 복사 */
#include <hip/hip_runtime.h>

int N = 1024 * 1024;
float *h_data = new float[N];  /* 호스트 메모리 */
float *d_data;                   /* 디바이스 (GPU VRAM) 포인터 */

/* VRAM 할당 */
hipMalloc(&d_data, N * sizeof(float));

/* 호스트 → GPU 복사 (PCIe DMA) */
hipMemcpy(d_data, h_data, N * sizeof(float), hipMemcpyHostToDevice);

/* GPU 내부 복사 (SDMA 엔진 사용) */
float *d_dst;
hipMalloc(&d_dst, N * sizeof(float));
hipMemcpy(d_dst, d_data, N * sizeof(float), hipMemcpyDeviceToDevice);

/* 페이지 고정(Pinned) 호스트 메모리: PCIe DMA 성능 향상 */
float *h_pinned;
hipHostMalloc(&h_pinned, N * sizeof(float), hipHostMallocDefault);
/* 해제 */
hipHostFree(h_pinned);

/* 메모리 0 초기화 */
hipMemset(d_data, 0, N * sizeof(float));

hipFree(d_data);

/* Unified(Managed) Memory: CPU-GPU 자동 마이그레이션 */
float *um_data;
hipMallocManaged(&um_data, N * sizeof(float));

/* CPU에서 초기화 */
for (int i = 0; i < N; ++i) um_data[i] = (float)i;

/* GPU 커널 실행 — 런타임이 자동으로 GPU로 마이그레이션 */
myKernel<<<grid, block>>>(um_data, N);
hipDeviceSynchronize();

/* CPU에서 다시 접근 — 런타임이 자동으로 CPU로 마이그레이션 */
printf("result[0] = %f\n", um_data[0]);

/* 마이그레이션 힌트: prefetch로 성능 개선 */
hipMemPrefetchAsync(um_data, N * sizeof(float),
                     0,  /* GPU device id */
                     0); /* stream */

hipFree(um_data);

/* 메모리 풀 (hipMallocFromPoolAsync): ROCm 5.x+ */
hipMemPool_t pool;
hipMemPoolProps props = {};
props.allocType       = hipMemAllocationTypePinned;
props.location.type   = hipMemLocationTypeDevice;
props.location.id     = 0;
hipMemPoolCreate(&pool, &props);

/* 스트림에서 풀 메모리 비동기 할당 */
float *d_buf;
hipStream_t stream;
hipStreamCreate(&stream);
hipMallocFromPoolAsync(&d_buf, N * sizeof(float), pool, stream);

/* 사용 ... */
myKernel<<<grid, block, 0, stream>>>(d_buf, N);

/* 스트림 완료 후 풀에 반환 (즉시 OS에 반환하지 않음) */
hipFreeAsync(d_buf, stream);
hipStreamSynchronize(stream);
hipMemPoolDestroy(pool);

SVM (Shared Virtual Memory) 심층 분석

AMD GPU의 SVM(Shared Virtual Memory)은 CPU와 GPU가 동일한 가상 주소 공간을 공유하여 포인터를 별도 복사 없이 두 장치 사이에서 직접 교환할 수 있게 하는 기술입니다. ROCm의 SVM 구현은 HSA 사양을 따르며 KFD(Kernel Fusion Driver)가 중심 역할을 합니다. hipMallocManaged()로 할당된 메모리는 처음에 CPU 또는 GPU 중 어느 쪽에도 물리적으로 고정되지 않고, 실제 접근이 발생하는 시점에 페이지 폴트 메커니즘을 통해 해당 장치의 메모리로 마이그레이션됩니다.

AMD GPU의 SVM은 IOMMU(Input-Output Memory Management Unit)와 깊게 연계됩니다. GPU가 아직 매핑되지 않은 페이지에 접근하면 IOMMU가 페이지 폴트를 발생시키고, KFD의 폴트 핸들러가 이를 처리합니다. 이 기능을 XNACK(eXtended Not Acknowledge)라 부르며, GPU가 메모리 접근 실패를 재시도할 수 있는 하드웨어 지원을 뜻합니다. XNACK 활성화 여부는 GFX9(Vega) 이후 아키텍처에서 지원하며, 운영 환경(bare-metal vs MES)과 커널 빌드 플래그에 따라 다릅니다.

SVM 페이지 폴트 메커니즘

XNACK 활성화 환경에서 GPU가 미매핑 페이지에 접근하면 다음 순서로 처리됩니다. (1) GPU 메모리 컨트롤러가 IOMMU로 XNACK 신호를 보냅니다. (2) IOMMU가 FAULT 인터럽트를 CPU에 전달합니다. (3) KFD의 kfd_iommu_device_init() 경로에서 등록된 폴트 핸들러가 호출됩니다. (4) 핸들러는 해당 가상 주소가 속한 SVM 범위를 조회하고 hipMemAdvise로 설정된 힌트를 참조합니다. (5) migrate_vma_pages() 또는 DMA를 통해 물리 페이지를 GPU VRAM으로 복사합니다. (6) GPU 페이지 테이블과 IOMMU TLB를 갱신합니다. (7) GPU가 메모리 접근을 재시도합니다.

// SVM hipMemAdvise 패턴 예제
#include <hip/hip_runtime.h>

void* ptr;
size_t size = 256 * 1024 * 1024;  // 256 MiB
hipMallocManaged(&ptr, size);

int gpuDevice = 0;

// 패턴 1: GPU 우선 배치 (GPU가 주로 읽기/쓰기)
hipMemAdvise(ptr, size,
    hipMemAdviseSetPreferredLocation, gpuDevice);
// 향후 GPU 접근 시 마이그레이션 없이 GPU VRAM에 배치 유지

// 패턴 2: CPU도 접근하지만 읽기 전용 (미러 페이지 생성)
hipMemAdvise(ptr, size,
    hipMemAdviseSetReadMostly, hipCpuDeviceId);
// CPU/GPU 모두 로컬 복사본 유지, 쓰기 시 미러 무효화

// 패턴 3: GPU가 매핑만 유지 (마이그레이션 없이 직접 접근)
hipMemAdvise(ptr, size,
    hipMemAdviseSetAccessedBy, gpuDevice);
// CPU DRAM에 두되 GPU 페이지 테이블에 직접 매핑
// PCIe/XGMI를 통한 원격 접근: 대역폭 낮지만 용량 활용 가능

// 패턴 4: GPU 실행 전 선제 프리페치 (폴트 방지)
hipStream_t stream;
hipStreamCreate(&stream);
hipMemPrefetchAsync(ptr, size, gpuDevice, stream);
// 비동기 DMA: 커널 실행 전에 GPU VRAM으로 이동 보장

// 사용 후 CPU로 복귀 (결과 수집)
hipMemPrefetchAsync(ptr, size, hipCpuDeviceId, stream);
hipStreamSynchronize(stream);
// 이제 CPU에서 ptr 직접 접근 가능

hipFree(ptr);

SVM 속성 타입 및 조회

hipMemRangeGetAttribute()를 통해 특정 메모리 범위의 현재 SVM 속성을 조회할 수 있습니다. 이는 런타임에 마이그레이션 결정 근거를 확인하거나 디버깅에 활용됩니다. KFD에서 관리하는 SVM 범위 정보는 /sys/kernel/debug/kfd/를 통해서도 부분적으로 확인할 수 있습니다.

// SVM 속성 조회 및 페이지 마이그레이션 상태 확인

// 현재 preferred location 조회
int location = -1;
hipMemRangeGetAttribute(
    &location, sizeof(int),
    hipMemRangeAttributePreferredLocation,
    ptr, size);
// location == hipCpuDeviceId(=-1) 또는 GPU device ID

// 페이지가 실제로 어디 있는지 확인 (디바이스 ID 배열 반환)
int locations[256];  // 페이지 수만큼
hipMemRangeGetAttribute(
    locations, size / (4096),  // 4KB 페이지 단위
    hipMemRangeAttributeAccessedBy,
    ptr, size);

// READ_MOSTLY 미러 페이지 생성 상태
int readMostly = 0;
hipMemRangeGetAttribute(
    &readMostly, sizeof(int),
    hipMemRangeAttributeReadMostly,
    ptr, size);

// KFD SVM ioctl을 통한 범위 속성 직접 설정 (고급)
#include <linux/kfd_ioctl.h>

struct kfd_ioctl_svm_args args = {};
args.start_addr = (uint64_t)ptr;
args.size        = size;
args.op          = KFD_IOCTL_SVM_OP_SET_ATTR;
args.nattr       = 2;

struct kfd_ioctl_svm_attribute attrs[2];
attrs[0].type  = KFD_IOCTL_SVM_ATTR_PREFERRED_LOC;
attrs[0].value = KFD_IOCTL_SVM_LOCATION_GPU(0);
attrs[1].type  = KFD_IOCTL_SVM_ATTR_SET_FLAGS;
attrs[1].value = KFD_IOCTL_SVM_FLAG_GPU_EXEC;
args.attrs = (uint64_t)attrs;

ioctl(kfd_fd, AMDKFD_IOC_SVM, &args);

XNACK 모드 설정과 성능 영향

XNACK 모드는 GPU 하드웨어 수준에서 활성화되며, AMD GPU 아키텍처별 지원 여부가 다릅니다. GFX9(Vega20/MI100) 이상에서 지원되며 RDNA(RX 5000-7000) 시리즈에서는 지원되지 않습니다. XNACK OFF 시 모든 SVM 메모리는 커널 실행 전에 사전 매핑되어야 하며, 이 경우 hipMallocManaged()는 실질적으로 항상 GPU VRAM 또는 시스템 메모리 중 하나에 고정(pinned) 방식으로 동작합니다.

## XNACK 모드 확인 및 설정

# 현재 XNACK 모드 확인
cat /sys/class/kfd/kfd/topology/nodes/*/properties | grep xnack
# xnack_enabled 1 → XNACK 지원 및 활성화

# 환경변수로 XNACK 활성화 (런타임 제어)
export HSA_XNACK=1
# GFX9 이상에서만 유효; RDNA에서는 무시됨

# 커널 부트 파라미터로 XNACK 강제 설정
# /etc/default/grub에서 GRUB_CMDLINE_LINUX에 추가:
# amdgpu.noretry=0   (XNACK 재시도 활성화)

# XNACK OFF 성능 프로파일: 폴트 핸들러 없음
#   장점: 커널 실행 중 예측 가능한 지연시간
#   단점: GPU VRAM 용량 초과 시 OOM 오류

# XNACK ON 성능 프로파일
#   장점: Oversubscription 가능 (VRAM > 물리 메모리처럼 사용)
#   단점: 첫 접근 페이지 폴트 → 마이그레이션 레이턴시 수십~수백 μs

# GPU 페이지 폴트 통계 확인
cat /sys/kernel/debug/kfd/proc/*/vm_fault_info 2>/dev/null

GPU 스케줄링과 컴퓨트 큐

ROCm의 GPU 작업 제출 경로는 사용자 공간에서 AQL 패킷을 큐 링 버퍼에 직접 쓰는 방식으로 동작합니다. 커널 경유 없이 사용자 공간에서 GPU 커맨드를 제출할 수 있으므로 제출 레이턴시가 매우 낮습니다. 이 설계는 HSA(Heterogeneous System Architecture)의 핵심 원칙 중 하나입니다.

/* 여러 스트림으로 병렬 GPU 작업 실행 */
const int NUM_STREAMS = 4;
hipStream_t streams[NUM_STREAMS];

for (int i = 0; i < NUM_STREAMS; ++i) {
    hipStreamCreateWithFlags(&streams[i],
        hipStreamNonBlocking);  /* 기본 스트림과 독립 */
}

/* 각 스트림에 독립적 작업 제출 */
for (int i = 0; i < NUM_STREAMS; ++i) {
    int offset = i * (N / NUM_STREAMS);
    hipMemcpyAsync(d_data + offset, h_data + offset,
                   (N / NUM_STREAMS) * sizeof(float),
                   hipMemcpyHostToDevice, streams[i]);
    myKernel<<<grid/NUM_STREAMS, block, 0, streams[i]>>>(
        d_data + offset, N / NUM_STREAMS);
}

/* 모든 스트림 완료 대기 */
for (int i = 0; i < NUM_STREAMS; ++i)
    hipStreamSynchronize(streams[i]);

/* 큐 우선순위 설정 (ROCm 확장) */
hipStream_t high_prio_stream, low_prio_stream;

/* 우선순위 범위 조회 */
int prio_low, prio_high;
hipDeviceGetStreamPriorityRange(&prio_low, &prio_high);
/* 일반적으로 prio_high=-1, prio_low=0 */

hipStreamCreateWithPriority(&high_prio_stream,
    hipStreamNonBlocking, prio_high);
hipStreamCreateWithPriority(&low_prio_stream,
    hipStreamNonBlocking, prio_low);

/* 우선순위 높은 스트림에 중요 작업 제출 */
criticalKernel<<<grid, block, 0, high_prio_stream>>>(d_data);
backgroundKernel<<<grid, block, 0, low_prio_stream>>>(d_buf);

ROCm 라이브러리 생태계

ROCm은 수치 연산, 딥러닝, 신호 처리, 희소 행렬 등 다양한 도메인을 위한 고도로 최적화된 라이브러리를 제공합니다. 이 라이브러리들은 AMD GPU의 Matrix Core, LDS 계층, 메모리 대역폭을 최대한 활용하도록 수작업 튜닝된 커널을 포함합니다. 대부분의 경우 직접 커널을 작성하기보다 이 라이브러리를 활용하는 것이 더 높은 성능을 얻을 수 있습니다.

/* rocBLAS: CUDA cuBLAS 호환 GEMM 예제 */
#include <rocblas/rocblas.h>

rocblas_handle handle;
rocblas_create_handle(&handle);

/* C = alpha * A * B + beta * C  (행렬 곱) */
const int M = 4096, N = 4096, K = 4096;
const float alpha = 1.0f, beta = 0.0f;

/* FP32 GEMM */
rocblas_sgemm(handle,
    rocblas_operation_none, rocblas_operation_none,
    M, N, K,
    &alpha,
    d_A, M,   /* A: M×K */
    d_B, K,   /* B: K×N */
    &beta,
    d_C, M);  /* C: M×N */

/* FP16 GEMM (Tensor Core / Matrix Core 활용) */
rocblas_hgemm(handle,
    rocblas_operation_none, rocblas_operation_none,
    M, N, K,
    &alpha_h16, d_A_h16, M, d_B_h16, K,
    &beta_h16,  d_C_h16, M);

rocblas_destroy_handle(handle);

/* rocFFT: FFT 실행 예제 */
#include <rocfft/rocfft.h>

rocfft_setup();  /* 라이브러리 초기화 */

rocfft_plan plan;
size_t lengths[1] = { 1024 };  /* 1D FFT 크기 */

rocfft_plan_create(&plan,
    rocfft_placement_inplace,      /* in-place FFT */
    rocfft_transform_type_complex_forward,
    rocfft_precision_single,       /* FP32 복소수 */
    1,      /* 차원 */
    lengths,
    1,      /* 배치 수 */
    nullptr);

/* 작업 버퍼 크기 조회 */
rocfft_plan_description desc;
size_t work_buf_size;
rocfft_plan_get_work_buffer_size(plan, &work_buf_size);

void *work_buf = nullptr;
if (work_buf_size) hipMalloc(&work_buf, work_buf_size);

/* FFT 실행 */
rocfft_execution_info info;
rocfft_execution_info_create(&info);
rocfft_execution_info_set_work_buffer(info, work_buf, work_buf_size);

void *in_out[] = { d_complex_data };
rocfft_execute(plan, in_out, nullptr, info);

rocfft_plan_destroy(plan);
rocfft_teardown();

/* MIOpen: 딥러닝 합성곱 예제 (cuDNN 호환) */
#include <miopen/miopen.h>

miopenHandle_t handle;
miopenCreate(&handle);

/* 텐서 디스크립터 생성 */
miopenTensorDescriptor_t inputDesc, filterDesc, outputDesc;
miopenCreateTensorDescriptor(&inputDesc);
miopenSet4dTensorDescriptor(inputDesc, miopenFloat,
    1, 64, 224, 224);  /* N, C, H, W */

miopenCreateTensorDescriptor(&filterDesc);
miopenSet4dTensorDescriptor(filterDesc, miopenFloat,
    64, 64, 3, 3);     /* K, C, R, S */

/* 합성곱 디스크립터 */
miopenConvolutionDescriptor_t convDesc;
miopenCreateConvolutionDescriptor(&convDesc);
miopenInitConvolutionDescriptor(convDesc,
    miopenConvolution, 1, 1, 1, 1, 1, 1);

/* 출력 크기 계산 → 합성곱 실행 */
int n, c, h, w;
miopenGetConvolutionForwardOutputDim(convDesc, inputDesc,
    filterDesc, &n, &c, &h, &w);

miopenConvolutionForward(handle, &alpha,
    inputDesc, d_input, filterDesc, d_filter,
    convDesc, algo, &beta, outputDesc, d_output,
    workspace, workspace_size);

miopenDestroy(handle);

Composable Kernel (CK) 라이브러리

AMD의 Composable Kernel(CK)은 C++ 템플릿 메타프로그래밍을 기반으로 한 GPU 커널 빌딩 블록 라이브러리입니다. rocBLAS, MIOpen, Flash Attention 등 AMD의 주요 고성능 라이브러리들이 내부적으로 CK를 사용합니다. CK의 핵심 가치는 타일(tile) 수준의 추상화입니다. 각 연산을 타일 단위로 분해하고, 타일별 크기·레이아웃·로딩 방식을 컴파일 타임에 결정하여 레지스터 및 LDS(Local Data Share) 활용을 최대화합니다.

CK는 장치별 마이크로커널(예: Matrix Core MFMA 명령어 래핑)을 직접 구현하며, 이를 Tensor Operation 레이어에서 GEMM, Convolution, Softmax, LayerNorm, Attention 등으로 조합합니다. 최종 사용자는 Device API 레이어만 호출하면 되고, 내부의 타일 스케줄링과 메모리 파이프라인은 CK가 자동으로 처리합니다.

CK 아키텍처와 타일 프로그래밍

CK 타일 프로그래밍 모델의 핵심은 TensorDescriptor입니다. 다차원 텐서의 논리적 형태(shape)와 물리적 메모리 레이아웃(stride, padding, transpose)을 컴파일 타임 타입으로 인코딩합니다. 이를 통해 런타임 분기 없이 다양한 행렬 레이아웃(RowMajor/ColumnMajor)을 동일한 커널 코드로 처리합니다. BlockwiseGemmXdlops는 하나의 워크그룹이 담당하는 출력 타일을 MFMA 명령어로 계산하는 핵심 빌딩 블록이며, M/N/K 방향의 타일 크기를 템플릿 파라미터로 지정합니다.

// CK GEMM 타일 크기 구성 예제
// C = A * B  (M×K) × (K×N) = M×N
#include <ck/tensor_operation/gpu/device/device_gemm_xdl.hpp>
#include <ck/tensor_operation/gpu/element/element_wise_operation.hpp>

using namespace ck;
using namespace ck::tensor_operation;
using namespace ck::tensor_operation::device;

// 타입 정의
using ADataType    = ck::half_t;
using BDataType    = ck::half_t;
using CDataType    = ck::half_t;
using AccDataType  = float;

// 레이아웃 지정 (행우선/열우선)
using ALayout = Row;
using BLayout = Col;  // B 전치
using CLayout = Row;

// 요소별 연산 (에필로그)
using AElementOp = element_wise::PassThrough;
using BElementOp = element_wise::PassThrough;
using CEElementOp = element_wise::PassThrough;

// GEMM Device 인스턴스 (타일 크기 템플릿 파라미터)
using DeviceGemmInstance = DeviceGemmXdl<
    ALayout, BLayout, CLayout,
    ADataType, BDataType, CDataType, AccDataType,
    AElementOp, BElementOp, CEElementOp,
    256,   // BlockSize: 워크그룹 스레드 수
    256,   // MPerBlock: M 방향 타일 크기
    128,   // NPerBlock: N 방향 타일 크기
    32,    // KPerBlock: K 방향 타일 크기 (LDS 사용)
    8,     // AK1: A 행렬 벡터화 단위
    8,     // BK1: B 행렬 벡터화 단위
    32,    // MPerXDL: MFMA M 타일
    32,    // NPerXDL: MFMA N 타일
    1, 2,  // MXdlPerWave, NXdlPerWave
    S<4,64,1>,  // ABlockTransfer ThreadCluster
    S<4,64,1>,  // BBlockTransfer ThreadCluster
    1, 1, 1   // CBlockTransfer 파라미터
>;

int run_ck_gemm(
    int M, int N, int K,
    const ADataType* dA,
    const BDataType* dB,
    CDataType* dC)
{
    auto gemm = DeviceGemmInstance{};
    auto invoker  = gemm.MakeInvoker();
    auto argument = gemm.MakeArgument(
        dA, dB, dC,
        M, N, K,
        K,   // StrideA
        K,   // StrideB (B가 전치이므로 K)
        N,   // StrideC
        AElementOp{}, BElementOp{}, CEElementOp{});

    if (!gemm.IsSupportedArgument(argument)) {
        return -1;  // 타일 크기가 M/N/K로 나누어지지 않음
    }
    invoker.Run(argument, StreamConfig{nullptr, false});
    return 0;
}

CK GEMM 성능 튜닝

CK 기반 GEMM은 MPerBlock, NPerBlock, KPerBlock, BlockSize 네 가지 타일 파라미터가 성능에 가장 큰 영향을 미칩니다. MI300X(CDNA3)에서는 256×128×32 타일 구성이 FP16 GEMM에서 가장 넓은 범위의 행렬 크기에 대해 최고 성능을 내는 경향이 있습니다. CK는 PerfConfig 열거로 주어진 행렬 크기에 대해 최적 구성을 자동 선택하는 커널 선택(kernel selection) 메커니즘도 제공합니다.

// CK 최적 커널 자동 선택 (kernel selection API)
#include <ck/tensor_operation/gpu/device/device_gemm_multiple_d_xdl_cshuffle.hpp>

// 컴파일 타임에 등록된 모든 구성을 런타임에 프로파일링 후 최적 선택
auto gemms = DeviceOperationInstanceFactory<
    DeviceGemmInstance>::GetInstances();

float best_time = std::numeric_limits<float>::max();
int   best_idx  = -1;

for (int i = 0; i < (int)gemms.size(); i++) {
    auto& gemm = *gemms[i];
    auto arg = gemm.MakeArgument(
        dA, dB, dC, M, N, K, K, K, N,
        AElementOp{}, BElementOp{}, CEElementOp{});
    if (!gemm.IsSupportedArgument(arg)) continue;

    auto invoker = gemm.MakeInvoker();
    float t = invoker.Run(
        arg, StreamConfig{stream, true}); // time_kernel=true
    if (t < best_time) { best_time = t; best_idx = i; }
}
// best_idx 인스턴스로 이후 실행

CK 퓨전 연산

CK의 강점 중 하나는 에필로그 퓨전(epilogue fusion)입니다. GEMM 결과에 적용되는 Bias 덧셈, ReLU/GELU 활성화, LayerNorm, Softmax를 출력 타일을 레지스터에 유지한 채 연속 적용하여 Global Memory 왕복을 줄입니다. Transformer 모델의 Attention 연산은 QKᵀ GEMM → Softmax → V GEMM 세 단계가 하나의 CK 퓨전 커널로 처리됩니다(Flash Attention 방식).

// CK 퓨전 연산: GEMM + BiasAdd + GELU in one kernel
#include <ck/tensor_operation/gpu/device/device_gemm_multiple_d_xdl_cshuffle.hpp>

// 에필로그: C = GELU(A*B + Bias)
using CDEElementOp = ck::tensor_operation::element_wise::AddGelu;

using DeviceFusedGemm = DeviceGemmMultipleD_Xdl_CShuffle<
    ALayout, BLayout,
    Tuple<BiasLayout>,  // D 텐서 목록 (Bias)
    ELayout,
    ADataType, BDataType,
    Tuple<BiasDataType>,
    EDataType,
    AccDataType,
    AElementOp, BElementOp, CDEElementOp,
    GemmDefault,
    256, 256, 128, 32, 8, 8, 32, 32, 4, 2
>;

auto gemm_fused = DeviceFusedGemm{};
auto arg = gemm_fused.MakeArgument(
    dA, dB,
    std::array<const void*, 1>{dBias},  // D 텐서 배열
    dE,
    M, N, K,
    K, K,
    std::array<int, 1>{0},  // Bias stride (broadcast)
    N,
    AElementOp{}, BElementOp{},
    CDEElementOp{});

auto invoker = gemm_fused.MakeInvoker();
invoker.Run(arg, StreamConfig{stream, false});

rocProfiler / rocTracer

ROCm 프로파일링 인프라는 GPU의 하드웨어 성능 카운터(PMC)를 수집하는 rocProfiler와 HIP/ROCr API 호출을 추적하는 rocTracer로 구성됩니다. omniperf는 이들 위에 구축된 고수준 분석 도구로, roofline 분석, 캐시 효율, 메모리 대역폭 활용률 등 종합적인 성능 리포트를 생성합니다. rocProf v2(ROCm 6.x)는 rocProfiler SDK를 기반으로 한 새 세대 프로파일러입니다.

# rocprof: 커널 실행 시간 프로파일링
rocprof --stats ./my_hip_app arg1 arg2
# 출력: results.csv (커널별 실행 횟수, 총 시간, 평균 시간)

# 특정 PMC 카운터 수집
rocprof --input metrics.txt ./my_hip_app
# metrics.txt 내용:
# pmc: SQ_WAVES, SQ_INSTS_VALU, SQ_INSTS_VMEM_RD
# range: 0:10000  (커널 N번째 실행 범위)

# HIP API 타임라인 추적 (roctracer 사용)
rocprof --hip-trace ./my_hip_app
# 출력: results.json (Chrome Trace 포맷)
# chrome://tracing 에서 시각화 가능

# HSA API 레벨 추적
rocprof --hsa-trace ./my_hip_app

# rocprof 주요 PMC 메트릭 예시 (metrics.txt)
# -- VALU 활용률: VALU 명령 / 전체 클럭 사이클
pmc: SQ_WAVES SQ_INSTS_VALU SQ_INSTS_SALU SQ_INSTS_VMEM_RD SQ_INSTS_VMEM_WR

# -- L1/L2 캐시 히트율
pmc: TCP_TOTAL_CACHE_ACCESSES TCP_TCP_LATENCY TCP_TOTAL_HIT

# -- 메모리 컨트롤러 대역폭
pmc: TCC_EA_RDREQ_32B TCC_EA_WRREQ_32B TCC_WRREQ_64B

# -- Wavefront 점유율
pmc: SQ_WAIT_INST_ANY SQ_ACTIVE_INST_ANY

/* roctracer API를 애플리케이션 코드에서 직접 사용 */
#include <roctracer/roctracer_hip.h>

/* HIP API 활동 콜백 등록 */
roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_API);
roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS);

/* 사용자 정의 영역 마킹 */
roctxRangePush("Forward Pass");
runForwardPass(...);
roctxRangePop();

roctxRangePush("Backward Pass");
runBackwardPass(...);
roctxRangePop();

# omniperf: 심층 성능 분석 (MI100/MI200/MI300 전용)
# 1단계: 프로파일 수집
omniperf profile -n my_profile -- ./my_hip_app

# 2단계: 분석 리포트 생성
omniperf analyze -p workloads/my_profile/mi300x_a1

# 특정 분석 모듈만 선택
omniperf analyze -p workloads/my_profile/mi300x_a1 \
    --list-metrics  # 가용 메트릭 목록

# Roofline 분석: 산술 강도 vs 실제 성능
omniperf analyze -p workloads/my_profile/mi300x_a1 \
    --roofline

omniperf 핵심 메트릭 심층 분석

omniperf는 단순한 실행 시간 측정을 넘어 GPU 하드웨어 카운터를 수집하고 Roofline 모델, 메모리 계층 효율, wavefront 스톨 원인까지 분석하는 종합 프로파일링 도구입니다. omniperf profile 단계에서 PMC를 수집하고 omniperf analyze에서 리포트를 생성합니다. GUI 모드(omniperf analyze --gui)는 웹 브라우저 기반 대화형 인터페이스를 제공합니다.

# omniperf 상세 메트릭 수집 및 분석

# 1단계: PMC 카운터 수집 (여러 패스 자동 분할)
omniperf profile   --name my_gemm_run   --hip-trace   --no-roof   -- ./my_gemm_app

# 2단계: 기본 분석 리포트 출력
omniperf analyze   --path workloads/my_gemm_run   --dispatch 0   --metric-list

# 3단계: 특정 메트릭 그룹만 조회
omniperf analyze   --path workloads/my_gemm_run   --dispatch 0   --block 2.1   # 2.1 = Compute Pipeline (VALU/MFMA 활용률)

# 4단계: Roofline 분석 (산술 강도 vs 대역폭 지붕)
omniperf profile --name roofline_run --roof-only -- ./my_app
omniperf analyze --path workloads/roofline_run --roof

# 5단계: 메모리 계층 분석 (L1/L2 히트율)
omniperf analyze   --path workloads/my_gemm_run   --dispatch 0   --block 5.1   # 5.1 = L1 캐시 통계
  --block 5.2   # 5.2 = L2 캐시 통계

# 6단계: CSV로 내보내기 (자동화 파이프라인 통합)
omniperf analyze   --path workloads/my_gemm_run   --dispatch 0   --output-format csv   --output my_metrics.csv

rocProfiler v2 SDK (rocprofiler-sdk)

ROCm 6.x에서 새롭게 도입된 rocprofiler-sdk는 기존 rocProfiler v1의 후속입니다. 콜백 기반 API로 HIP API 호출, 커널 실행, 메모리 복사 등 다양한 이벤트를 프로그래밍 방식으로 추적할 수 있습니다. v1보다 멀티 에이전트(multi-GPU) 지원과 스레드 안전성이 크게 개선되었으며, Perfetto, OpenTelemetry와의 통합도 지원합니다.

/* rocprofiler-sdk 콜백 API 사용 예시 */
#include <rocprofiler-sdk/rocprofiler.h>

/* 커널 dispatch 콜백 함수 */
void kernel_dispatch_callback(
    rocprofiler_callback_tracing_record_t record,
    rocprofiler_user_data_t              *user_data,
    void                                 *callback_data)
{
    if (record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API)
        return;

    auto *api_data = (const rocprofiler_callback_tracing_hip_api_data_t *)record.payload;

    if (record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) {
        /* API 호출 진입 시점 */
        printf("HIP API 진입: %s
",
               rocprofiler_get_hip_runtime_api_name(record.operation));
    } else if (record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) {
        /* API 호출 종료 시점 (반환값 등 확인 가능) */
        printf("HIP API 완료: %s, 시간=%lu ns
",
               rocprofiler_get_hip_runtime_api_name(record.operation),
               record.end_timestamp - record.start_timestamp);
    }
}

/* 커널 실행 완료 콜백 (하드웨어 타이밍 포함) */
void kernel_completion_callback(
    rocprofiler_callback_tracing_record_t record,
    rocprofiler_user_data_t              *user_data,
    void                                 *callback_data)
{
    if (record.kind != ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH)
        return;

    auto *kern = (const rocprofiler_callback_tracing_kernel_dispatch_data_t *)
                  record.payload;

    printf("커널: %s
", kern->kernel_name);
    printf("  GPU 실행 시간: %lu ns
", kern->end_timestamp - kern->start_timestamp);
    printf("  그리드: %u x %u x %u
",
           kern->dispatch_info.grid_size.x,
           kern->dispatch_info.grid_size.y,
           kern->dispatch_info.grid_size.z);
}

/* 초기화: 콜백 서비스 등록 */
void setup_profiling()
{
    rocprofiler_context_id_t ctx_id;
    rocprofiler_create_context(&ctx_id);

    /* HIP Runtime API 추적 활성화 */
    rocprofiler_configure_callback_tracing_service(
        ctx_id,
        ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API,
        nullptr, 0,           /* 모든 API 추적 (필터 없음) */
        kernel_dispatch_callback,
        nullptr
    );

    /* 커널 dispatch 완료 추적 활성화 */
    rocprofiler_configure_callback_tracing_service(
        ctx_id,
        ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH,
        nullptr, 0,
        kernel_completion_callback,
        nullptr
    );

    rocprofiler_start_context(ctx_id);
    /* ... 애플리케이션 실행 ... */
    rocprofiler_stop_context(ctx_id);
    rocprofiler_destroy_context(ctx_id);
}

Multi-GPU 프로그래밍

ROCm은 여러 GPU를 연결하는 두 가지 인터커넥트를 지원합니다. XGMI(eXtended Global Memory Interface)는 AMD의 Infinity Fabric 기반 고대역폭 GPU 간 직접 링크로, 피어 GPU의 메모리를 원격 접근하거나 RDMA 방식으로 데이터를 전송합니다. PCIe는 표준 버스를 통한 연결로 대역폭이 낮지만 범용성이 높습니다. 서버 플랫폼에서 MI300X 8개는 XGMI 완전 연결(All-to-All) 토폴로지를 구성합니다.

/* P2P(Peer-to-Peer) 직접 GPU 간 메모리 접근 */
int canAccessPeer;

/* GPU 0 → GPU 1 P2P 가능 여부 확인 */
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
if (canAccessPeer) {
    hipSetDevice(0);
    hipDeviceEnablePeerAccess(1, 0);  /* GPU0가 GPU1에 접근 허용 */
}

/* GPU 0의 메모리 할당 */
hipSetDevice(0);
float *d_src;
hipMalloc(&d_src, N * sizeof(float));

/* GPU 1의 메모리 할당 */
hipSetDevice(1);
float *d_dst;
hipMalloc(&d_dst, N * sizeof(float));

/* GPU 0에서 GPU 1로 직접 P2P 복사 (PCIe/XGMI DMA) */
hipMemcpyPeer(d_dst, 1, d_src, 0, N * sizeof(float));

/* GPU 1의 커널에서 GPU 0의 메모리 직접 접근 (XGMI 필요) */
hipSetDevice(1);
crossGpuKernel<<<grid, block>>>(d_src, d_dst, N);
/* d_src는 GPU 0의 메모리지만 XGMI로 직접 접근 가능 */

/* RCCL AllReduce: Multi-GPU 그래디언트 합산 예제 */
#include <rccl/rccl.h>

int nGPUs = 8;
ncclComm_t comms[8];
int devList[8] = {0,1,2,3,4,5,6,7};

/* RCCL 통신자 초기화 */
ncclCommInitAll(comms, nGPUs, devList);

/* 각 GPU에 스트림 생성 */
hipStream_t streams[8];
for (int i = 0; i < nGPUs; ++i) {
    hipSetDevice(i);
    hipStreamCreate(&streams[i]);
}

/* AllReduce: 모든 GPU의 그래디언트를 합산하여 모두에게 배포 */
ncclGroupStart();
for (int i = 0; i < nGPUs; ++i) {
    ncclAllReduce(
        d_grad[i],        /* 각 GPU의 그래디언트 버퍼 */
        d_grad[i],        /* 결과도 같은 위치에 */
        N,                /* 요소 수 */
        ncclFloat,        /* FP32 */
        ncclSum,          /* 합산 연산 */
        comms[i],
        streams[i]);
}
ncclGroupEnd();

for (int i = 0; i < nGPUs; ++i) {
    ncclCommDestroy(comms[i]);
}

# XGMI 토폴로지 및 링크 상태 확인
rocm-smi --showtopo
# 출력 예 (4 GPU, 완전 연결):
# GPU[0]: GPU[1]=XGMI GPU[2]=XGMI GPU[3]=XGMI
# GPU[1]: GPU[0]=XGMI GPU[2]=XGMI GPU[3]=XGMI
# ...

# XGMI 링크 대역폭 측정
rocm-smi --showxgmibw
# 또는 p2pBandwidthLatencyTest 도구 사용
p2pBandwidthLatencyTest

# GPU 토폴로지 상세 (Linux NUMA 노드 포함)
rocminfo | grep -A5 "Topology"
cat /sys/class/kfd/kfd/topology/nodes/*/io_links/*/type

GPUDirect RDMA / ROCnRDMA

GPUDirect RDMA는 GPU VRAM과 네트워크 인터페이스 카드(NIC) 사이에서 CPU와 시스템 메모리를 경유하지 않고 직접 데이터를 전송하는 기술입니다. 전통적인 GPU↔네트워크 전송 경로는 GPU VRAM → CPU DRAM(복사) → NIC(DMA)였으나, GPUDirect RDMA를 사용하면 GPU VRAM → NIC(DMA) 경로로 단축되어 CPU 개입과 중간 복사 오버헤드가 제거됩니다. AMD에서는 이 기술의 구현을 ROCnRDMA라 부르며, amdgpu 커널 드라이버의 amdgpu_dmabuf_ops 및 PEER Direct 인터페이스를 통해 제공합니다.

HPC 클러스터에서 AllReduce 같은 집합 통신(collective)이 초대형 모델 학습의 병목이 됩니다. MI250X/MI300X + InfiniBand HDR(200 Gb/s)/NDR(400 Gb/s) 조합에서 ROCnRDMA는 노드 간 통신 지연시간을 크게 낮추고, RCCL이 이를 자동으로 활용합니다.

ROCnRDMA 개요와 원리

ROCnRDMA는 Linux 커널의 PEER Direct(또는 peer_mem) 인터페이스를 활용합니다. NIC 드라이버(예: Mellanox OFED의 mlx5_core)가 GPU 메모리 영역에 대한 DMA 매핑을 요청할 때, 표준 경로에서는 시스템 메모리만 허용됩니다. peer_mem 인터페이스는 서드파티 메모리 공급자(amdgpu)가 등록되면 NIC 드라이버가 GPU VRAM 영역을 직접 DMA 매핑할 수 있도록 합니다. ROCm 6.x 이후 DMA-BUF(Linux 5.12+)를 기반으로 한 방식으로 전환 중입니다.

## ROCnRDMA 환경 설정

# 커널 모듈 확인
lsmod | grep -E 'amdgpu|ib_core|mlx5'

# ROCm + OFED 설치 (ROCm 6.x → MLNX_OFED 5.8+)
apt-get install -y rocm-dev mlnx-ofed-all

# 구형: peer_mem 방식
modprobe amdgpu_peer_mem && dmesg | grep peer_mem

# 신형: DMA-BUF (ROCm 6.x 자동)
dmesg | grep -i "dmabuf\|p2p"

# BAR1 크기 확인 (VRAM 크기와 일치 필요)
lspci -vv -s <GPU_BDF> | grep "Memory at"

# 대역폭 테스트
ib_write_bw -d mlx5_0 --use_rocm=0

ROCnRDMA libibverbs 연동

GPU와 NIC가 동일한 NUMA 노드에 위치하는 것이 P2P DMA 성능에 중요합니다. ibv_reg_mr()에 GPU VRAM 포인터를 전달하면 ROCnRDMA/DMA-BUF가 내부적으로 처리합니다.

// ROCnRDMA: GPU 버퍼를 InfiniBand로 직접 전송
#include <infiniband/verbs.h>
#include <hip/hip_runtime.h>

ibv_mr* register_gpu_buf(
    ibv_pd* pd, void* gpu_ptr, size_t size)
{
    int flags = IBV_ACCESS_LOCAL_WRITE
              | IBV_ACCESS_REMOTE_WRITE
              | IBV_ACCESS_REMOTE_READ;
    return ibv_reg_mr(pd, gpu_ptr, size, flags);
}

void rdma_write_gpu(
    ibv_qp* qp, ibv_mr* src_mr,
    void* src, uint64_t dst_addr,
    uint32_t rkey, size_t size)
{
    ibv_sge sge = {
        .addr   = (uint64_t)src,
        .length = (uint32_t)size,
        .lkey   = src_mr->lkey,
    };
    ibv_send_wr wr = {
        .opcode     = IBV_WR_RDMA_WRITE,
        .send_flags = IBV_SEND_SIGNALED,
        .wr.rdma    = { .remote_addr = dst_addr, .rkey = rkey },
        .sg_list    = &sge,
        .num_sge    = 1,
    };
    ibv_send_wr* bad;
    ibv_post_send(qp, &wr, &bad);
}

MPI GPU-Aware 집합 통신

ROCm 환경에서 GPU-aware MPI는 통신 버퍼가 GPU VRAM에 있어도 MPI가 직접 처리합니다. Open MPI + UCX 조합이 ROCnRDMA를 자동으로 활용하며, CPU DRAM 복사 단계가 생략됩니다.

// GPU-Aware MPI AllReduce
#include <mpi.h>
#include <hip/hip_runtime.h>

int main(int argc, char** argv) {
    MPI_Init(&argc, &argv);
    int rank;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    hipSetDevice(rank % 8);

    const int N = 1 << 24;
    float* dGrad;
    hipMalloc(&dGrad, N * sizeof(float));

    computeGradients<<<grid, block>>>(dGrad, N);
    hipDeviceSynchronize();

    // GPU VRAM 포인터를 직접 MPI에 전달
    MPI_Allreduce(MPI_IN_PLACE, dGrad, N,
                  MPI_FLOAT, MPI_SUM, MPI_COMM_WORLD);

    hipFree(dGrad);
    MPI_Finalize();
}

## Open MPI + UCX 환경 설정
export OMPI_MCA_pml=ucx
export UCX_TLS=rc,rocm_copy,rocm_ipc,gdr_copy
mpirun -np 32 --map-by ppr:8:node ./my_app

컨테이너와 가상화

ROCm은 Docker 컨테이너, Kubernetes(K8s), 그리고 제한적으로 가상화 환경에서도 동작합니다. AMD는 DockerHub에서 공식 ROCm 베이스 이미지를 제공하며, K8s에서는 AMD GPU Device Plugin을 통해 GPU 자원 할당을 자동화합니다. 가상화 환경에서는 SR-IOV(Single Root I/O Virtualization)나 MxGPU 기술로 GPU를 가상 기계에 노출합니다.

# Docker with ROCm: AMD 공식 베이스 이미지 사용
docker pull rocm/rocm-terminal:6.1
docker pull rocm/pytorch:latest

# ROCm 컨테이너 실행 (GPU 디바이스 전달)
docker run -it --rm \
    --device=/dev/kfd \           # KFD 디바이스 (필수)
    --device=/dev/dri \           # DRI 디바이스 (렌더 노드)
    --group-add video \           # video 그룹 추가
    --group-add render \          # render 그룹 추가
    --ipc=host \                  # RCCL 멀티 GPU를 위한 공유 메모리
    --cap-add=SYS_PTRACE \        # 디버거 사용 시
    --security-opt seccomp=unconfined \
    -v /home/user/data:/data \
    rocm/pytorch:latest \
    python3 train.py

# 특정 GPU만 컨테이너에 전달
# ROCR_VISIBLE_DEVICES=0,2 docker run ... (GPU 0, 2만 노출)
docker run --device=/dev/kfd --device=/dev/dri/renderD128 \
    -e ROCR_VISIBLE_DEVICES=0 \
    rocm/pytorch:latest python3 -c "import torch; print(torch.cuda.device_count())"

# Kubernetes: AMD GPU Device Plugin 배포
# 1. Device Plugin 설치
# kubectl apply -f https://raw.githubusercontent.com/RadeonOpenCompute/k8s-device-plugin/main/k8s-ds-amdgpu-dp.yaml

# 2. GPU를 요청하는 Pod 스펙
apiVersion: v1
kind: Pod
metadata:
  name: rocm-training-pod
spec:
  containers:
  - name: trainer
    image: rocm/pytorch:latest
    command: ["python3", "train.py"]
    resources:
      limits:
        amd.com/gpu: 4          # GPU 4개 요청
    env:
    - name: HIP_VISIBLE_DEVICES
      value: "0,1,2,3"
    volumeMounts:
    - name: dataset
      mountPath: /data
  volumes:
  - name: dataset
    persistentVolumeClaim:
      claimName: training-data-pvc

# SR-IOV 설정 (AMD MxGPU 가상화)
# amdgpu 드라이버에서 SR-IOV 활성화
modprobe amdgpu sriov_enabled=1

# VF(Virtual Function) 생성
echo 4 > /sys/bus/pci/devices/0000:03:00.0/sriov_numvfs
# 4개의 VF 생성 (각각 /dev/dri/renderD128~D131)

# VF를 VM에 passthrough (libvirt/QEMU)
virsh nodedev-detach pci_0000_03_08_0   # VF PCI 주소
# virsh XML에서 hostdev 항목으로 VM에 할당

# ROCm 환경에서 VF 확인
rocm-smi --showproductname
# VF에서는 물리 GPU 스펙의 1/N으로 표시됨

AI/ML 프레임워크 통합

ROCm은 주요 AI/ML 프레임워크의 공식 백엔드로 지원됩니다. PyTorch는 torch.cuda API를 AMD GPU에서 그대로 사용할 수 있으며, TensorFlow-ROCm은 XLA-HLO 컴파일러가 ROCm 백엔드를 사용합니다. JAX는 XLA를 통해 ROCm을 지원합니다. 이 프레임워크들은 내부적으로 MIOpen, rocBLAS, RCCL을 호출하여 실제 GPU 연산을 수행합니다.

# PyTorch with ROCm: torch.cuda API는 변경 없이 AMD GPU에서 동작
import torch

# ROCm에서는 torch.cuda.is_available()이 True 반환
print(torch.cuda.is_available())      # True (ROCm 설치 시)
print(torch.cuda.get_device_name(0))  # AMD Instinct MI300X
print(torch.version.hip)              # ROCm HIP 버전

# 모델 학습 예제 (코드 변경 불필요)
device = torch.device("cuda")  # ROCm도 "cuda" 디바이스 명칭 사용
model = torch.nn.Linear(1024, 1024).to(device)
optimizer = torch.optim.Adam(model.parameters())

x = torch.randn(256, 1024, device=device)
y = model(x)
loss = y.sum()
loss.backward()
optimizer.step()

# Mixed Precision (FP16/BF16) with GradScaler
from torch.cuda.amp import autocast, GradScaler
scaler = GradScaler()

with autocast(dtype=torch.bfloat16):  # BF16 (MI300X 권장)
    output = model(x)
    loss = output.sum()
scaler.scale(loss).backward()
scaler.step(optimizer)
scaler.update()

# TensorFlow-ROCm 사용
import tensorflow as tf

# GPU 확인
gpus = tf.config.list_physical_devices('GPU')
print(f"ROCm GPUs: {gpus}")

# tf.data와 Keras로 학습
model = tf.keras.Sequential([
    tf.keras.layers.Dense(512, activation='relu'),
    tf.keras.layers.Dense(10, activation='softmax')
])
model.compile(optimizer='adam', loss='sparse_categorical_crossentropy')

# Mixed precision (BF16)
tf.keras.mixed_precision.set_global_policy('mixed_bfloat16')

# Multi-GPU 분산 학습
strategy = tf.distribute.MirroredStrategy()
with strategy.scope():
    dist_model = tf.keras.Sequential([...])

# JAX with ROCm backend
import jax
import jax.numpy as jnp

# ROCm 백엔드 활성화 (환경 변수 설정 필요)
# export JAX_BACKEND_TARGET=rocm
# 또는 xla_flags 설정

print(jax.devices())  # [RocmDevice(id=0), RocmDevice(id=1), ...]

# JIT 컴파일 함수
@jax.jit
def matmul(a, b):
    return jnp.dot(a, b)

a = jnp.ones((4096, 4096))
b = jnp.ones((4096, 4096))
c = matmul(a, b)

# pmap: 데이터 병렬화 (Multi-GPU)
@jax.pmap
def parallel_fn(x):
    return x * 2

xs = jnp.stack([a] * jax.device_count())
ys = parallel_fn(xs)

LLM 추론 최적화

AMD MI300X(192 GB HBM3)는 현존하는 가속기 중 단일 카드 메모리 용량이 가장 크며, 이는 대형 언어 모델(LLM) 추론에 근본적인 이점을 제공합니다. NVIDIA H100 SXM5(80 GB HBM3e)와 비교하면 단일 카드 기준 2.4배 더 많은 모델 파라미터를 수용할 수 있어, 70B~405B 파라미터 모델을 더 적은 수의 카드로 서빙할 수 있습니다. ROCm 생태계에서 주요 LLM 추론 프레임워크들이 공식 AMD GPU 지원을 추가하고 있습니다.

LLM 추론의 성능은 크게 두 단계로 나뉩니다. 프리필(Prefill) 단계는 프롬프트 토큰 전체를 병렬 처리하여 KV 캐시를 생성하는 컴퓨트 집약 단계이고, 디코드(Decode) 단계는 토큰을 하나씩 자기회귀적으로 생성하는 메모리 대역폭 집약 단계입니다. MI300X의 HBM3는 5.2 TB/s 이상의 메모리 대역폭을 제공하여 디코드 단계의 병목을 완화합니다.

LLM 서빙 프레임워크

vLLM은 PagedAttention 알고리즘으로 KV 캐시 단편화를 제거하고 연속 배칭(continuous batching)을 구현한 고성능 서빙 엔진으로, ROCm 공식 지원이 vLLM 0.2.x 이후 추가되어 MI200/MI300 시리즈에서 작동합니다. llama.cpp는 GGML 양자화 포맷을 지원하며 HIP 백엔드로 ROCm GPU에서 실행됩니다.

## vLLM on ROCm 설치 및 구성

# 1. ROCm 지원 vLLM 설치
pip install vllm --extra-index-url \
    https://download.pytorch.org/whl/rocm6.0

# 2. 기본 서버 실행 (MI300X, Llama-3-70B BF16)
python -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3-70B-Instruct \
    --dtype bfloat16 \
    --max-model-len 8192 \
    --tensor-parallel-size 1 \
    --gpu-memory-utilization 0.90

# 3. FP8 KV 캐시 + FP8 양자화 (MI300X 전용)
python -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3-70B-Instruct \
    --dtype bfloat16 \
    --kv-cache-dtype fp8 \
    --quantization fp8 \
    --max-model-len 131072 \
    --tensor-parallel-size 1

# 4. 멀티 GPU 텐서 병렬 (8× MI300X, 405B 모델)
python -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3.1-405B-Instruct \
    --dtype bfloat16 \
    --tensor-parallel-size 8 \
    --max-model-len 32768 \
    --gpu-memory-utilization 0.95

KV 캐시 메모리 관리

MI300X의 192 GB HBM은 LLM 추론에서 메모리 용량의 병목을 크게 완화합니다. Llama-3-70B(BF16)는 약 140 GB를 소비하여 단일 MI300X에서 완전히 수용할 수 있으며, 나머지 ~50 GB를 KV 캐시와 활성화 버퍼에 활용할 수 있습니다.

## Flash Attention ROCm 구성

# ROCm용 Flash Attention 설치
pip install flash-attn --no-build-isolation \
    --extra-index-url https://download.pytorch.org/whl/rocm6.0

# 소스 빌드 (ROCm)
export ROCM_HOME=/opt/rocm
export GPU_ARCHS="gfx942"
git clone https://github.com/ROCm/flash-attention
cd flash-attention && pip install -e . --no-build-isolation

## Flash Attention 사용 (BF16, causal)
import torch
from flash_attn import flash_attn_func

Q = torch.randn(2, 4096, 32, 128,
                dtype=torch.bfloat16, device='cuda')
K = torch.randn_like(Q)
V = torch.randn_like(Q)

out = flash_attn_func(Q, K, V, dropout_p=0.0, causal=True)
# HBM 왕복 최소화: softmax+matmul 온칩 퓨전

양자화 옵션

MI300X는 FP8(E4M3/E5M2)을 하드웨어 수준에서 지원합니다(CDNA3 아키텍처). FP8 양자화는 BF16 대비 모델 크기와 KV 캐시를 절반으로 줄이면서도 대부분의 모델에서 2% 이하의 정확도 손실을 보입니다. llama.cpp는 GGUF 포맷의 Q4_K_M, Q8_0 등 다양한 정밀도 옵션을 제공합니다.

## llama.cpp HIP 빌드 및 양자화

git clone https://github.com/ggerganov/llama.cpp
cd llama.cpp
mkdir build-hip && cd build-hip
cmake .. -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx942 \
         -DCMAKE_BUILD_TYPE=Release
cmake --build . --config Release -j$(nproc)

# Q4_K_M: 4-bit 혼합 양자화 (품질 우선)
./quantize /path/to/llama3-70b-f16.gguf \
           llama3-70b-q4_k_m.gguf Q4_K_M

# GPU 추론 서버 (전 레이어 GPU 적재)
./llama-server \
    -m llama3-70b-q4_k_m.gguf \
    -ngl 80 --host 0.0.0.0 --port 8080 \
    -c 4096 --threads 16

성능 비교 및 벤치마크

MI300X는 메모리 대역폭 우위(5.2 TB/s vs H100의 3.35 TB/s)로 디코드 단계 처리량에서 유리하며, 메모리 용량 우위(192 GB vs 80 GB)로 대형 모델의 단일 카드 수용 및 긴 컨텍스트 서빙에서 두드러집니다.

디버깅과 트러블슈팅

ROCm 디버깅은 사용자 공간(HIP 런타임 오류, 메모리 접근 위반)부터 커널 공간(amdgpu/KFD 오류, GPU hang) 까지 여러 레벨에 걸쳐 있습니다. 주요 도구로는 rocgdb(GDB 기반 GPU 디버거), umr(amdgpu 레지스터 접근), dmesg(커널 GPU 에러 로그), 다양한 환경 변수를 통한 런타임 디버그 모드가 있습니다.

# rocgdb: HIP 애플리케이션 GPU 디버깅
# 컴파일 시 디버그 심볼 포함
hipcc -g -O0 -o myapp_debug myapp.cpp

# rocgdb 시작
rocgdb ./myapp_debug

# rocgdb 명령어 (GDB 호환)
# (rocgdb) break vectorAdd        # GPU 커널 브레이크포인트
# (rocgdb) run                    # 실행
# (rocgdb) info threads           # 모든 wavefront/스레드 표시
# (rocgdb) thread 1.1:0           # wavefront 1의 스레드 0 선택
# (rocgdb) print threadIdx.x      # GPU 빌트인 변수 출력
# (rocgdb) x/16f $s_p             # 스칼라 레지스터 메모리 덤프
# (rocgdb) info registers         # 레지스터 상태

# umr: amdgpu 레지스터 직접 접근 (디버그용)
# 설치: apt install umr 또는 소스 빌드

# GPU 레지스터 읽기 (예: GFX 상태)
umr -asic gfx942 -read 0x2100  # GRBM_STATUS 읽기

# 파워 레지스터 상태 확인
umr -asic gfx942 -pp          # 전원 상태 덤프

# 클럭 정보
umr -asic gfx942 -clocks

# GPU hang 분석: dmesg 로그 확인
dmesg -T | grep -E "amdgpu|kfd" | grep -i "error\|hang\|reset\|timeout"

# 정상 복구 로그 예:
# [Mon Mar 10 15:23:41] amdgpu 0000:03:00.0: amdgpu: GPU reset begin!
# [Mon Mar 10 15:23:42] amdgpu 0000:03:00.0: amdgpu: GPU reset succeeded

# GPU hang 상세 정보 (amdgpu_fence_info)
cat /sys/kernel/debug/dri/0/amdgpu_fence_info

# GPU 렌더링 진행 중인 작업 상태
cat /sys/kernel/debug/dri/0/amdgpu_ring_gfx_0

# GPU 메모리 누수 확인
cat /sys/kernel/debug/dri/0/amdgpu_gtt_mm
cat /sys/kernel/debug/dri/0/amdgpu_vram_mm

# ROCm 디버그 환경 변수

# HIP 런타임 오류 검사 활성화 (느리지만 상세 오류 정보)
export AMD_LOG_LEVEL=4         # 0=오류, 1=경고, 2=정보, 3=디버그, 4=상세
export HSA_ENABLE_SDMA=0       # SDMA 비활성화 (오류 격리)
export HIP_LAUNCH_BLOCKING=1   # 모든 커널을 동기적으로 실행 (오류 위치 특정)
export HIPCC_VERBOSE=1         # 컴파일러 상세 출력

# Address Sanitizer (메모리 오류 감지)
hipcc -fsanitize=address -o myapp myapp.cpp
export ASAN_OPTIONS=detect_leaks=1

# KFD 이벤트 추적
export ROCR_DEBUG_API=1

# GPU 코어 덤프 활성화 (amdgpu hang 시)
export AMD_DEBUG=vm_fault_info
ulimit -c unlimited

성능 최적화

ROCm GPU 성능 최적화는 roofline 모델에서 시작합니다. 커널이 메모리 바운드인지 컴퓨트 바운드인지 먼저 파악하고, 각 경우에 맞는 최적화 전략을 적용합니다. AMD GPU에서의 핵심 최적화 요소는 CU occupancy(동시 실행 wavefront 수), VGPR 사용량, LDS 활용, 메모리 접근 패턴, 분기 발산(divergence)입니다.

/* occupancy 계산 및 최적 블록 크기 선택 */
#include <hip/hip_runtime.h>

/* 특정 커널의 최적 블록 크기 자동 계산 */
int minGridSize, blockSize;
hipOccupancyMaxPotentialBlockSize(
    &minGridSize,
    &blockSize,
    myKernel,         /* 커널 함수 포인터 */
    0,                 /* 동적 LDS 크기 */
    0);                /* 블록 크기 제한 없음 */
printf("Optimal block size: %d\n", blockSize);

/* occupancy 수치 조회 */
int numBlocks;
hipOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocks, myKernel, blockSize, 0);
printf("Blocks per CU: %d\n", numBlocks);
printf("Wavefronts per CU: %d\n",
       numBlocks * (blockSize / 64));

/* __launch_bounds__: 컴파일러에 실행 바운드 힌트 제공 */
/* 이로써 컴파일러가 더 공격적인 레지스터 할당 최적화 가능 */
__global__
__attribute__((amdgpu_flat_work_group_size(64, 256)))  /* min, max 블록 크기 */
void optimizedKernel(float *data, int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        /* ... */
    }
}

/* CUDA 스타일 launch_bounds (HIP 호환) */
__global__ __launch_bounds__(256, 2)  /* maxThreadsPerBlock, minBlocksPerMP */
void anotherKernel(float *data) { /* ... */ }

/* 메모리 접근 최적화: 합체 접근(Coalescing) 패턴 */

/* 나쁜 패턴: 열 우선 접근 → VRAM 접근이 산발적 */
__global__ void bad_access(float *matrix, int rows, int cols)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    for (int row = 0; row < rows; ++row)
        matrix[row * cols + col] += 1.0f;  /* stride=cols, 비합체 */
}

/* 좋은 패턴: 행 우선 접근 → VRAM 합체 접근 */
__global__ void good_access(float *matrix, int rows, int cols)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < rows && col < cols)
        matrix[row * cols + col] += 1.0f;  /* 연속 접근, 합체 */
}

# AMDGCN 어셈블리 확인 (컴파일러 출력 검사)
hipcc -O3 --offload-arch=gfx942 \
    -save-temps \          # 중간 파일 보존
    -o myapp myapp.cpp

# 생성된 어셈블리 확인
ls -la *.s *.ll           # .s=어셈블리, .ll=LLVM IR

# 커널 레지스터 사용량 확인
hipcc -O3 --offload-arch=gfx942 \
    -Rpass-analysis=regalloc \
    myapp.cpp 2>&1 | grep "spilled\|VGPR"

# 어셈블리에서 VGPR 수 확인
grep "NumVgprs" *.s

메모리 코얼레싱 심층 분석

메모리 코얼레싱(coalescing)은 wavefront 내 64개 스레드의 메모리 접근을 가능한 한 적은 수의 VRAM 트랜잭션으로 병합하는 최적화입니다. AMD GPU의 최소 메모리 트랜잭션 단위는 64 바이트 또는 128 바이트입니다. 64개 스레드가 각각 4 바이트(float)에 접근할 때 완전 코얼레싱되면 256 바이트 = 4 트랜잭션으로 처리되지만, 완전 비코얼레싱(각 스레드가 임의 주소 접근)이면 최대 64 트랜잭션이 발생합니다.

/* 코얼레싱 분석 코드 패턴과 rocprof 메트릭 수집 */

/* 패턴 1: 완전 코얼레싱 (최적) - 연속 주소 접근 */
__global__ void coalesced_access(float *data, float *result, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        /* lane i가 data[tid]에 접근: 연속 → 완전 코얼레싱 */
        result[tid] = data[tid] * 2.0f;
    }
}

/* 패턴 2: 스트라이드 접근 (부분 코얼레싱) */
__global__ void strided_access(float *data, float *result, int n, int stride)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        /* lane i가 data[tid * stride] 접근: stride=2이면 50% 효율 */
        result[tid] = data[tid * stride];  /* stride > 1: 비효율적 */
    }
}

/* 패턴 3: 전치 행렬 접근 최적화 예 (LDS를 통한 코얼레싱) */
#define TILE_DIM 32
#define PAD 1  /* LDS 뱅크 충돌 방지 패딩 */

__global__ void transpose_coalesced(float *out, const float *in, int width, int height)
{
    __shared__ float tile[TILE_DIM][TILE_DIM + PAD];  /* PAD로 뱅크 충돌 방지 */

    int xin  = blockIdx.x * TILE_DIM + threadIdx.x;
    int yin  = blockIdx.y * TILE_DIM + threadIdx.y;
    int xout = blockIdx.y * TILE_DIM + threadIdx.x;
    int yout = blockIdx.x * TILE_DIM + threadIdx.y;

    /* 입력: 행 방향 코얼레싱 (연속 읽기) */
    if (xin < width && yin < height)
        tile[threadIdx.y][threadIdx.x] = in[yin * width + xin];

    __syncthreads();

    /* 출력: LDS에서 열 방향으로 읽고 행 방향으로 쓰기 (코얼레싱 유지) */
    if (xout < height && yout < width)
        out[yout * height + xout] = tile[threadIdx.x][threadIdx.y];
}

/* rocprof로 코얼레싱 효율 측정 */
/* input.txt 카운터 파일: */
/*   pmc: SQ_INSTS_VMEM_WR SQ_INSTS_VMEM_RD               */
/*        TCP_TOTAL_CACHE_ACCESSES_sum TCP_TOTAL_READ_sum  */

# bash 명령으로 수집:
# rocprof --pmc -i coalescing_counters.txt ./my_kernel

/* 코얼레싱 비율 계산: */
/* coalesce_ratio = (이론적 트랜잭션 수) / (실제 트랜잭션 수) */
/* SQ_INSTS_VMEM_RD: 발행된 벡터 메모리 읽기 명령 수            */
/* TCP_TOTAL_CACHE_ACCESSES: L1 캐시 라인 접근 수               */
/* 비율이 낮을수록(1에 가까울수록) 코얼레싱이 잘 됨              */

커널 융합 (Kernel Fusion)

커널 융합은 여러 개의 GPU 커널 호출을 하나로 합치는 최적화 기법입니다. 각 커널 실행에는 런치 오버헤드(~수 μs)와 중간 결과를 VRAM에 쓰고 다시 읽는 메모리 대역폭 낭비가 발생합니다. 융합 커널은 중간 결과를 VRAM 대신 레지스터나 LDS에 유지하므로 메모리 트래픽을 크게 줄입니다. 특히 element-wise 연산 + reduction, activation + normalization 같은 패턴에서 효과적입니다.

/* 커널 융합 예: Element-wise 연산 + Reduction 합산 */
/* 비융합: (1) scale 커널 → VRAM → (2) reduce 커널 = 2× VRAM BW 소모 */
/* 융합: scale + reduce 동시 처리, 중간 결과 레지스터 유지 */

constexpr int BLOCK_SIZE = 256;

__global__ void fused_scale_reduce(
    const float *__restrict__ input,
    float       *__restrict__ output,
    float                    scale,
    int                      n)
{
    __shared__ float smem[BLOCK_SIZE];

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int lid = threadIdx.x;

    /* 1단계 (융합): 로드 + 스케일 적용 → 중간 결과를 레지스터에 유지 */
    float val = 0.0f;
    for (int i = tid; i < n; i += gridDim.x * blockDim.x) {
        /* element-wise: ReLU + scale (비융합이면 별도 커널 필요) */
        float x = input[i] * scale;
        x = (x > 0.0f) ? x : 0.0f;  /* ReLU */
        val += x;                       /* 로컬 reduce: 레지스터 내 누산 */
    }

    /* 2단계: 블록 내 shared memory reduce */
    smem[lid] = val;
    __syncthreads();

    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (lid < s) smem[lid] += smem[lid + s];
        __syncthreads();
    }

    /* 3단계: 블록 결과를 원자적으로 글로벌 accumulator에 추가 */
    if (lid == 0) atomicAdd(output, smem[0]);
}

/* Flash Attention 스타일의 Online Softmax + Attention 융합 */
/* 비융합: QK^T → softmax → @ V = 3개 커널, O(N²) VRAM 트래픽  */
/* 융합:   타일 단위 온라인 softmax로 O(N) VRAM 트래픽            */
__global__ void online_softmax_attn(
    const float *__restrict__ Q,  /* [seq, dim] */
    const float *__restrict__ K,
    const float *__restrict__ V,
    float       *__restrict__ O,
    int seq_len, int dim)
{
    int row = blockIdx.x;  /* 이 블록이 담당하는 쿼리 행 */

    float m_i = -INFINITY;  /* 온라인 max (수치 안정성) */
    float l_i = 0.0f;       /* 누적 exp 합 */
    float acc[64] = {};    /* 출력 accumulator (레지스터 내 유지) */

    /* 키/값 타일을 순회하며 온라인으로 attention 계산 */
    for (int j = 0; j < seq_len; j++) {
        float score = 0.0f;
        /* Q[row] · K[j] 내적 계산 */
        for (int d = threadIdx.x; d < dim; d += blockDim.x)
            score += Q[row * dim + d] * K[j * dim + d];
        /* reduce score (warp shuffle) ... */

        float m_new = fmaxf(m_i, score);
        float alpha = expf(m_i - m_new);
        /* 기존 acc 재조정 + 현재 V 타일 가중치 누산 */
        l_i = l_i * alpha + expf(score - m_new);
        m_i = m_new;
        /* acc 업데이트 (V[j] * exp(score - m_new)) ... */
    }
    /* 최종 정규화: acc / l_i → O[row] 저장 */
}

커널 설정과 빌드

ROCm이 제대로 동작하려면 Linux 커널이 amdgpu, KFD, IOMMU 등 필수 기능과 함께 빌드되어야 합니다. 대부분의 배포판 커널에는 이미 amdgpu가 모듈로 포함되어 있지만, 서버 환경이나 커스텀 커널을 빌드하는 경우 아래 설정을 확인해야 합니다. DKMS를 사용하면 커널 업그레이드 시 amdgpu 모듈을 자동으로 재빌드할 수 있습니다.

# 필수 커널 Kconfig 옵션 (menuconfig 또는 .config 파일)

# DRM/GPU 기본
CONFIG_DRM=m
CONFIG_DRM_AMDGPU=m
CONFIG_DRM_AMDGPU_SI=y          # GFX6/GFX7 Southern Islands 지원
CONFIG_DRM_AMDGPU_CIK=y         # GFX7/8 Sea Islands 지원

# KFD (ROCm 컴퓨트 인터페이스)
CONFIG_HSA_AMD=m                 # KFD 모듈 (필수)
CONFIG_HSA_AMD_SVM=y             # SVM(Shared Virtual Memory) 지원
CONFIG_HSA_AMD_P2P=y             # GPU P2P 접근 지원

# IOMMU (GPU 메모리 격리)
CONFIG_AMD_IOMMU=y
CONFIG_AMD_IOMMU_V2=y            # IOMMUv2 (PASID/PRI 지원)
CONFIG_IOMMU_SVA=y               # SVA (PASID 기반 공유 가상 주소)

# 메모리 관리
CONFIG_HMM_MIRROR=y              # HMM (Heterogeneous Memory Management)
CONFIG_MEMORY_HOTPLUG=y          # CXL/HBM 메모리 핫플러그
CONFIG_DEVICE_PRIVATE=y          # 디바이스 전용 메모리 (SVM)

# 네트워킹 (ROCm RDMA)
CONFIG_INFINIBAND=m
CONFIG_INFINIBAND_USER_ACCESS=m  # uverbs (GPUDirect RDMA)

# DKMS를 통한 amdgpu 모듈 관리
# ROCm 설치 후 DKMS로 모듈 상태 확인
dkms status

# 수동으로 amdgpu-dkms 설치 (Ubuntu)
apt install amdgpu-dkms linux-headers-$(uname -r)
# amdgpu-dkms는 ROCm 리포지터리에서 제공

# DKMS 빌드 강제 실행
dkms build amdgpu/6.7.8 -k $(uname -r)
dkms install amdgpu/6.7.8 -k $(uname -r)

# 모듈 로드 확인
modprobe amdgpu
lsmod | grep amdgpu

# udev 규칙: /dev/kfd와 /dev/dri/renderD* 권한 설정
# /etc/udev/rules.d/70-amdgpu.rules

KERNEL=="kfd", GROUP="render", MODE="0660"
SUBSYSTEM=="drm", KERNEL=="renderD*", GROUP="render", MODE="0660"
SUBSYSTEM=="drm", KERNEL=="card*", GROUP="video", MODE="0660"

# udev 규칙 즉시 적용
udevadm control --reload-rules
udevadm trigger --type=subsystems --action=change
udevadm trigger --type=devices --subsystem-match=drm

# GRUB 부팅 파라미터 설정 (/etc/default/grub)

# IOMMU 활성화 (AMD IOMMU)
GRUB_CMDLINE_LINUX="amd_iommu=on iommu=pt"
# iommu=pt: passthrough 모드 (KVM VM에 GPU passthrough 시)

# amdgpu 특정 파라미터
GRUB_CMDLINE_LINUX="amdgpu.ppfeaturemask=0xffffffff"
# 전력 관리 모든 기능 활성화

# 디스플레이 없는 서버 GPU (CDNA)
GRUB_CMDLINE_LINUX="amdgpu.dc=0 nomodeset"
# KMS/디스플레이 코드 비활성화

# 적용
update-grub
reboot

ROCm 완전 설치 가이드

ROCm을 처음 설치할 때는 운영체제별로 AMD 공식 리포지터리를 설정하고 커널 드라이버(amdgpu-dkms), ROCm 런타임, 라이브러리를 순서대로 설치해야 합니다. 설치 후에는 사용자 그룹 설정, 환경 변수 설정, 기능 검증까지 완료해야 HIP 프로그램과 AI 프레임워크가 정상적으로 GPU를 인식합니다.

##############################################################
# Ubuntu 22.04 LTS (Jammy) - ROCm 6.x 완전 설치 절차
##############################################################

# 1. 사전 조건: 커널 헤더 및 빌드 도구 설치
sudo apt update
sudo apt install -y linux-headers-$(uname -r) build-essential

# 2. AMD GPU 드라이버 설치 스크립트 다운로드
wget https://repo.radeon.com/amdgpu-install/6.3/ubuntu/jammy/amdgpu-install_6.3.60300-1_all.deb
sudo dpkg -i amdgpu-install_6.3.60300-1_all.deb
sudo apt update

# 3. ROCm 전체 스택 설치 (드라이버 + 런타임 + 라이브러리)
#    --usecase=rocm: 컴퓨트 전용 (그래픽 불필요 시)
#    --usecase=graphics: 그래픽 + 컴퓨트
#    --usecase=rocm,graphics,opencl: 다목적
sudo amdgpu-install --usecase=rocm --no-dkms-upgrade
# 또는 선택적 설치:
# sudo amdgpu-install --usecase=hiplibsdk  # HIP SDK만
# sudo amdgpu-install --usecase=mlsdk      # ML 프레임워크 포함

# 4. 사용자 그룹 추가 (필수! 로그아웃 후 재로그인 필요)
sudo usermod -aG render,video $USER
# 변경 즉시 적용 (재로그인 없이):
newgrp render

# 5. 환경 변수 설정 (~/.bashrc 또는 /etc/environment에 추가)
echo 'export PATH=$PATH:/opt/rocm/bin:/opt/rocm/hip/bin' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/opt/rocm/lib' >> ~/.bashrc
echo 'export ROCM_PATH=/opt/rocm' >> ~/.bashrc
source ~/.bashrc

# 6. DKMS 상태 확인 (커널 모듈 빌드 완료 여부)
dkms status
# 예상 출력: amdgpu/6.7.0, 6.5.0-45-generic/x86_64: installed

# 7. 설치 검증
rocminfo           # GPU 목록, 아키텍처, 메모리 용량 확인
rocm-smi           # 실시간 GPU 상태 (클럭, 온도, 메모리)
hipcc --version    # HIP 컴파일러 버전 확인
clinfo             # OpenCL 장치 목록 (rocm-opencl 설치 시)

##############################################################
# RHEL 9 / Rocky Linux 9 / AlmaLinux 9 - ROCm 6.x 설치
##############################################################

# 1. 사전 조건
sudo dnf install -y kernel-headers kernel-devel gcc make
sudo dnf install -y epel-release
sudo crb enable  # CodeReady Builder (RHEL9) 또는 PowerTools (Rocky8)

# 2. AMD ROCm 리포지터리 추가
sudo tee /etc/yum.repos.d/rocm.repo << 'EOF'
[ROCm-6.3]
name=ROCm 6.3
baseurl=https://repo.radeon.com/rocm/rhel9/6.3/main
enabled=1
gpgcheck=1
gpgkey=https://repo.radeon.com/rocm/rocm.gpg.key
EOF

# amdgpu 드라이버 리포지터리
sudo tee /etc/yum.repos.d/amdgpu.repo << 'EOF'
[amdgpu]
name=amdgpu
baseurl=https://repo.radeon.com/amdgpu/6.3/rhel/9.4/main/x86_64/
enabled=1
gpgcheck=1
gpgkey=https://repo.radeon.com/rocm/rocm.gpg.key
EOF

# 3. 드라이버 + ROCm 설치
sudo dnf install -y amdgpu-dkms
sudo dnf install -y rocm-hip-sdk rocm-opencl-sdk

# 4. 사용자 그룹 추가
sudo usermod -aG render,video $USER

# 5. SELinux 정책 (필요 시)
# ROCm이 SELinux enforcing 모드에서 동작하려면 avcstat로 거부 확인:
# sudo ausearch -m avc -ts recent | audit2allow -M rocm_policy
# sudo semodule -i rocm_policy.pp

# 6. 재부팅 후 검증
sudo reboot
rocminfo
rocm-smi --showproductname

##############################################################
# 설치 후 검증 체크리스트
##############################################################

# [1] GPU 인식 여부
rocminfo | grep -A3 "Name:"
# 예: Name: gfx942 (MI300X), gfx1100 (RX 7900 XTX)

# [2] /dev/kfd 존재 및 권한
ls -la /dev/kfd /dev/dri/render*
# crw-rw---- 1 root render 0 Mar 11 /dev/kfd
# crw-rw---- 1 root render /dev/dri/renderD128

# [3] 현재 사용자 그룹 확인
groups | grep -E "render|video"

# [4] HIP Hello World 컴파일 및 실행
cat > /tmp/hip_test.cpp << 'EOF'
#include <hip/hip_runtime.h>
#include <cstdio>
__global__ void hello() { printf("GPU thread %d
", threadIdx.x); }
int main() {
  hello<<<1, 4>>>();
  hipDeviceSynchronize();
}
EOF
hipcc -o /tmp/hip_test /tmp/hip_test.cpp
/tmp/hip_test  # 예상 출력: GPU thread 0~3

# [5] rocBLAS 동작 확인
rocblas-bench -f gemm --transpA N --transpB N -m 4096 -n 4096 -k 4096   --a_type f16_r --b_type f16_r --c_type f32_r --d_type f32_r   --compute_type f32_r
# 성능 수치(TFLOPS) 출력되면 정상

# [6] PyTorch ROCm 동작 확인
python3 -c "
import torch
print('PyTorch 버전:', torch.__version__)
print('CUDA(ROCm) 사용 가능:', torch.cuda.is_available())
print('GPU 장치 이름:', torch.cuda.get_device_name(0))
x = torch.randn(1000, 1000, device='cuda')
y = torch.mm(x, x)
print('행렬 곱 성공:', y.shape)
"

# [7] 메모리 대역폭 벤치마크
rocm-bandwidth-test  # H2D, D2H, D2D 대역폭 측정

# [8] GPU 상태 종합 확인
rocm-smi --showallinfo

관련 문서

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

• GPU 서브시스템 (DRM/KMS) — amdgpu DRM 드라이버, GEM/TTM 메모리, DMA-BUF 버퍼 공유
• DRM GPU 스케줄러 — GPU 작업 큐, fence 기반 동기화, hang 복구
• HMM (Heterogeneous Memory Management) — CPU-GPU 통합 가상 메모리, SVM 페이지 폴트
• NPU (Neural Processing Unit) — AI 가속기 드라이버 모델, DRM Accel 서브시스템
• PCI / PCIe — PCIe 버스 아키텍처, MSI 인터럽트, BAR 매핑
• DMA — DMA 매핑과 전송, IOMMU, 메모리 주소 변환
• 디바이스 드라이버 — 리눅스 드라이버 모델, platform/PCI 드라이버
• 전원 관리 — runtime PM, DVFS, GPU 전력 관리