CUDA 동적 병렬성

21.1 일이 어디서 터질지 모를 때

지금까지 우리가 다룬 모든 커널은 사실상 같은 가정을 깔고 있었다. "호스트가 미리 격자(grid)를 결정한다." 행렬이 1024x1024라면 블록 16x16짜리를 64x64로 깐다. 입력 크기를 보고 호스트가 한 번에 격자를 정하고, GPU는 그 격자에 맞춰 일을 해치운다. 데이터가 균질하다면 이 모델은 훌륭하다. 모든 픽셀에 같은 필터를 적용하거나, 모든 원소에 같은 활성 함수를 걸 때, 한 번의 launch가 깔끔하게 끝난다.

그런데 세상에는 일이 어디서 터질지 모르는 문제가 있다. 난류 시뮬레이션을 떠올려 보자. 유체가 잔잔한 영역에서는 거친 격자로도 충분하지만, 충격파가 지나가는 자리, 와류가 도는 자리에서는 격자를 100배 촘촘하게 잡아도 모자라다. 이걸 한 번에 가장 촘촘한 격자로 깔면 메모리도 시간도 폭발한다. "거친 데서는 거칠게, 미세한 데서는 미세하게" — 이 적응 격자(adaptive mesh refinement, AMR) 아이디어가 진작부터 있었던 이유다.

비슷한 이야기는 여기저기에 있다. 광선 추적(ray tracing)에서 한 픽셀의 광선은 한두 번 반사하고 끝나지만, 옆 픽셀의 광선은 거울방 안에 들어가서 30번 튀고 나올 수 있다. 그래프 탐색에서 한 정점은 차수가 3인데 옆 정점은 100만 짜리 허브일 수 있다. 트리 탐색이나 분할 정복 알고리즘은 한 호출이 자식 호출을 또 만들고, 그 자식이 또 자식을 만든다. 이런 작업량은 입력을 보지 않고는 알 수 없다.

21.2 동적 병렬성: 디바이스가 직접 부른다

CUDA 5.0(컴퓨트 능력 3.5, Kepler GK110)부터 도입된 동적 병렬성(dynamic parallelism)은 이 호스트 왕복을 끊어 버리는 기능이다. 한 줄로 요약하면 "디바이스 코드 안에서 또 다른 __global__ 커널을 launch할 수 있다." 부모(parent) 커널이 돌다가 어떤 영역에 일이 더 필요하다고 판단하면, 그 자리에서 자식(child) 커널을 호출한다. 호스트는 그 사이에 끼지 않는다.

문법 자체는 우리가 이미 아는 것과 똑같다. kernel<<<grid, block>>>(args)를 디바이스 함수 안에서 그냥 쓰면 된다. 다만 몇 가지 새로운 개념이 따라온다. 첫째, 부모 스레드와 자식 grid는 비동기다. launch한 시점에 부모는 다음 줄로 진행하고, 자식이 다 끝났는지 확인하려면 cudaDeviceSynchronize()를 명시적으로 불러야 한다. 둘째, 자식 grid는 부모와 같은 스트림을 공유할 수도 있고 새 스트림으로 분리할 수도 있다. 셋째, "부모의 부모의 부모…"가 얼마나 깊어질 수 있는지 — nesting depth — 가 하드웨어/런타임 제한에 걸린다. 보통 24단계까지 허용되지만, 동기화를 거치는 깊이는 그보다 훨씬 얕게(보통 2~3단계) 잡는 것이 안전하다.

왜 동적 병렬성이 매력적인가? 호스트 왕복 제거 외에도, 코드의 표현력이 자연스러워진다. 재귀 알고리즘을 재귀처럼 쓸 수 있다. 트리 탐색을 트리답게 쓸 수 있다. 분할 정복을 분할 정복답게 쓸 수 있다. 적응 격자는 격자가 적응하는 자리에 직접 launch를 박을 수 있다. "GPU에서 도는 코드가 GPU 일거리를 만든다"는 한 단계 더 자율적인 모델이다.

21.3 예제 1 — 곡률 적응 베지에 곡선

가장 깔끔한 예제 하나를 먼저 보자. 베지에(Bezier) 곡선들을 화면에 그리는데, 곡선마다 곡률(curvature)이 다르다. 거의 직선에 가까운 베지에는 4~8개의 점만으로도 매끈해 보이지만, 급격히 휘는 곡선은 256개 점으로 분할해도 각져 보일 수 있다. 모든 곡선을 일률적으로 가장 미세하게 분할하면 직선 같은 곡선에서 메모리와 시간이 낭비된다.

해결책은 두 단계 launch다. 부모 커널은 곡선당 한 스레드씩 받아서 곡률을 계산하고, 곡률에 비례한 분할 수 nTess를 정한다. 그러고 나서 각 스레드가 자식 커널을 launch해 그 곡선 하나를 nTess개의 점으로 평가한다. 호스트는 한 번 부모만 부르면 된다.

// 베지에 곡선 한 개를 표현하는 구조체
struct Bezier {
    float2 p0, p1, p2;     // 제어점 3개 (2차 곡선)
    float2 *vertices;      // 평가된 점들이 들어갈 버퍼
    int nTess;             // 자식 launch 시 결정
};

__device__ float compute_curvature(const Bezier& b) {
    // 제어점 사이 거리 비율로 거친 곡률 추정
    float2 v0 = make_float2(b.p1.x - b.p0.x, b.p1.y - b.p0.y);
    float2 v1 = make_float2(b.p2.x - b.p1.x, b.p2.y - b.p1.y);
    float cross = v0.x * v1.y - v0.y * v1.x;
    float len  = sqrtf(v0.x*v0.x + v0.y*v0.y) + sqrtf(v1.x*v1.x + v1.y*v1.y) + 1e-6f;
    return fabsf(cross) / (len * len);
}

__global__ void tessellate_bezier(Bezier* curves, int nCurves);
__global__ void evaluate_bezier(Bezier* curve);

부모 커널은 곡선당 한 스레드를 받아 분할 수를 결정하고, 결정된 격자 크기로 자식을 launch한다.

__global__ void tessellate_bezier(Bezier* curves, int nCurves) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= nCurves) return;

    float kappa = compute_curvature(curves[i]);
    int nTess = max(4, min(256, (int)(8.0f + 256.0f * kappa)));
    curves[i].nTess = nTess;

    // 메모리는 디바이스 malloc으로 할당 가능
    curves[i].vertices = (float2*)malloc(sizeof(float2) * nTess);

    // 자식 grid: 한 곡선당 nTess개의 평가점
    int threadsPerBlock = 32;
    int blocks = (nTess + threadsPerBlock - 1) / threadsPerBlock;
    evaluate_bezier<<<blocks, threadsPerBlock>>>(&curves[i]);
    // 부모는 자식 종료를 굳이 기다리지 않아도 launch 자체로 OK.
    // 단, 부모 커널이 끝나기 전에 자식이 다 안 끝나면 호스트 시점에서는 자식도 진행 중.
}

__global__ void evaluate_bezier(Bezier* curve) {
    int t = blockIdx.x * blockDim.x + threadIdx.x;
    if (t >= curve->nTess) return;

    float u = (float)t / (float)(curve->nTess - 1);
    float w = 1.0f - u;
    // 2차 베지에: B(u) = w^2 P0 + 2wu P1 + u^2 P2
    curve->vertices[t].x = w*w*curve->p0.x + 2*w*u*curve->p1.x + u*u*curve->p2.x;
    curve->vertices[t].y = w*w*curve->p0.y + 2*w*u*curve->p1.y + u*u*curve->p2.y;
}

호스트는 한 번 tessellate_bezier만 부르면 된다. 곡선 1만 개가 있고, 그중 7천 개가 거의 직선, 3천 개가 격렬한 곡선이라고 해 보자. 일률적으로 256분할하면 256만 점을 만든다. 곡률 적응이면 7천 곡선은 평균 8점, 3천 곡선은 평균 128점 — 합쳐도 50만 점 안쪽이다. 메모리/시간을 5배 가까이 절약한다.

21.4 재귀 예제 — Quadtree 구축

동적 병렬성의 진짜 힘이 드러나는 곳은 재귀다. 2차원 점 분포를 quadtree로 분할하는 문제를 보자. 영역(bounding box) 안에 점이 임계값(가령 32개) 이하로 있으면 더 이상 나누지 않는다. 그보다 많으면 영역을 4등분하고, 각 사분면에 들어가는 점들을 골라내 재귀적으로 같은 일을 한다. 점 분포가 한쪽에 몰려 있으면 그쪽은 깊게, 다른 쪽은 얕게 들어간다 — 작업량이 데이터에 따라 다르다.

호스트 주도로 풀면 BFS로 한 단계씩 까며 단계마다 launch + 호스트 복귀를 반복해야 한다. 동적 병렬성에서는 한 노드 처리 커널이 자식 4개를 자기가 launch한다. 코드 골격은 이렇다.

struct QNode {
    float2 bbox_min, bbox_max;
    int* point_indices;     // 이 노드에 속한 점 인덱스 배열
    int n_points;
    QNode* children[4];      // 자식 4개 (NULL이면 leaf)
    int depth;
};

__global__ void build_quadtree(QNode* node, float2* all_points,
                              int threshold, int max_depth) {
    // 한 노드를 한 블록이 처리한다고 하자.
    if (threadIdx.x == 0) {
        // leaf 조건 검사
        if (node->n_points <= threshold || node->depth >= max_depth) {
            for (int k = 0; k < 4; ++k) node->children[k] = nullptr;
            return;
        }
    }
    __syncthreads();

    // 1) 사분면별 점 개수 카운트 (블록 내 reduce/scan)
    // 2) 사분면별 점 인덱스 분배 (compaction)
    // 3) 자식 4개 메모리 할당
    // (생략 — 블록 내 협력으로 분할)

    if (threadIdx.x == 0) {
        for (int k = 0; k < 4; ++k) {
            QNode* child = (QNode*)malloc(sizeof(QNode));
            child->depth = node->depth + 1;
            // child->bbox, point_indices, n_points 채우기
            node->children[k] = child;

            // 자식 노드 처리도 GPU에서 — 동적 병렬성!
            cudaStream_t s;
            cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
            build_quadtree<<<1, 128, 0, s>>>(child, all_points, threshold, max_depth);
            cudaStreamDestroy(s);
        }
    }
}

네 자식을 각각 별도 스트림에 띄운 점에 주목하자. 같은 스트림에 모두 보내면 사분면들이 직렬로 처리된다. 별도 스트림이면 GPU가 자원만 있으면 동시에 진행할 수 있다. 단, 스트림 핸들 자체를 디바이스에서 만들어야 하므로 cudaStreamCreateWithFlags를 디바이스 런타임이 지원해야 한다(Compute Capability 3.5+).

깊이 관리도 까다롭다. 1억 개의 점을 잘게 나누면 quadtree 깊이가 20을 넘을 수 있다. nesting depth 제한을 넘으면 launch가 실패한다. 대안은 깊이가 어느 수준 이상이면 동적 병렬성을 멈추고 그 노드 안의 점들을 한 블록 안에서 직렬/병렬로 처리해 끝내는 것이다 — "재귀의 깊이 제한 + 잎(leaf) 가까이서 평탄화"는 GPU 재귀의 일반적 패턴이다.

   호스트 주도(BFS)              동적 병렬성(GPU 재귀)
   ─────────────                  ──────────────────
   level 0:  [root] ─launch─▶ K   level 0: parent K──┐
              ↓ copy back to host             │ self-launches
   level 1:  [c0,c1,c2,c3] ──▶ K              ▼
              ↓ copy back                     [c0]K  [c1]K  [c2]K  [c3]K
   level 2:  ... ──▶ K                          │      │      │      │
              ↓ copy back                       ▼      ▼      ▼      ▼
   level d:  ... ──▶ K                        ...    ...    ...    ...

   GPU↔CPU 왕복 d번                          호스트는 부모 1번만 launch

그림 21.1 — 트리 처리에서 호스트 주도 BFS와 디바이스 재귀의 작동 차이.

21.5 실전 고려사항

동적 병렬성은 우아하지만 공짜는 아니다. 실전에서 부딪히는 함정 몇 가지를 정리한다.

(1) Pending launch pool. 부모가 launch한 자식 grid는 하드웨어가 즉시 실행하지 않을 수 있다. 큐에 들어갔다가 자원이 비면 실행된다. 이 큐의 크기가 기본값에서는 작아서(보통 2048개), 자식을 우다다 launch하면 큐가 차서 launch가 실패하기 시작한다. cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, N)으로 늘려야 한다.

(2) 메모리 가시성 규칙. 자식 grid는 부모의 로컬 변수(레지스터, 스택)를 볼 수 없다. 부모의 글로벌 메모리, 그리고 동적 할당한 힙은 자식이 보지만, 부모 함수의 지역 변수 주소를 자식에게 넘기는 것은 정의되지 않은 동작이다. 이 때문에 자식에게 데이터를 넘길 때는 글로벌 메모리에 한 번 써 두거나, malloc으로 힙에 잡은 메모리를 통해 넘긴다.

(3) 부모/자식 동기화. 한 부모가 여러 자식을 띄우고, 자식들이 끝나기를 기다리려면 cudaDeviceSynchronize()를 부른다. 이때 부모 스레드가 그 호출에서 정말로 막히기 때문에, 같은 워프의 다른 스레드들이 불필요하게 대기할 수 있다. 가능한 한 동기화는 한 워프의 한 스레드(보통 lane 0)에서만 부르고, 결과를 워프에 브로드캐스트하는 패턴이 일반적이다.

(4) Launch 오버헤드. 호스트 launch는 보통 5~10us 정도지만, 디바이스 launch도 1~3us는 든다. 자식이 하는 일이 워프 한두 개의 1us짜리 일이면 launch 비용이 일 자체보다 크다. 동적 병렬성의 손익분기는 "자식의 일이 적어도 수십 us는 되어야 한다"가 경험칙이다. 그렇지 않다면 자식 launch 대신 부모 안에서 직접 처리하는 평탄화가 낫다.

// 동적 병렬성 사용 시 자주 같이 부르는 한도 설정
cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 3);              // 동기화 깊이
cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 1<<14); // 큐 크기 16K
cudaDeviceSetLimit(cudaLimitMallocHeapSize, 256*1024*1024);    // 디바이스 malloc 힙 256MB

// 컴파일 시 -rdc=true (relocatable device code) + cudalibrary 링크 필요
// nvcc -rdc=true -arch=sm_70 -lcudadevrt your_kernels.cu

(5) 디버깅과 프로파일링. 자식 grid는 부모와 다른 launch 이벤트로 기록된다. Nsight Systems에서는 한 줄 한 줄 child kernel이 보인다. 너무 많은 작은 자식이 보이면 평탄화 후보다. 반대로 어떤 부모가 자식을 잘못된 메모리로 띄워 죽으면 에러는 자식 launch 시점이 아니라 그 다음 동기화 시점에 보고된다 — 디버그할 때 추적하기가 까다롭다.

21.6 정리

동적 병렬성은 GPU 프로그래밍의 표현력을 한 단계 끌어올리는 기능이다. 호스트가 모든 launch를 결정하던 시대에서, GPU 안의 코드가 자기 일거리를 만드는 시대로 옮겨 가는 다리다. 적응 격자, 광선 추적, 트리 알고리즘, 그래프 탐색에서 호스트 왕복을 제거하고 코드를 자연스럽게 만들 수 있다.

그러나 그것이 만능 망치는 아니다. Launch 오버헤드, 메모리 가시성 규칙, pending launch pool, nesting depth 제한, 디버깅의 어려움 — 모두 무게를 더하는 요소다. 가장 좋은 자세는 "고전 모델로 우아하게 풀리지 않을 때만 동적 병렬성을 꺼내는 것"이다. 기능을 알지만 함부로 쓰지 않는 것 — 이게 GPU 프로그래머의 미덕 중 하나다.

다음 장에서는 시야를 더 넓혀, CUDA 프로그래밍 환경 자체가 어떻게 진화해 왔고 앞으로 어디로 갈지를 살펴본다. 메모리 모델의 변천부터 컴파일러, 라이브러리, 인터커넥트, 그리고 미래의 지평까지.