GPU 서브시스템 (DRM/KMS)

Linux GPU 서브시스템(DRM/KMS)을 디스플레이 파이프라인과 연산 가속 경로를 함께 보는 관점에서 심층 분석합니다. DRM core와 KMS 객체 모델(CRTC/plane/encoder/connector), GEM/TTM 메모리 관리자, DMA-BUF 기반 장치 간 버퍼 공유, atomic modesetting과 vblank 동기화, fence 기반 GPU 작업 순서 제어, 전원·클럭·열 관리, userspace(Mesa/Wayland)와의 ioctl 계약, debugfs/tracepoint로 프레임 드롭과 hang을 진단하는 방법까지 GPU 드라이버 개발 핵심을 다룹니다.

DRM (Direct Rendering Manager) 개요

DRM은 Linux 커널의 GPU 접근을 관리하는 서브시스템입니다. 원래 3D 그래픽 가속을 위해 도입되었으나, 현재는 디스플레이 출력(KMS), GPU 메모리 관리(GEM/TTM), GPU 작업 스케줄링까지 포괄하는 핵심 프레임워크로 발전했습니다.

디바이스 노드와 권한 모델

현대 DRM UAPI는 GPU 하나를 단일 문자 디바이스로만 노출하지 않습니다. 같은 하드웨어라도 화면 제어, 비특권 렌더링, 비그래픽 연산을 서로 다른 노드로 나누어 권한 경계와 사용자 공간 스택을 분리합니다. 이 구분을 이해하지 못하면 drm_file의 인증 상태, client capability, 버퍼 공유 규칙, compositor와 compute runtime의 역할이 한꺼번에 흐려집니다.

/* 유저 공간은 node 종류와 client capability를 먼저 협상한다 */
int fd = open("/dev/dri/card0", O_RDWR | O_CLOEXEC);
drmSetClientCap(fd, DRM_CLIENT_CAP_UNIVERSAL_PLANES, 1);
drmSetClientCap(fd, DRM_CLIENT_CAP_ATOMIC, 1);
drmSetClientCap(fd, DRM_CLIENT_CAP_WRITEBACK_CONNECTORS, 1);

/* para-virtual GPU에서는 커서 plane hotspot capability가 추가로 필요할 수 있다 */
drmSetClientCap(fd, DRM_CLIENT_CAP_CURSOR_PLANE_HOTSPOT, 1);

DRM 아키텍처

유저 공간 그래픽 스택

DRM/KMS 위에서 동작하는 유저 공간 소프트웨어 스택의 구조입니다. 각 계층이 DRM을 어떻게 활용하는지 이해하면 전체 그래픽 파이프라인을 파악할 수 있습니다.

DRM 핵심 구조체

drm_device & drm_driver

drm_device는 하나의 GPU 디바이스를 나타내고, drm_driver는 해당 GPU 유형의 오퍼레이션을 정의합니다.

#include <drm/drm_drv.h>
#include <drm/drm_gem.h>
#include <drm/drm_ioctl.h>

/* DRM 드라이버 정의 — 드라이버 전체에 하나 */
static const struct drm_driver my_gpu_driver = {
    .driver_features    = DRIVER_GEM | DRIVER_MODESET | DRIVER_ATOMIC |
                          DRIVER_RENDER,
    /* GEM 오퍼레이션 */
    .gem_prime_import   = drm_gem_prime_import,
    .gem_prime_import_sg_table = drm_gem_shmem_prime_import_sg_table,

    /* DRM ioctl 테이블 (드라이버 고유 ioctl) */
    .ioctls             = my_gpu_ioctls,
    .num_ioctls         = ARRAY_SIZE(my_gpu_ioctls),

    /* /sys, /dev 노드 이름 */
    .name               = "my-gpu",
    .desc               = "My GPU DRM driver",
    .date               = "20260101",
    .major              = 1,
    .minor              = 0,

    .fops               = &my_gpu_fops,
};

drm_file (파일 프라이빗)

유저 프로세스가 DRM/accel 노드를 open하면 drm_file 구조체가 생성됩니다. 이 구조체는 “프로세스별 GPU 컨텍스트”라기보다 파일 디스크립터별 세션 상태에 가깝고, 인증 여부, master 여부, client capability, 핸들 테이블, 이벤트 큐를 함께 들고 있습니다.

/* drm_file 핵심 필드 발췌 (include/drm/drm_file.h) */
struct drm_file {
    bool               authenticated; /* 레거시 primary 노드 렌더 권한 */
    bool               universal_planes;
    bool               atomic;
    bool               writeback_connectors;
    bool               is_master;
    bool               supports_virtualized_cursor_plane;
    struct drm_minor   *minor;        /* primary/render/accel 중 어떤 노드를 열었는지 */
    struct idr         object_idr;    /* GEM 핸들 → drm_gem_object 매핑 */
    struct idr         syncobj_idr;   /* drm_syncobj 핸들 테이블 */
    struct list_head   pending_event_list;
    struct list_head   event_list;    /* vblank, page flip, out-fence 이벤트 */
    void              *driver_priv;   /* 드라이버별 프라이빗 데이터 */
};

KMS (Kernel Mode Setting)

KMS는 커널에서 디스플레이 모드를 설정하는 프레임워크입니다. 기존의 유저 공간 모드 설정(UMS)과 달리, 커널이 직접 해상도/주사율/디스플레이 파이프라인을 제어하여 콘솔↔그래픽 전환 시 깜박임이 없고 다중 모니터를 안정적으로 관리합니다.

KMS 핵심 오브젝트

Vblank 처리와 페이지 플립

Vblank(수직 귀선 기간)은 디스플레이가 한 프레임을 마치고 다음 프레임을 시작하기 직전의 시간 구간입니다. 이 시점에 프레임버퍼를 교체(page flip)해야 화면 찢어짐(tearing)이 발생하지 않습니다. DRM 코어는 drm_vblank 인프라로 정확한 vblank 타이밍과 시퀀스 번호를 관리합니다.

/* Vblank 관련 드라이버 콜백 */
static const struct drm_crtc_funcs my_crtc_funcs = {
    /* vblank 카운터/타임스탬프 제공 */
    .get_vblank_counter  = drm_crtc_vblank_count,
    .enable_vblank       = my_enable_vblank,   /* vblank IRQ 활성화 */
    .disable_vblank      = my_disable_vblank,  /* vblank IRQ 비활성화 */
    .get_vblank_timestamp = drm_crtc_vblank_helper_get_vblank_timestamp,
};

/* 드라이버 vblank IRQ 핸들러 */
static irqreturn_t my_irq_handler(int irq, void *data)
{
    struct drm_device *dev = data;
    u32 status = readl(regs + IRQ_STATUS);

    if (status & VBLANK_IRQ) {
        drm_crtc_handle_vblank(&my_crtc);  /* vblank 카운터 증가 + 이벤트 전달 */
    }
    return IRQ_HANDLED;
}

/* 페이지 플립 완료 알림 (vblank IRQ에서 호출) */
if (page_flip_pending) {
    drm_crtc_send_vblank_event(&my_crtc, event);
    drm_crtc_vblank_put(&my_crtc);
    page_flip_pending = false;
}

Atomic Modesetting

Atomic modesetting은 KMS의 현대적 커밋 모델입니다. 여러 디스플레이 속성 변경을 하나의 원자적 트랜잭션으로 묶어 중간 상태 없이 적용하거나 전체를 롤백합니다.

Atomic 상태 객체 그래프

실제 atomic UAPI의 핵심은 “property를 바로 레지스터에 쓰는 것”이 아니라, drm_atomic_state 아래에 CRTC/plane/connector/private object의 새 상태를 모아 둔 뒤 한 번에 검증하고 커밋하는 것입니다. driver의 atomic_check는 이 그래프 전체를 보고 대역폭, scaler 수, watermark, 링크 훈련 가능 여부, color LUT 크기, writeback 출력 버퍼 충돌까지 판단합니다.

/* Atomic 커밋 흐름 (커널 내부) */

/* 1. 유저가 DRM_IOCTL_MODE_ATOMIC 호출 */
/*    → drm_mode_atomic_ioctl() 진입 */

/* 2. 상태 복제 */
struct drm_atomic_state *state = drm_atomic_state_alloc(dev);
/* 각 오브젝트의 현재 상태를 복제하여 new_state 생성 */

/* 3. 속성 적용 (유저가 요청한 변경사항) */
drm_atomic_set_crtc_for_connector(new_conn_state, crtc);
new_crtc_state->mode_blob = mode;
new_plane_state->fb = framebuffer;
new_plane_state->crtc = crtc;

/* 4. 검증 (atomic_check) — 하드웨어 제약 확인 */
ret = drm_atomic_check_only(state);
/* 각 CRTC/Plane/Connector의 atomic_check 콜백 호출 */
/* 실패 시 -EINVAL 반환, 유저에게 거부 통보 */

/* 5. TEST_ONLY 플래그면 여기서 종료 (검증만 수행) */
if (flags & DRM_MODE_ATOMIC_TEST_ONLY)
    return ret;

/* 6. 커밋 (atomic_commit) — 실제 하드웨어 적용 */
ret = drm_atomic_commit(state);
/* NONBLOCK 플래그: 비동기 커밋 (vsync 대기 안 함) */
/* PAGE_FLIP_EVENT: 완료 시 유저에게 이벤트 전달 */

DRM Properties 시스템

KMS 오브젝트(CRTC, Plane, Connector)의 모든 설정 가능한 속성은 DRM Property로 표현됩니다. Atomic modesetting에서는 property ID와 값을 쌍으로 전달하여 디스플레이 상태를 변경합니다.

Property 타입

표준 KMS Property

/* === CRTC Properties === */
"ACTIVE"           /* bool: CRTC 활성화 (atomic에서 DPMS 대체) */
"MODE_ID"          /* blob: 디스플레이 모드 (해상도/주사율) */
"OUT_FENCE_PTR"    /* ptr: 커밋 완료 fence fd 반환 위치 */
"VRR_ENABLED"      /* bool: 가변 주사율 활성화 */
"DEGAMMA_LUT"      /* blob: 디감마 LUT (Color Management) */
"CTM"              /* blob: 색 변환 매트릭스 (3x3, S31.32 고정소수) */
"GAMMA_LUT"        /* blob: 감마 LUT */

/* === Plane Properties === */
"FB_ID"            /* object: 표시할 프레임버퍼 */
"CRTC_ID"         /* object: 연결할 CRTC */
"SRC_X/Y/W/H"     /* range: 소스 영역 (16.16 고정소수점) */
"CRTC_X/Y/W/H"    /* range: 화면 위 위치/크기 (스케일링) */
"rotation"         /* bitmask: ROTATE_0/90/180/270 | REFLECT_X/Y */
"alpha"            /* range: 투명도 (0=투명, 0xFFFF=불투명) */
"pixel blend mode" /* enum: None/Pre-multiplied/Coverage */
"zpos"             /* range: 레이어 순서 (높을수록 위) */
"IN_FENCE_FD"      /* range: explicit sync 입력 fence fd */
"COLOR_ENCODING"   /* enum: YCbCr BT.601/709/2020 */
"COLOR_RANGE"      /* enum: YCbCr Limited/Full Range */

/* === Connector Properties === */
"CRTC_ID"          /* object: 연결할 CRTC */
"DPMS"             /* enum: On/Standby/Suspend/Off (레거시) */
"link-status"      /* enum: Good/Bad — 링크 학습 실패 시 Bad */
"EDID"             /* blob: 모니터 EDID 바이너리 (읽기 전용) */
"content type"     /* enum: No Data/Graphics/Photo/Cinema/Game */
"max bpc"          /* range: 최대 색 깊이 (bits per component) */
"HDR_OUTPUT_METADATA" /* blob: HDR 정적/동적 메타데이터 */
"vrr_capable"      /* range: 모니터 VRR 지원 여부 (읽기 전용) */

/* 드라이버에서 커스텀 property 등록 */
struct drm_property *prop;

/* Enum property 생성 */
static const struct drm_prop_enum_list my_scaling_modes[] = {
    { 0, "None" },
    { 1, "Full" },
    { 2, "Center" },
    { 3, "Full aspect" },
};
prop = drm_property_create_enum(dev, 0, "scaling mode",
        my_scaling_modes, ARRAY_SIZE(my_scaling_modes));

/* CRTC에 property 연결 */
drm_object_attach_property(&crtc->base, prop, 0);

/* atomic_check에서 property 값 읽기 */
static int my_crtc_atomic_check(struct drm_crtc *crtc,
                                struct drm_atomic_state *state)
{
    struct drm_crtc_state *crtc_state =
        drm_atomic_get_new_crtc_state(state, crtc);
    /* crtc_state에서 property 값에 접근 */
    ...
}

Color Management 파이프라인

KMS는 CRTC 레벨에서 3단계 색상 처리 파이프라인을 제공합니다. 이 파이프라인으로 HDR 톤매핑, 색 공간 변환, 감마 보정 등을 하드웨어 가속으로 수행합니다.

/* Color Management 관련 구조체 */

/* 감마/디감마 LUT 엔트리 (유저 공간에서 설정) */
struct drm_color_lut {
    __u16 red;      /* 0~0xFFFF */
    __u16 green;
    __u16 blue;
    __u16 reserved;
};

/* 색 변환 매트릭스 (CTM) — S31.32 고정소수점 */
struct drm_color_ctm {
    __u64 matrix[9];   /* 3x3, 부호 비트 + 31.32 고정소수 */
};

/* 유저 공간에서 감마 LUT 설정 예시 (libdrm) */
struct drm_color_lut lut[256];
for (int i = 0; i < 256; i++) {
    /* sRGB 감마 커브: 선형 → sRGB 비선형 */
    double v = (double)i / 255.0;
    double srgb = (v <= 0.0031308) ?
        v * 12.92 : 1.055 * pow(v, 1.0/2.4) - 0.055;
    __u16 val = (__u16)(srgb * 0xFFFF);
    lut[i].red = lut[i].green = lut[i].blue = val;
}
/* blob property로 설정 */
drmModeCreatePropertyBlob(fd, lut, sizeof(lut), &blob_id);
/* atomic 커밋에 GAMMA_LUT=blob_id 추가 */

VRR (Variable Refresh Rate) / Adaptive Sync

VRR은 GPU 렌더링 속도에 맞춰 디스플레이 주사율을 동적으로 조절하는 기술입니다. 화면 찢어짐(tearing)과 스터터링(stuttering) 없이 부드러운 프레임 전달을 가능하게 합니다.

/* VRR 활성화 흐름 (커널 내부) */

/* 1. Connector가 VRR 지원하는지 확인 */
/*    모니터 EDID의 Adaptive-Sync range 파싱 */
connector->vrr_capable = true;  /* 드라이버가 EDID 기반으로 설정 */

/* 2. 유저 공간에서 CRTC의 VRR_ENABLED property 설정 */
/*    → atomic commit에 포함 */
new_crtc_state->vrr_enabled = true;

/* 3. 드라이버 atomic_check에서 VRR 설정 검증 */
static int my_crtc_atomic_check(struct drm_crtc *crtc,
                                struct drm_atomic_state *state)
{
    struct drm_crtc_state *crtc_state =
        drm_atomic_get_new_crtc_state(state, crtc);

    if (crtc_state->vrr_enabled) {
        /* VRR 범위 확인 (min_vfreq ~ max_vfreq) */
        struct drm_connector *conn = get_connector(crtc_state);
        if (!conn->display_info.monitor_range.max_vfreq)
            return -EINVAL;

        /* 하드웨어 VRR 타이밍 설정 */
        my_state->vmin = conn->display_info.monitor_range.min_vfreq;
        my_state->vmax = conn->display_info.monitor_range.max_vfreq;
    }
    return 0;
}

/* 4. 페이지 플립 시 VRR 타이밍 적용 */
/*    GPU가 프레임을 완료하면 즉시 vsync 발생 */
/*    → 가변 vblank 주기 */

Format Modifier (타일링/압축)

GPU는 메모리 접근 효율을 위해 픽셀 데이터를 단순 행(linear) 방식이 아닌 타일(tile) 형태로 배치합니다. Format Modifier는 프레임버퍼의 메모리 레이아웃을 기술하는 64비트 토큰입니다.

/* Format Modifier 구조 (include/uapi/drm/drm_fourcc.h) */
/*
 * [63:56] = vendor (DRM_FORMAT_MOD_VENDOR_*)
 * [55:0]  = vendor별 modifier 값
 *
 * 0 = DRM_FORMAT_MOD_LINEAR (행 순차 배치, 모든 디바이스 공통)
 */

#define DRM_FORMAT_MOD_LINEAR  0

/* Intel 타일링 modifier 예시 */
#define I915_FORMAT_MOD_X_TILED    /* X-tiling (레거시) */
#define I915_FORMAT_MOD_Y_TILED    /* Y-tiling (Gen9+) */
#define I915_FORMAT_MOD_4_TILED    /* Tile-4 (Gen12.5+, DG2/MTL) */
#define I915_FORMAT_MOD_Y_TILED_CCS /* Y-tiling + CCS 압축 */

/* AMD modifier 예시 */
#define AMD_FMT_MOD                /* 타일 버전, pipe_xor_bits, DCC 등 인코딩 */

/* ARM AFBC (Arm Frame Buffer Compression) */
#define DRM_FORMAT_MOD_ARM_AFBC    /* Mali GPU 프레임버퍼 무손실 압축 */

/* Broadcom */
#define DRM_FORMAT_MOD_BROADCOM_VC4_T_TILED /* VideoCore T-tiling */

/* NVIDIA */
#define DRM_FORMAT_MOD_NVIDIA_BLOCK_LINEAR_2D /* GOB 기반 타일링 */

/* 프레임버퍼 생성 시 modifier 지정 (ADDFB2) */
struct drm_mode_fb_cmd2 cmd = {
    .width   = 1920,
    .height  = 1080,
    .pixel_format = DRM_FORMAT_XRGB8888,
    .handles[0] = gem_handle,
    .pitches[0] = pitch,
    .modifier[0] = I915_FORMAT_MOD_Y_TILED, /* modifier 지정 */
    .flags   = DRM_MODE_FB_MODIFIERS,       /* modifier 사용 플래그 */
};
drmIoctl(fd, DRM_IOCTL_MODE_ADDFB2, &cmd);

/* Plane이 지원하는 modifier 목록 조회 */
/* DRM_IOCTL_MODE_GETPLANE2 또는 IN_FORMATS blob property */
/* → (format, modifier) 쌍의 목록 반환 */

/* 드라이버: Plane에 지원 modifier 등록 */
static const uint64_t my_modifiers[] = {
    DRM_FORMAT_MOD_LINEAR,
    I915_FORMAT_MOD_Y_TILED,
    I915_FORMAT_MOD_Y_TILED_CCS,
    DRM_FORMAT_MOD_INVALID  /* 종료 마커 */
};
drm_universal_plane_init(dev, plane, ...,
    formats, num_formats, my_modifiers, ...);

GEM (Graphics Execution Manager) 메모리 관리

GEM은 GPU 메모리 버퍼 객체를 관리하는 프레임워크입니다. 유저 공간에서는 정수 핸들로 버퍼를 참조하고, 커널은 내부적으로 drm_gem_object로 관리합니다.

/* drm_gem_object 핵심 필드 발췌 (include/drm/drm_gem.h) */
struct drm_gem_object {
    struct kref           refcount;      /* 참조 카운팅 */
    unsigned             handle_count;  /* 유저 핸들 수 */
    struct drm_device     *dev;          /* 소속 디바이스 */
    struct file           *filp;         /* shmem backing store 또는 NULL */
    struct drm_vma_offset_node vma_node; /* mmap 오프셋 관리 */
    size_t                size;          /* 버퍼 크기 */
    int                   name;          /* flink 이름 (레거시) */
    struct dma_buf        *dma_buf;      /* PRIME export/import 연결점 */
    struct dma_buf_attachment *import_attach;
    struct dma_resv       *resv;         /* 공유 fence 컨테이너 */
    const struct drm_gem_object_funcs *funcs;
};

/* GEM 오퍼레이션 콜백 */
struct drm_gem_object_funcs {
    void (*free)(struct drm_gem_object *obj);
    int  (*open)(struct drm_gem_object *obj, struct drm_file *file);
    void (*close)(struct drm_gem_object *obj, struct drm_file *file);
    int  (*pin)(struct drm_gem_object *obj);
    int  (*vmap)(struct drm_gem_object *obj, struct iosys_map *map);
    struct sg_table *(*get_sg_table)(struct drm_gem_object *obj);
};

GEM 버퍼 생명주기

GEM SHMEM 헬퍼

대부분의 임베디드/모바일 GPU 드라이버는 drm_gem_shmem_object를 사용합니다. 시스템 메모리(shmem)에 버퍼를 할당하며, CMA(Contiguous Memory Allocator) 통합도 지원합니다.

#include <drm/drm_gem_shmem_helper.h>

/* GEM SHMEM 기반 드라이버의 dumb_create 구현 */
static int my_dumb_create(struct drm_file *file,
                           struct drm_device *dev,
                           struct drm_mode_create_dumb *args)
{
    /* pitch 정렬 (하드웨어 요구사항에 맞춤) */
    args->pitch = ALIGN(args->width * DIV_ROUND_UP(args->bpp, 8), 64);
    args->size  = args->pitch * args->height;

    /* shmem 헬퍼가 할당 + 핸들 반환 */
    return drm_gem_shmem_dumb_create(file, dev, args);
}

/* 드라이버에서 vmap으로 커널 가상 주소 획득 */
struct drm_gem_shmem_object *shmem = to_drm_gem_shmem_obj(gem_obj);
struct iosys_map map;
drm_gem_shmem_vmap(shmem, &map);
/* map.vaddr로 커널에서 버퍼 접근 */

TTM (Translation Table Manager)

TTM은 전용 VRAM(비디오 메모리)이 있는 GPU를 위한 메모리 관리자입니다. 버퍼를 VRAM과 시스템 메모리 사이에서 이동(eviction/migration)시키는 고급 기능을 제공합니다.

/* TTM 메모리 도메인 기술자 */
struct ttm_place {
    uint32_t  mem_type;   /* TTM_PL_SYSTEM, TTM_PL_VRAM, TTM_PL_TT */
    uint32_t  flags;      /* TTM_PL_FLAG_CONTIGUOUS 등 */
    uint64_t  fpfn;       /* 시작 페이지 프레임 번호 */
    uint64_t  lpfn;       /* 끝 페이지 프레임 번호 (0=제한 없음) */
};

/* ttm_buffer_object 핵심 필드 발췌 */
struct ttm_buffer_object {
    struct drm_gem_object    base;         /* GEM 상속 */
    struct ttm_device        *bdev;         /* TTM 디바이스 */
    enum ttm_bo_type        type;
    uint32_t                 page_alignment;
    void (*destroy)(struct ttm_buffer_object *bo);
    struct kref              kref;
    struct ttm_resource      *resource;     /* 현재 메모리 위치 */
    struct ttm_tt            *ttm;          /* 시스템 메모리 페이지 */
    bool                     deleted;
    unsigned                 priority;
    unsigned                 pin_count;
    struct sg_table         *sg;
};

/* TTM 메모리 도메인 상수 */
#define TTM_PL_SYSTEM  0  /* 시스템 RAM (커널 관리) */
#define TTM_PL_TT      1  /* GART/GTT 매핑된 시스템 메모리 */
#define TTM_PL_VRAM    2  /* 디바이스 전용 비디오 메모리 */

DMA-BUF (버퍼 공유)

DMA-BUF는 디바이스 간 버퍼 공유를 위한 커널 프레임워크입니다. GPU에서 렌더링한 버퍼를 디스플레이 컨트롤러, 비디오 인코더, 카메라 등 다른 디바이스와 복사 없이 공유할 수 있습니다.

/* DMA-BUF exporter 오퍼레이션 (드라이버 구현) */
static const struct dma_buf_ops my_dmabuf_ops = {
    .attach         = my_dmabuf_attach,       /* importer 디바이스 등록 */
    .detach         = my_dmabuf_detach,
    .map_dma_buf    = my_dmabuf_map,          /* sg_table 반환 (DMA 주소) */
    .unmap_dma_buf  = my_dmabuf_unmap,
    .release        = my_dmabuf_release,      /* 버퍼 해제 */
    .mmap           = my_dmabuf_mmap,         /* 유저 공간 매핑 */
    .vmap           = my_dmabuf_vmap,         /* 커널 가상 주소 매핑 */
};

/* GEM → DMA-BUF export (PRIME) */
struct dma_buf *dmabuf = drm_gem_prime_export(gem_obj, flags);
/* → 유저 공간에서 fd = DRM_IOCTL_PRIME_HANDLE_TO_FD(handle) */

/* DMA-BUF → GEM import (PRIME) */
struct drm_gem_object *obj = drm_gem_prime_import(dev, dmabuf);
/* → 유저 공간에서 handle = DRM_IOCTL_PRIME_FD_TO_HANDLE(fd) */

dma_fence & dma_resv (동기화)

GPU 작업은 비동기로 실행되므로, 버퍼를 공유하는 디바이스들은 작업 완료를 동기화해야 합니다.

/* Fence 기본 사용 패턴 */
struct dma_fence *fence;

/* GPU 작업 제출 시 fence 생성 */
fence = my_gpu_submit_job(job);

/* 버퍼의 reservation object에 fence 등록 */
dma_resv_add_fence(bo->base.resv, fence, DMA_RESV_USAGE_WRITE);

/* 다른 디바이스가 이 버퍼를 사용하기 전에 대기 */
ret = dma_resv_wait_timeout(bo->base.resv, DMA_RESV_USAGE_WRITE,
                            true, msecs_to_jiffies(5000));

/* Fence 시그널 (GPU 인터럽트 핸들러에서 호출) */
dma_fence_signal(fence);
dma_fence_put(fence);

Implicit Sync vs Explicit Sync

GPU 동기화 모델은 두 가지로 나뉩니다. Linux DRM은 전통적으로 implicit sync를 사용했으나, Vulkan과 Android를 중심으로 explicit sync로 전환하고 있습니다.

/* Explicit Sync: KMS IN/OUT fence 사용 */

/* 유저 공간: GPU 렌더링 완료 fence fd를 Plane에 전달 */
drmModeAtomicAddProperty(req, plane_id, IN_FENCE_FD_prop, gpu_fence_fd);

/* 유저 공간: 커밋 완료 fence fd를 받을 위치 지정 */
int64_t out_fence_fd = -1;
drmModeAtomicAddProperty(req, crtc_id, OUT_FENCE_PTR_prop,
                         (uint64_t)(uintptr_t)&out_fence_fd);

/* Atomic 커밋 */
drmModeAtomicCommit(fd, req, DRM_MODE_ATOMIC_NONBLOCK |
                             DRM_MODE_PAGE_FLIP_EVENT, NULL);
/* out_fence_fd에 유효한 fd가 설정됨
 * → 다음 렌더링 시작 전에 이 fence를 GPU에 제출하여 대기
 * → 완전한 파이프라인 동기화 */

GPU 커맨드 서브미션

GPU 렌더링 작업은 커맨드 버퍼(command buffer)를 통해 제출됩니다. 유저 공간(Mesa/Vulkan 드라이버)이 GPU 명령어를 버퍼에 기록하고, 커널 드라이버가 이를 GPU 하드웨어에 전달합니다.

서브미션 모델 비교

/* amdgpu 커맨드 서브미션 예시 (유저 공간, libdrm_amdgpu) */
struct amdgpu_cs_request cs_req = {};
struct amdgpu_cs_ib_info ib_info = {
    .ib_mc_address = ib_gpu_addr,  /* IB의 GPU 가상 주소 */
    .size          = ib_size_dw,   /* 더블워드 단위 크기 */
};
cs_req.ip_type       = AMDGPU_HW_IP_GFX;   /* GFX 엔진 */
cs_req.number_of_ibs = 1;
cs_req.ibs           = &ib_info;

/* 의존 fence 목록 (이전 작업 완료 대기) */
struct amdgpu_cs_fence_info fence_info = { ... };
cs_req.fence_info = fence_info;

/* 서브미션 */
amdgpu_cs_submit(ctx, 0, &cs_req, 1);

/* 완료 대기 (선택적) */
struct amdgpu_cs_fence fence = {
    .context = ctx,
    .ip_type = AMDGPU_HW_IP_GFX,
    .fence   = cs_req.seq_no,
};
amdgpu_cs_query_fence_status(&fence, timeout_ns, 0, &expired);

GPU 스케줄러 (drm_sched)

DRM GPU 스케줄러(drivers/gpu/drm/scheduler/)는 여러 프로세스의 GPU 작업을 공정하게 스케줄링하고, GPU 행(hang) 감지 및 복구를 처리합니다.

#include <drm/gpu_scheduler.h>

/* 스케줄러 초기화 (드라이버 probe에서) */
struct drm_gpu_scheduler sched;

drm_sched_init(&sched,
    &my_sched_ops,              /* 드라이버 콜백 */
    NULL,                       /* workqueue (NULL=기본) */
    num_hw_rings,               /* 하드웨어 큐 수 */
    0,                          /* credit limit */
    msecs_to_jiffies(10000),   /* 타임아웃 (10초) */
    NULL,                       /* workqueue for timeout */
    NULL,                       /* score */
    "my-gpu",                   /* 이름 */
    dev->dev);

/* 스케줄러 오퍼레이션 */
static const struct drm_sched_backend_ops my_sched_ops = {
    .run_job    = my_run_job,        /* GPU에 작업 제출 */
    .timedout_job = my_timedout_job, /* 타임아웃 시 GPU 리셋 */
    .free_job   = my_free_job,       /* 완료된 작업 해제 */
};

/* run_job 구현 예시 */
static struct dma_fence *my_run_job(struct drm_sched_job *job)
{
    struct my_job *mj = to_my_job(job);

    /* GPU 하드웨어 레지스터에 커맨드 버퍼 주소 기록 */
    writel(mj->cmd_buf_addr, gpu->regs + CMD_SUBMIT);

    /* 완료 fence 반환 (GPU 인터럽트에서 시그널) */
    return dma_fence_get(mj->fence);
}

DRM 드라이버 기본 구조

최소한의 KMS DRM 드라이버 (디스플레이만 지원, GPU 가속 없음)의 골격입니다.

#include <linux/module.h>
#include <linux/platform_device.h>
#include <drm/drm_drv.h>
#include <drm/drm_gem_shmem_helper.h>
#include <drm/drm_atomic_helper.h>
#include <drm/drm_simple_kms_helper.h>
#include <drm/drm_fbdev_shmem.h>

struct my_drm {
    struct drm_device              drm;
    struct drm_simple_display_pipe pipe;
    struct drm_connector           connector;
    void __iomem                  *regs;
};

/* --- Connector --- */
static int my_connector_get_modes(struct drm_connector *conn)
{
    /* 고정 해상도 모드 추가 (실제로는 EDID 파싱) */
    return drm_add_modes_noedid(conn, 1920, 1080);
}

static const struct drm_connector_helper_funcs my_conn_helpers = {
    .get_modes = my_connector_get_modes,
};

/* --- Simple Display Pipe --- */
static void my_pipe_enable(struct drm_simple_display_pipe *pipe,
                           struct drm_crtc_state *crtc_state,
                           struct drm_plane_state *plane_state)
{
    struct my_drm *my = container_of(pipe, struct my_drm, pipe);
    struct drm_framebuffer *fb = plane_state->fb;

    /* 하드웨어 레지스터에 프레임버퍼 주소/포맷 설정 */
    writel(drm_fb_dma_get_gem_addr(fb, plane_state, 0),
           my->regs + FB_ADDR_REG);
    writel(fb->pitches[0], my->regs + FB_PITCH_REG);
    /* 디스플레이 엔진 활성화 */
    writel(1, my->regs + DISPLAY_ENABLE_REG);
}

static void my_pipe_update(struct drm_simple_display_pipe *pipe,
                           struct drm_plane_state *old_state)
{
    struct drm_plane_state *state = pipe->plane.state;
    if (state->fb)
        my_pipe_enable(pipe, NULL, state);
}

static const struct drm_simple_display_pipe_funcs my_pipe_funcs = {
    .enable  = my_pipe_enable,
    .update  = my_pipe_update,
};

/* --- DRM Driver --- */
DEFINE_DRM_GEM_FOPS(my_fops);

static const struct drm_driver my_driver = {
    .driver_features = DRIVER_GEM | DRIVER_MODESET | DRIVER_ATOMIC,
    .fops            = &my_fops,
    DRM_GEM_SHMEM_DRIVER_OPS,
    .name  = "my-display",
    .desc  = "My simple display driver",
    .date  = "20260101",
    .major = 1,
    .minor = 0,
};

/* --- Platform Driver --- */
static int my_probe(struct platform_device *pdev)
{
    struct my_drm *my;
    struct drm_device *drm;
    static const uint32_t formats[] = { DRM_FORMAT_XRGB8888 };
    int ret;

    my = devm_drm_dev_alloc(&pdev->dev, &my_driver,
                            struct my_drm, drm);
    if (IS_ERR(my))
        return PTR_ERR(my);
    drm = &my->drm;

    /* MMIO 매핑 */
    my->regs = devm_platform_ioremap_resource(pdev, 0);
    if (IS_ERR(my->regs))
        return PTR_ERR(my->regs);

    /* Connector 초기화 */
    drm_connector_init(drm, &my->connector,
                       &my_conn_funcs, DRM_MODE_CONNECTOR_Unknown);
    drm_connector_helper_add(&my->connector, &my_conn_helpers);

    /* Simple display pipe 초기화 (CRTC+Plane+Encoder 한 번에) */
    ret = drm_simple_display_pipe_init(drm, &my->pipe,
            &my_pipe_funcs, formats, ARRAY_SIZE(formats),
            NULL, &my->connector);
    if (ret)
        return ret;

    drm_mode_config_reset(drm);
    ret = drm_dev_register(drm, 0);
    if (ret)
        return ret;

    /* fbdev 에뮬레이션 (콘솔 출력) */
    drm_fbdev_shmem_setup(drm, 32);
    return 0;
}

MODULE_LICENSE("GPL");

주요 오픈소스 DRM 드라이버

최근 DRM 코어의 방향

버전 번호 자체보다 중요한 것은 메인라인이 어떤 설계 축으로 이동하는가입니다. 최근 DRM 문서와 헤더를 보면 다음 네 가지 흐름이 거의 모든 드라이버 설계에 공통으로 반영됩니다.

DRM Bridge & Panel 프레임워크

임베디드/모바일 SoC에서는 디스플레이 출력이 여러 개의 하드웨어 블록 체인으로 구성됩니다. drm_bridge는 이 체인의 중간 단계(MIPI-DSI to HDMI 변환기, LVDS 직렬화기 등)를, drm_panel은 최종 디스플레이 패널을 추상화합니다.

/* DRM Bridge 드라이버 구현 예시 */
#include <drm/drm_bridge.h>

static const struct drm_bridge_funcs my_bridge_funcs = {
    /* 디스플레이 파이프라인 활성화 순서 */
    .pre_enable  = my_bridge_pre_enable,  /* PLL/클럭 설정 */
    .enable      = my_bridge_enable,      /* 출력 활성화 */
    .disable     = my_bridge_disable,
    .post_disable = my_bridge_post_disable,

    /* 모드 검증 */
    .mode_valid  = my_bridge_mode_valid,  /* 지원 해상도 필터링 */
    .mode_fixup  = my_bridge_mode_fixup,  /* adjusted_mode 수정 */

    /* Atomic 지원 (권장) */
    .atomic_pre_enable = my_bridge_atomic_pre_enable,
    .atomic_enable     = my_bridge_atomic_enable,
    .atomic_disable    = my_bridge_atomic_disable,

    /* 출력 감지 (DRM_BRIDGE_OP_DETECT) */
    .detect      = my_bridge_detect,
    .get_modes   = my_bridge_get_modes,   /* EDID 기반 모드 목록 */
    .get_edid    = my_bridge_get_edid,    /* EDID 읽기 */
};

static int my_bridge_probe(struct i2c_client *client)
{
    struct my_bridge *mb = devm_kzalloc(&client->dev, sizeof(*mb), GFP_KERNEL);

    mb->bridge.funcs = &my_bridge_funcs;
    mb->bridge.of_node = client->dev.of_node;
    mb->bridge.type = DRM_MODE_CONNECTOR_HDMIA;  /* 출력 커넥터 타입 */
    mb->bridge.ops = DRM_BRIDGE_OP_DETECT |
                     DRM_BRIDGE_OP_EDID |
                     DRM_BRIDGE_OP_MODES;

    drm_bridge_add(&mb->bridge);
    return 0;
}

/* DRM Panel 드라이버 구현 예시 */
#include <drm/drm_panel.h>

static const struct drm_panel_funcs my_panel_funcs = {
    .prepare    = my_panel_prepare,     /* 전원 ON, 리셋 해제 */
    .enable     = my_panel_enable,      /* 백라이트 ON, 디스플레이 ON */
    .disable    = my_panel_disable,     /* 백라이트 OFF */
    .unprepare  = my_panel_unprepare,   /* 전원 OFF */
    .get_modes  = my_panel_get_modes,   /* 고정 모드 반환 */
};

static int my_panel_probe(struct mipi_dsi_device *dsi)
{
    struct my_panel *panel = devm_kzalloc(&dsi->dev, sizeof(*panel), GFP_KERNEL);

    drm_panel_init(&panel->panel, &dsi->dev, &my_panel_funcs,
                   DRM_MODE_CONNECTOR_DSI);
    drm_panel_add(&panel->panel);
    return 0;
}

/* 디스플레이 드라이버에서 bridge + panel 연결 */
struct drm_bridge *bridge;
bridge = devm_drm_of_get_bridge(dev, dev->of_node, 0, 0);
/* drm_panel은 자동으로 panel-bridge로 래핑됨 */
drm_bridge_attach(encoder, bridge, NULL, DRM_BRIDGE_ATTACH_NO_CONNECTOR);

/* 체인의 마지막 bridge에서 connector 생성 */
connector = drm_bridge_connector_init(drm, encoder);

GPU 가상 메모리 (GPUVM)

현대 GPU는 자체 MMU를 가지며, 프로세스별 독립적인 GPU 가상 주소 공간을 제공합니다. 커널 DRM 프레임워크는 drm_gpuvm(drivers/gpu/drm/drm_gpuvm.c)으로 이를 추상화합니다.

GPU 전원 관리

GPU는 시스템에서 가장 전력을 많이 소비하는 디바이스 중 하나이므로, 정교한 전원 관리가 매우 중요합니다.

# GPU Runtime PM 상태 확인
cat /sys/class/drm/card0/device/power/runtime_status
# active / suspended / suspending

# Runtime PM 자동 절전 지연 설정 (밀리초)
echo 5000 > /sys/class/drm/card0/device/power/autosuspend_delay_ms

# GPU 클럭 주파수 확인 (amdgpu 예시)
cat /sys/class/drm/card0/device/pp_dpm_sclk
# 0: 300Mhz
# 1: 600Mhz *
# 2: 900Mhz

# PRIME GPU 오프로드 (외장 GPU로 렌더링)
DRI_PRIME=1 glxinfo | grep "OpenGL renderer"

# vga_switcheroo 상태 (하이브리드 GPU)
cat /sys/kernel/debug/vgaswitcheroo/switch

GPU 리셋 및 복구

GPU는 잘못된 셰이더, 무한 루프, 하드웨어 결함 등으로 응답 불능(hang) 상태에 빠질 수 있습니다. DRM 프레임워크는 hang 감지, GPU 리셋, 작업 재제출의 3단계 복구 메커니즘을 제공합니다.

Hang 감지

리셋 수준

/* GPU 리셋 수준 (세밀한 것부터 거친 순서) */

/* Level 1: Per-engine 리셋 (권장, 최소 영향) */
/* 특정 GPU 엔진(GFX, SDMA, 비디오 등)만 리셋 */
/* 다른 엔진은 계속 동작 */
amdgpu_device_reset_engine(adev, ring);  /* amdgpu 예시 */

/* Level 2: Per-context 리셋 (하드웨어 지원 필요) */
/* 특정 GPU 컨텍스트만 무효화 */
/* i915: 잘못된 컨텍스트를 ban (금지) 처리 */
intel_context_ban(ce, NULL);

/* Level 3: 전체 GPU 리셋 */
/* 모든 엔진/컨텍스트 중단 후 GPU 완전 재초기화 */
amdgpu_device_gpu_recover(adev, NULL, false);

/* Level 4: FLR (Function Level Reset) - PCIe */
/* PCIe 레벨에서 디바이스 전체 리셋 */
pci_reset_function(pdev);

/* Level 5: BACO (Bus Active, Chip Off) - AMD */
/* GPU 전원을 완전히 차단 후 재투입 */
amdgpu_device_baco_enter(adev_to_drm(adev));
amdgpu_device_baco_exit(adev_to_drm(adev));

복구 흐름

/* drm_sched 타임아웃 → 복구 흐름 */

/* 1. 타임아웃 콜백 호출 */
static enum drm_gpu_sched_stat
my_timedout_job(struct drm_sched_job *job)
{
    struct my_device *dev = job_to_dev(job);

    /* 2. GPU 상태 덤프 (디버깅용) */
    dev_coredump_snapshot(dev);  /* devcoredump 프레임워크 */

    /* 3. GPU 리셋 수행 */
    my_gpu_reset(dev);

    /* 4. 스케줄러에 리셋 완료 통보 */
    /*    → 대기 중인 fence들을 에러로 시그널 */
    /*    → guilty 작업의 entity에 guilty 플래그 설정 */
    return DRM_GPU_SCHED_STAT_NOMINAL;
}

/* 5. 유저 공간 알림 */
/*    - 리셋 후 fence가 에러(-EIO)로 시그널됨 */
/*    - Vulkan: VK_ERROR_DEVICE_LOST 반환 */
/*    - GL: GL_CONTEXT_LOST 또는 glGetGraphicsResetStatusARB() */
/*    - amdgpu: AMDGPU_CTX_QUERY2 ioctl로 리셋 감지 */

# GPU 리셋 관련 sysfs/debugfs

# amdgpu: 수동 GPU 복구 트리거
echo 1 > /sys/kernel/debug/dri/0/amdgpu_gpu_recover

# i915: 엔진 리셋 카운트
cat /sys/kernel/debug/dri/0/i915_reset_info

# devcoredump: GPU 상태 덤프 읽기
cat /sys/class/devcoredump/devcd0/data > gpu_dump.bin

# devcoredump 삭제 (5분 후 자동 삭제)
echo 1 > /sys/class/devcoredump/devcd0/data

# dmesg에서 GPU 리셋 로그 확인
dmesg | grep -i "gpu\|reset\|hang\|timed out"

커널 설정 (Kconfig)

# DRM 핵심 설정
CONFIG_DRM=m                      # DRM 코어 모듈
CONFIG_DRM_KMS_HELPER=m           # KMS 헬퍼 함수
CONFIG_DRM_GEM_SHMEM_HELPER=m     # GEM shmem 헬퍼
CONFIG_DRM_SCHED=m                # GPU 스케줄러
CONFIG_DRM_TTM=m                  # TTM 메모리 관리자
CONFIG_DRM_DISPLAY_HELPER=m       # DP/HDMI 디스플레이 헬퍼

# 드라이버별 설정
CONFIG_DRM_I915=m                 # Intel i915
CONFIG_DRM_XE=m                   # Intel Xe
CONFIG_DRM_AMDGPU=m               # AMD GPU
CONFIG_DRM_NOUVEAU=m              # NVIDIA (오픈소스)
CONFIG_DRM_PANFROST=m             # ARM Mali (Midgard/Bifrost)
CONFIG_DRM_LIMA=m                 # ARM Mali (Utgard)
CONFIG_DRM_VC4=m                  # Broadcom VideoCore
CONFIG_DRM_MSM=m                  # Qualcomm Adreno
CONFIG_DRM_VIRTIO_GPU=m           # 가상 GPU (QEMU)
CONFIG_DRM_SIMPLEDRM=y            # EFI/BIOS 프레임버퍼

# fbdev 에뮬레이션 (DRM 위에서 /dev/fb0 제공)
CONFIG_DRM_FBDEV_EMULATION=y

# DMA-BUF
CONFIG_DMA_SHARED_BUFFER=y

DRM 디버깅

디버깅 도구

커널 디버그 메시지

# DRM 디버그 레벨 설정 (비트마스크)
echo 0x1ff > /sys/module/drm/parameters/debug

# 비트마스크 의미:
#   0x001 = DRM_UT_CORE     — DRM 코어
#   0x002 = DRM_UT_DRIVER   — 드라이버
#   0x004 = DRM_UT_KMS      — 모드 설정
#   0x008 = DRM_UT_PRIME    — PRIME/DMA-BUF
#   0x010 = DRM_UT_ATOMIC   — Atomic modesetting
#   0x020 = DRM_UT_VBL      — Vblank
#   0x040 = DRM_UT_STATE    — Atomic 상태
#   0x080 = DRM_UT_LEASE    — DRM 리스
#   0x100 = DRM_UT_DP       — DisplayPort
#   0x200 = DRM_UT_DRMRES   — 리소스 관리

# 부팅 시 커널 파라미터로 설정
# drm.debug=0x1e (DRIVER+KMS+PRIME+ATOMIC)

# DRM debugfs 정보 확인
ls /sys/kernel/debug/dri/0/
# clients  gem_names  name  state  ...

# GEM 객체 목록 (메모리 사용량 확인)
cat /sys/kernel/debug/dri/0/gem_names

# 현재 KMS 상태 (모든 CRTC/Plane/Connector)
cat /sys/kernel/debug/dri/0/state

# GPU 행(hang) 시 ftrace로 추적
echo 1 > /sys/kernel/debug/tracing/events/drm/enable
cat /sys/kernel/debug/tracing/trace_pipe

GPU 성능 프로파일링

# 프레임 드롭 디버깅 — vblank 이벤트 추적
echo 1 > /sys/kernel/debug/tracing/events/drm/drm_vblank_event/enable
echo 1 > /sys/kernel/debug/tracing/events/drm/drm_vblank_event_delivered/enable
cat /sys/kernel/debug/tracing/trace_pipe
# 출력: vblank 시퀀스 번호와 타임스탬프 → 누락된 프레임 확인

# Atomic commit 성능 추적
echo 1 > /sys/kernel/debug/tracing/events/drm/drm_atomic_state/enable
# commit 소요 시간, 포함된 CRTC/Plane 수 확인

# GPU 메모리 사용량 상세 (드라이버별 debugfs)
# amdgpu:
cat /sys/kernel/debug/dri/0/amdgpu_gem_info       # 프로세스별 GEM 사용량
cat /sys/kernel/debug/dri/0/amdgpu_vram_mm        # VRAM 할당 맵
cat /sys/kernel/debug/dri/0/amdgpu_vm_info        # GPU 가상 메모리 통계

# i915:
cat /sys/kernel/debug/dri/0/i915_gem_objects      # GEM 객체 통계
cat /sys/kernel/debug/dri/0/i915_frequency_info   # GPU 주파수 정보

DRM ioctl 요약

DRM Lease (디스플레이 리스)

DRM Lease는 DRM Master가 자신의 KMS 리소스(CRTC, Connector, Plane) 일부를 다른 프로세스에게 독점적으로 위임하는 메커니즘입니다. VR 헤드셋, 멀티시트 디스플레이, 게임 전용 출력 등에서 사용됩니다.

/* DRM Lease 생성 (유저 공간, libdrm) */
uint32_t objects[] = {
    crtc_id,       /* 리스할 CRTC */
    connector_id,  /* 리스할 Connector */
    plane_id,      /* 리스할 Plane */
};

struct drm_mode_create_lease lease = {
    .object_ids   = (uint64_t)(uintptr_t)objects,
    .object_count = 3,
    .flags        = O_CLOEXEC | O_NONBLOCK,
};

drmIoctl(master_fd, DRM_IOCTL_MODE_CREATE_LEASE, &lease);
/* lease.lessee_id = 리스 ID */
/* lease.fd = lessee가 사용할 새 DRM fd */

/* Lessee는 lease.fd를 통해 리스된 오브젝트만 접근 가능 */
/* → 독립적인 atomic commit, 페이지 플립 수행 */

Content Protection (HDCP)

HDCP(High-bandwidth Digital Content Protection)는 디스플레이 출력의 콘텐츠를 암호화하여 무단 복제를 방지하는 DRM(Digital Rights Management) 기술입니다. Linux DRM 서브시스템은 KMS property를 통해 HDCP를 지원합니다.

/* HDCP KMS Properties (Connector) */

/* "Content Protection" property */
/* 값: Undesired (0) → Desired (1) → Enabled (2) */
/*     유저가 Desired로 설정하면 드라이버가 HDCP 인증 시작 */
/*     성공 시 Enabled로 전환 (읽기 전용) */

/* "HDCP Content Type" property */
/* 값: Type 0 — HDCP 1.4+ 허용 */
/*     Type 1 — HDCP 2.2+ 필수 (4K UHD 콘텐츠) */

/* Connector에 HDCP property 등록 (드라이버) */
drm_connector_attach_content_protection_property(connector, true);
/* true = HDCP Content Type property도 함께 등록 */

/* HDCP 인증 흐름 (드라이버 내부) */
/* 1. 유저가 "Content Protection" = Desired 설정 */
/* 2. atomic_check에서 HDCP 요청 감지 */
/* 3. 드라이버가 수신기(모니터)와 HDCP 핸드셰이크 */
/*    - AKE_Init / AKE_Send_Cert / AKE_Send_Km ... */
/* 4. 인증 성공 → "Content Protection" = Enabled */
/* 5. 링크 실패/핫플러그 → "Content Protection" = Desired */
/*    → 유저 공간에 uevent 통보 → 재인증 시도 */

디스플레이 연결 프로토콜

DisplayPort MST (Multi-Stream Transport)

DP MST는 하나의 DP 출력에서 여러 디스플레이를 데이지 체인으로 연결하거나 MST 허브를 통해 분기하는 기술입니다. 커널은 drm_dp_mst_topology_mgr로 MST 토폴로지를 관리합니다.

/* MST 토폴로지 매니저 초기화 (드라이버) */
drm_dp_mst_topology_mgr_init(&my->mst_mgr, &my->aux,
    16,   /* 최대 페이로드 수 */
    4,    /* 최대 레인 수 */
    conn_base_id);

/* MST 허브 감지 시 (HPD IRQ에서) */
drm_dp_mst_hpd_irq_handle_event(&my->mst_mgr, esi);
/* → 토폴로지 변경 시 connector 추가/제거 이벤트 */

/* MST 대역폭 계산 예시 */
/* DP 1.4 HBR3: 8.1 Gbps/레인 × 4 레인 = 32.4 Gbps 총 대역폭 */
/* 4K@60Hz (XRGB8888): ~12.5 Gbps 필요 → 63개 시간 슬롯 중 ~25개 사용 */
/* 나머지 슬롯으로 2번째 모니터 (FHD@60Hz) 연결 가능 */

DSC (Display Stream Compression)

DSC는 VESA가 표준화한 시각적 무손실(visually lossless) 디스플레이 압축입니다. 고해상도(8K, 4K@120Hz+)에서 DP/HDMI 대역폭 한계를 극복하기 위해 사용됩니다. Linux DRM에서는 drm_dsc_config로 DSC 파라미터를 관리합니다.

/* DSC 설정 예시 (드라이버 atomic_check에서) */
struct drm_dsc_config dsc_cfg = {};

dsc_cfg.line_buf_depth     = 13;
dsc_cfg.bits_per_component = 8;
dsc_cfg.bits_per_pixel     = 128;    /* 8 bpp × 16 (4비트 소수부) */
dsc_cfg.slice_width        = 1920;  /* 슬라이스 너비 */
dsc_cfg.slice_height       = 108;   /* 슬라이스 높이 */
dsc_cfg.pic_width          = 3840;
dsc_cfg.pic_height         = 2160;

/* RC (Rate Control) 파라미터 자동 계산 */
drm_dsc_compute_rc_parameters(&dsc_cfg);

/* PPS (Picture Parameter Set) 생성 → DP SDP 또는 HDMI infoframe으로 전송 */
drm_dsc_pps_payload_pack(pps_payload, &dsc_cfg);

fbdev 에뮬레이션

DRM/KMS는 /dev/fb0 (fbdev) 인터페이스를 에뮬레이션하여 레거시 애플리케이션과 콘솔 출력을 지원합니다. 커널 콘솔(fbcon), Plymouth(부팅 스플래시), 일부 임베디드 UI 프레임워크가 fbdev에 의존합니다.

GPU 컴퓨트 (GPGPU)

GPU는 렌더링 외에도 대규모 병렬 연산(GPGPU)에 사용됩니다. Linux 커널은 render node(/dev/dri/renderD128)를 통해 비특권 GPU 컴퓨트 접근을 제공하며, AMD는 KFD(Kernel Fusion Driver)로 HSA 컴퓨트를 추가 지원합니다.

KFD (Kernel Fusion Driver) — AMD HSA

SVM (Shared Virtual Memory)

SVM은 CPU와 GPU가 동일한 가상 주소 공간을 공유하는 기술입니다. malloc()으로 할당한 메모리를 GPU가 동일한 포인터로 접근할 수 있습니다.

/* 커널 SVM 지원 (amdgpu/KFD) */
#include <linux/hmm.h>

/* GPU 페이지 폴트 핸들러 */
/* 1. GPU가 매핑되지 않은 주소 접근 → 인터럽트 */
/* 2. 커널이 CPU 페이지 테이블에서 물리 주소 조회 */
/*    (페이지 없으면 CPU 페이지 폴트도 처리) */
/* 3. GPU 페이지 테이블에 매핑 추가 */
/* 4. GPU 작업 재개 */

# GPU 컴퓨트 관련 확인 명령

# Render node 확인
ls -la /dev/dri/renderD*

# KFD 디바이스 확인 (AMD)
ls -la /dev/kfd

# GPU 토폴로지 (AMD ROCm)
cat /sys/class/kfd/kfd/topology/nodes/0/properties

# GPU 메모리 사용량 (amdgpu)
cat /sys/class/drm/card0/device/mem_info_vram_used
cat /sys/class/drm/card0/device/mem_info_gtt_used

# clinfo (OpenCL 디바이스 정보)
clinfo

# vulkaninfo (Vulkan 컴퓨트 능력)
vulkaninfo --summary

DRM Accel 서브시스템 (AI/NPU 가속기)

최근 커널은 AI/ML/NPU 같은 비그래픽 가속기를 위해 DRM Accel 서브시스템을 제공합니다. 핵심 아이디어는 GPU DRM이 이미 갖고 있는 파일 디스크립터별 세션 상태, 버퍼 객체, 동기화, ioctl 디스패치를 재사용하되, 디스플레이와 렌더 노드 개념은 제거하고 /dev/accel/accel0 같은 전용 노드에 compute UAPI만 싣는 것입니다.

/* DRM Accel 드라이버 등록 (최소 골격) */
static const struct drm_driver my_accel_driver = {
    .driver_features = DRIVER_GEM | DRIVER_COMPUTE_ACCEL,
    .fops            = &my_accel_fops,
    .ioctls          = my_accel_ioctls,
    .num_ioctls      = ARRAY_SIZE(my_accel_ioctls),
    .name  = "my-npu",
    .desc  = "My NPU Accelerator",
    .date  = "20260101",
    .major = 1, .minor = 0,
};
/* drm_dev_register() 시 /dev/accel/accel0 자동 생성 */
/* (DRIVER_COMPUTE_ACCEL 플래그에 의해 accel 네임스페이스 사용) */

CUDA / NVIDIA — GPU 컴퓨트

CUDA(Compute Unified Device Architecture)는 NVIDIA가 2006년에 도입한 GPU 범용 컴퓨팅 플랫폼으로, GPU의 수천 개 코어를 C/C++ 확장 문법으로 프로그래밍할 수 있게 합니다. Linux에서 CUDA는 독점 커널 모듈(nvidia.ko)과 사용자 공간 런타임(libcuda.so, libcudart.so)으로 구성되며, 딥러닝(cuDNN, TensorRT), 과학 계산(cuBLAS, cuFFT), 고성능 컴퓨팅(NCCL, GPUDirect RDMA) 생태계가 핵심 경쟁력입니다.

CUDA 프로그래밍 모델의 핵심은 이종 컴퓨팅(Heterogeneous Computing)입니다. CPU(호스트)가 프로그램 흐름을 제어하고 GPU(디바이스)에 병렬 작업을 위임하는 구조로, 호스트 코드는 표준 C/C++ 컴파일러(gcc/clang)가, 디바이스 코드는 NVIDIA의 nvcc가 처리합니다. CUDA Runtime API(libcudart.so)는 고수준 추상화를, Driver API(libcuda.so)는 컨텍스트·모듈·함수 수준의 세밀한 제어를 제공합니다. 대부분의 애플리케이션은 Runtime API를 사용하며, 멀티 GPU·JIT 컴파일 등 고급 시나리오에서 Driver API를 활용합니다.

NVIDIA GPU 아키텍처 진화

NVIDIA GPU 아키텍처는 2006년 Tesla(CUDA 최초 지원)부터 시작하여 세대마다 SM(Streaming Multiprocessor) 구조, 메모리 계층, 인터커넥트, 전용 하드웨어 유닛을 혁신해 왔습니다. 각 세대의 Compute Capability는 지원하는 CUDA 기능 집합을 결정하며, nvcc의 -arch=sm_XX 옵션으로 대상 아키텍처를 지정합니다.

NVIDIA Linux 드라이버 스택

NVIDIA GPU를 CUDA로 활용하려면 커널 모듈 + 사용자 라이브러리가 함께 설치되어야 합니다. 아래 다이어그램은 CUDA 애플리케이션에서 GPU 하드웨어까지의 전체 소프트웨어 스택을 보여줍니다.

# NVIDIA 커널 모듈 확인
lsmod | grep nvidia
# nvidia               61440000  5 nvidia_uvm,nvidia_modeset
# nvidia_uvm            3280896  0
# nvidia_modeset        1282048  1 nvidia_drm
# nvidia_drm              94208  3
# nvidia_peermem         16384  0

# 디바이스 노드 확인
ls -la /dev/nvidia*
# crw-rw-rw- 1 root root 195,   0 ... /dev/nvidia0
# crw-rw-rw- 1 root root 195, 255 ... /dev/nvidiactl
# crw-rw-rw- 1 root root 511,   0 ... /dev/nvidia-uvm

CUDA는 두 가지 API 레벨을 제공합니다. Runtime API(libcudart.so)는 cudaMalloc(), cudaMemcpy() 등 간결한 함수로 대부분의 사용 사례를 커버합니다. Driver API(libcuda.so)는 cuCtxCreate(), cuModuleLoad() 등 더 세밀한 제어를 제공하며, PTX JIT 컴파일이나 멀티 컨텍스트 관리에 필수적입니다.

Linux CUDA 설치 및 환경 설정

Linux에서 CUDA 환경을 구성하는 방법은 크게 세 가지입니다: 배포판 패키지 매니저(apt/dnf), NVIDIA CUDA 저장소(cuda-keyring), runfile 직접 설치. 프로덕션 환경에서는 NVIDIA 공식 저장소를 통한 설치가 버전 관리와 업데이트 측면에서 권장됩니다.

# === 방법 1: NVIDIA 공식 저장소 (Ubuntu/Debian) ===
# 저장소 키링 패키지 설치
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt-get update

# CUDA Toolkit + 드라이버 설치 (메타 패키지)
sudo apt-get install cuda-toolkit-12-4
sudo apt-get install nvidia-open  # 오픈 커널 모듈 (Turing+)
# 또는: sudo apt-get install cuda-drivers  (독점 모듈)

# === 방법 2: RHEL/Rocky/AlmaLinux ===
sudo dnf config-manager --add-repo \
  https://developer.download.nvidia.com/compute/cuda/repos/rhel9/x86_64/cuda-rhel9.repo
sudo dnf install cuda-toolkit-12-4 nvidia-open

# === 환경 변수 설정 (~/.bashrc) ===
export PATH=/usr/local/cuda-12.4/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-12.4/lib64:$LD_LIBRARY_PATH

# === 설치 검증 ===
nvcc --version        # CUDA 컴파일러 버전
nvidia-smi            # 드라이버 및 GPU 상태
cuda-install-samples-12.4.sh ~/cuda-samples  # 샘플 코드 설치
cd ~/cuda-samples/Samples/1_Utilities/deviceQuery
make && ./deviceQuery  # GPU 정보 상세 출력

CUDA 프로그래밍 모델

CUDA는 SIMT(Single Instruction Multiple Thread) 실행 모델을 사용합니다. 프로그래머는 __global__ 함수(커널)를 정의하고, 호스트에서 <<<gridDim, blockDim>>> 구문으로 수천~수백만 스레드를 동시에 실행합니다.

/* 벡터 덧셈 — CUDA 커널 기본 예제 */
__global__ void vecAdd(const float *A, const float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N)
        C[i] = A[i] + B[i];
}

int main(void) {
    int N = 1 << 20;  /* 1M 원소 */
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, N * sizeof(float));
    cudaMalloc(&d_B, N * sizeof(float));
    cudaMalloc(&d_C, N * sizeof(float));

    /* 호스트→디바이스 전송 (생략) */

    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    cudaDeviceSynchronize();
    /* 디바이스→호스트 전송 (생략) */
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    return 0;
}

CUDA 커널 실행의 핵심 개념을 정리하면 다음과 같습니다:

/* Cooperative Groups — 워프 수준 리덕션 (CC 7.0+) */
#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

__global__ void warpReduce(const float *input, float *output, int N) {
    cg::thread_block block = cg::this_thread_block();
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = (idx < N) ? input[idx] : 0.0f;

    /* 워프 내 셔플 리덕션 — 레지스터 수준, 공유메모리 불필요 */
    float sum = cg::reduce(warp, val, cg::plus<float>());

    if (warp.thread_rank() == 0)
        atomicAdd(output, sum);
}

워프 실행과 분기 다이버전스

GPU의 SIMT(Single Instruction, Multiple Thread) 실행 모델에서 32개 스레드로 구성된 워프(Warp)는 동일한 프로그램 카운터(PC)를 공유합니다. 워프 내 모든 스레드가 같은 분기 경로를 따르면 최대 효율이지만, 서로 다른 경로를 택하면 분기 다이버전스(Branch Divergence)가 발생하여 각 경로를 순차적으로 실행해야 합니다.

분기 다이버전스를 최소화하는 것은 CUDA 최적화의 기본입니다. 워프 내 스레드들이 서로 다른 경로를 따를 때 SM의 실행 유닛 활용률이 떨어지며, 최악의 경우(32개 스레드 모두 다른 경로) 성능이 1/32로 저하될 수 있습니다.

/* 워프 셔플 리덕션 — 분기 없이 워프 합 계산 */
__device__ float warpReduceSum(float val) {
    for (int offset = warpSize / 2; offset > 0; offset /= 2)
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    return val;  /* lane 0에 합 결과 */
}

/* 워프 투표 함수 활용 */
__global__ void earlyExit(const int *data, int *result) {
    int val = data[blockIdx.x * blockDim.x + threadIdx.x];
    /* 워프 전체가 0이면 조기 종료 — 분기 다이버전스 없음 */
    if (__all_sync(0xFFFFFFFF, val == 0))
        return;
    /* ... 실제 연산 ... */
}

CUDA 스트림과 비동기 실행

CUDA 스트림(Stream)은 순서가 보장되는 GPU 명령 큐입니다. 서로 다른 스트림의 명령은 하드웨어가 허용하는 한 동시에 실행될 수 있으며, 이를 통해 커널 실행, 메모리 전송, 호스트 연산을 오버랩하여 GPU 파이프라인 활용률을 극대화합니다.

/* 3-스트림 파이프라인 — 전송과 커널 오버랩 */
const int nStreams = 3;
cudaStream_t streams[nStreams];
for (int i = 0; i < nStreams; i++)
    cudaStreamCreate(&streams[i]);

int chunkSize = N / nStreams;
for (int i = 0; i < nStreams; i++) {
    int offset = i * chunkSize;
    /* 비동기 H→D 전송 (핀드 메모리 필수) */
    cudaMemcpyAsync(d_in + offset, h_in + offset,
        chunkSize * sizeof(float), cudaMemcpyHostToDevice, streams[i]);

    /* 커널 실행 — 같은 스트림이므로 전송 완료 후 자동 실행 */
    myKernel<<<chunkSize/256, 256, 0, streams[i]>>>(d_in + offset, d_out + offset);

    /* 비동기 D→H 전송 */
    cudaMemcpyAsync(h_out + offset, d_out + offset,
        chunkSize * sizeof(float), cudaMemcpyDeviceToHost, streams[i]);
}

/* 모든 스트림 완료 대기 */
cudaDeviceSynchronize();

/* 이벤트로 경과 시간 측정 */
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start, streams[0]);
myKernel<<<grid, block, 0, streams[0]>>>(d_in, d_out);
cudaEventRecord(stop, streams[0]);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
/* ms = 커널 실행 시간 (밀리초) */

/* CUDA 그래프 — 반복 실행 워크플로 최적화 (CC 7.0+) */
cudaGraph_t graph;
cudaGraphExec_t graphExec;

/* 1. 스트림 캡처로 그래프 기록 */
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
cudaMemcpyAsync(d_in, h_in, size, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(d_in, d_out);
cudaMemcpyAsync(h_out, d_out, size, cudaMemcpyDeviceToHost, stream);
cudaStreamEndCapture(stream, &graph);

/* 2. 그래프 인스턴스화 (1회) */
cudaGraphInstantiate(&graphExec, graph, 0);

/* 3. 반복 실행 — 실행 오버헤드 대폭 감소 */
for (int iter = 0; iter < 1000; iter++)
    cudaGraphLaunch(graphExec, stream);

cudaGraphExecDestroy(graphExec);
cudaGraphDestroy(graph);

CUDA 메모리 계층

GPU 성능 최적화의 핵심은 메모리 계층을 이해하고 활용하는 것입니다. CUDA 메모리는 크게 레지스터 → 공유 메모리(Shared) → L1/L2 캐시 → 글로벌 메모리(VRAM) 순으로 용량이 커지고 지연 시간이 증가합니다.

/* 공유 메모리 타일링 — 행렬 곱셈 최적화 */
#define TILE 16

__global__ void matMul(const float *A, const float *B, float *C, int N) {
    __shared__ float sA[TILE][TILE], sB[TILE][TILE];
    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < N / TILE; t++) {
        sA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col];
        __syncthreads();  /* 블록 내 동기화 */

        for (int k = 0; k < TILE; k++)
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

/* 공유 메모리 뱅크 충돌 회피 — 패딩 기법 */
/* 뱅크 충돌: 같은 뱅크를 동시 접근하면 순차 처리 */
/* 32개 뱅크, 4B 인터리빙: bank = (addr / 4) % 32 */

/* ❌ 뱅크 충돌 발생 — 열 접근 */
__shared__ float tile[32][32];  /* tile[0][0], tile[1][0]은 같은 뱅크 */
float val = tile[threadIdx.x][0];  /* 32-way 뱅크 충돌! */

/* ✅ 패딩으로 해결 */
__shared__ float tile[32][33];  /* +1 패딩 → 뱅크 오프셋 이동 */
float val = tile[threadIdx.x][0];  /* 충돌 없음 */

/* ✅ swizzle 기법 (CUTLASS 스타일) */
int swizzled_col = col ^ (row & 0x1F);
float val = tile[row][swizzled_col];

UVM (Unified Virtual Memory) 심층 분석

Unified Virtual Memory(UVM)는 CPU와 GPU가 동일한 가상 주소 공간을 공유하는 메커니즘입니다. nvidia-uvm.ko 커널 모듈이 페이지 폴트 기반 마이그레이션을 처리하여, 프로그래머가 명시적 cudaMemcpy() 없이도 양쪽에서 데이터에 접근할 수 있습니다.

/* UVM 최적화 패턴 — 프리페치 + 힌트 */
float *data;
cudaMallocManaged(&data, size);

/* CPU에서 데이터 초기화 */
initDataOnCPU(data, N);

/* GPU로 선제적 마이그레이션 (폴트 없이 전송) */
cudaMemPrefetchAsync(data, size, deviceId, stream);

/* 읽기 전용 데이터: 양쪽에 복제본 유지 */
cudaMemAdvise(readOnlyData, size, cudaMemAdviseSetReadMostly, deviceId);

/* GPU 커널 실행 */
myKernel<<<grid, block, 0, stream>>>(data);

/* CPU로 다시 마이그레이션 */
cudaMemPrefetchAsync(data, size, cudaCpuDeviceId, stream);
cudaStreamSynchronize(stream);
processOnCPU(data, N);

Tensor Core 연산

Tensor Core는 NVIDIA GPU에 내장된 행렬 연산 전용 하드웨어 유닛으로, Volta(V100) 세대에서 처음 도입되었습니다. 단일 사이클에 작은 행렬의 FMA(Fused Multiply-Add) 연산 D = A × B + C를 수행하며, 일반 CUDA 코어 대비 4~16× 높은 연산 처리량을 달성합니다.

/* WMMA API — Tensor Core 프로그래밍 (CC 7.0+) */
#include <mma.h>
using namespace nvcuda::wmma;

__global__ void tensorGemm(const half *A, const half *B, float *C) {
    /* 16×16×16 타일 단위 MMA */
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> c_frag;

    fill_fragment(c_frag, 0.0f);

    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / warpSize;
    int warpN = blockIdx.y;

    /* A와 B 타일 로드 */
    load_matrix_sync(a_frag, A + warpM * 16 * K, K);
    load_matrix_sync(b_frag, B + warpN * 16, N);

    /* Tensor Core MMA: D = A × B + C */
    mma_sync(c_frag, a_frag, b_frag, c_frag);

    /* 결과 저장 */
    store_matrix_sync(C + warpM * 16 * N + warpN * 16, c_frag, N, mem_row_major);
}

오큐펀시 최적화

오큐펀시(Occupancy)는 SM이 동시에 유지할 수 있는 활성 워프 수 대비 실제 실행 중인 워프의 비율입니다. 높은 오큐펀시는 메모리 지연 시간을 워프 스위칭으로 효과적으로 숨길 수 있게 합니다. 오큐펀시를 결정하는 3대 제약 요소는 레지스터 사용량, 공유 메모리 사용량, 블록당 스레드 수입니다.

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

cudnnTensorDescriptor_t inputDesc, outputDesc;
cudnnFilterDescriptor_t filterDesc;
cudnnConvolutionDescriptor_t convDesc;

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

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

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

cudnnDestroy(cudnn);

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

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

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

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

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

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

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

    float sum = 0.0f;

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

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

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

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

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

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

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

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

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

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

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

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

#version 450

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

layout(local_size_x = 256) in;

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

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

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

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

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

    barrier();

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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