CUDA로 구현하는 병렬 그래프 탐색과 BFS·DFS 설계

frontier를 동시에 확장하는 그래프 탐색

그래프 병렬 처리는 정점과 간선에 대한 연산을 여러 스레드, 코어, GPU 워프로 나눠 동시에 실행하는 방식이다. 데이터 병렬성과 작업 병렬성이 함께 얽히며, 그래프의 희소성·차수 분포·동적성에 따라 병렬화 난이도와 효율이 달라진다.

BFS는 동일 깊이의 정점 집합인 frontier를 병렬로 확장하기에 적합하다. 레벨 동기화 방식으로 실행하고, 활성 정점 중심으로 확장하는 push와 미방문 정점이 이웃의 방문 여부를 확인하는 pull을 전환해 간선 접근량을 줄일 수 있다. 방문 상태를 기록하고 다음 frontier를 만들 때는 중복 삽입을 막기 위한 원자 연산이 필요하다.

DFS는 깊이 우선이라는 성격 때문에 직렬성이 높다. 다만 다중 시작점, 서브트리 병렬 확장, 작업 훔치기(work stealing)를 활용하면 병렬 처리 대상으로 구성할 수 있다. 이때 경로 의존성, 스택 경쟁, 결과 비결정성을 제어할 스케줄링이 필요하다.

CUDA에서는 다수의 경량 스레드를 블록과 워프 단위로 스케줄하는 SIMT 모델을 사용한다. 메모리 결합(coalescing)과 공유 메모리 사용 방식이 성능에 큰 영향을 주며, 희소 그래프에는 CSR·COO 표현과 프리픽스-섬(prefix-sum) 기반 워크큐 구성이 널리 쓰인다.

그래프 표현과 작업 분배를 함께 설계하기

탐색용 그래프는 CSR(rowPtr, colInd) 형식으로 구성하면 간선을 순차 접근하기 쉽고 메모리 결합에도 유리하다. COO는 삽입과 정렬에는 유연하지만 탐색 과정에서는 캐시 효율이 떨어질 수 있다.

frontier는 큐나 벡터로 유지하면서 비트셋을 함께 둘 수 있다. 비트셋은 방문 여부 확인 비용을 줄이고, 큐는 다음 레벨의 정점을 구성하는 데 유리하다.

차수 편차를 다루려면 정점 병렬과 간선 병렬을 섞어야 한다. 고차수 허브에는 전용 커널을 두거나 warp-centric 확장을 적용해 스레드 다이버전스를 완화할 수 있다. 동적 워크큐와 prefix-sum 분배는 불균형한 차수에 대응하는 수단이며, 매우 큰 인접 리스트는 CTA(block) 협력 방식으로 나눠 처리한다.

멀티 GPU 환경에서는 파티션의 cut을 줄이고 고스트 노드 동기화 주기를 관리하는 일이 핵심이 된다. CPU 전처리, H2D 전송, 반복 커널 실행, D2H 결과 수집으로 이어지는 파이프라인을 만들고 CUDA 스트림으로 전송과 커널을 겹쳐 실행할 수 있다. OOM이나 실행 오류에 대비해 체크포인트와 부분 레벨 재시도 경로도 필요하다.

GPU BFS가 반복되는 흐름

⤢✕

입력은 CSR(rowPtr, colInd), 시작점, 선택적 가중치로 구성한다. 처리 단계에서는 레벨을 반복하며 push/pull 전환, 로드밸런싱, 원자 연산 최소화를 적용한다. 결과로 거리 배열, 부모 배열, 도달성 정보를 생성한다.

OOM이 발생하면 배치 크기를 줄이거나 frontier를 분할하고, 방향 전환을 재시도하는 절차를 둘 수 있다.

CUDA 커널로 보는 레벨 동기화 BFS

전제조건은 CUDA 12.x, GPU 아키텍처 sm_80 이상 권장, 입력 CSR 형식 가정이다. 소규모 그래프에서 검증한 뒤 대규모 데이터에 적용한다.

컴파일에는 nvcc -O3 -arch=sm_80 bfs.cu -o bfs를 사용한다. 실행 시 CUDA_LAUNCH_BLOCKING=1 환경 변수를 설정하면 디버그가 용이하다.

// bfs.cu
// 최소 실행 가능 예시, 단일 소스 BFS, 방향은 push 모드 중심
#include <cuda_runtime.h>
#include <vector>
#include <cstdio>

__global__ void bfs_expand_kernel(
    const int* __restrict__ rowPtr,
    const int* __restrict__ colInd,
    const int* __restrict__ frontier,
    int frontierSize,
    int level,
    int* __restrict__ visited,   // -1: unvisited, >=0: level
    int* __restrict__ parent,
    int* __restrict__ nextFrontier,
    int* __restrict__ nextCount)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= frontierSize) return;

    int u = frontier[tid];
    int start = rowPtr[u];
    int end   = rowPtr[u + 1];

    // 고차수 정점 처리: warp 협력 최적화는 단순화하여 생략
    for (int e = start; e < end; ++e) {
        int v = colInd[e];
        // 방문 체크: -1 에서 level 로 마킹
        if (atomicCAS(&visited[v], -1, level + 1) == -1) {
            parent[v] = u;
            int pos = atomicAdd(nextCount, 1);
            nextFrontier[pos] = v; // 중복 삽입 최소화
        }
    }
}

void check(cudaError_t err, const char* msg) {
    if (err != cudaSuccess) {
        fprintf(stderr, "CUDA Error at %s: %s\n", msg, cudaGetErrorString(err));
        exit(1);
    }
}

int main() {
    // 예시 그래프: 0-1-2-3 선형
    std::vector<int> h_rowPtr = {0, 1, 2, 3, 3};
    std::vector<int> h_colInd = {1, 2, 3};
    int n = 4, m = 3, src = 0;

    int *d_rowPtr, *d_colInd, *d_frontier, *d_nextFrontier, *d_visited, *d_parent, *d_nextCount;
    check(cudaMalloc(&d_rowPtr, (n+1)*sizeof(int)), "malloc rowPtr");
    check(cudaMalloc(&d_colInd, m*sizeof(int)), "malloc colInd");
    check(cudaMemcpy(d_rowPtr, h_rowPtr.data(), (n+1)*sizeof(int), cudaMemcpyHostToDevice), "copy rowPtr");
    check(cudaMemcpy(d_colInd, h_colInd.data(), m*sizeof(int), cudaMemcpyHostToDevice), "copy colInd");

    check(cudaMalloc(&d_frontier, n*sizeof(int)), "malloc frontier");
    check(cudaMalloc(&d_nextFrontier, n*sizeof(int)), "malloc nextFrontier");
    check(cudaMalloc(&d_visited, n*sizeof(int)), "malloc visited");
    check(cudaMalloc(&d_parent, n*sizeof(int)), "malloc parent");
    check(cudaMalloc(&d_nextCount, sizeof(int)), "malloc nextCount");

    std::vector<int> h_visited(n, -1), h_parent(n, -1), h_frontier(n, -1);
    h_frontier[0] = src; h_visited[src] = 0;
    check(cudaMemcpy(d_frontier, h_frontier.data(), n*sizeof(int), cudaMemcpyHostToDevice), "copy frontier");
    check(cudaMemcpy(d_visited, h_visited.data(), n*sizeof(int), cudaMemcpyHostToDevice), "copy visited");
    check(cudaMemcpy(d_parent, h_parent.data(), n*sizeof(int), cudaMemcpyHostToDevice), "copy parent");

    int frontierSize = 1, level = 0;
    dim3 blk(256);

    while (frontierSize > 0) {
        check(cudaMemset(d_nextCount, 0, sizeof(int)), "memset nextCount");

        dim3 grd((frontierSize + blk.x - 1) / blk.x);
        bfs_expand_kernel<<<grd, blk>>>(
            d_rowPtr, d_colInd, d_frontier, frontierSize, level,
            d_visited, d_parent, d_nextFrontier, d_nextCount);
        check(cudaGetLastError(), "kernel launch");
        check(cudaDeviceSynchronize(), "kernel sync");

        int h_nextCount = 0;
        check(cudaMemcpy(&h_nextCount, d_nextCount, sizeof(int), cudaMemcpyDeviceToHost), "copy nextCount");
        if (h_nextCount == 0) break;

        // frontier 스왑
        check(cudaMemcpy(d_frontier, d_nextFrontier, h_nextCount*sizeof(int), cudaMemcpyDeviceToDevice), "swap frontier");
        frontierSize = h_nextCount;
        level++;
    }

    // 결과 수집
    check(cudaMemcpy(h_visited.data(), d_visited, n*sizeof(int), cudaMemcpyDeviceToHost), "copy visited");
    check(cudaMemcpy(h_parent.data(), d_parent, n*sizeof(int), cudaMemcpyDeviceToHost), "copy parent");

    for (int i = 0; i < n; ++i) {
        printf("v=%d, dist=%d, parent=%d\n", i, h_visited[i], h_parent[i]);
    }

    cudaFree(d_rowPtr); cudaFree(d_colInd);
    cudaFree(d_frontier); cudaFree(d_nextFrontier);
    cudaFree(d_visited); cudaFree(d_parent); cudaFree(d_nextCount);
    return 0;
}

frontier 밀도를 기준으로 활성 정점 비율이 10%를 넘을 때 pull 모드로 전환하는 로직을 추가할 수 있다. 고차수 정점은 warp-cooperative strip-mining으로 인접 리스트를 32스레드가 함께 처리하도록 설계할 수 있다. 방문 배열 외에 비트셋이나 블룸필터를 보조로 사용하면 원자 연산 빈도를 더 낮출 수 있다.

병렬 탐색이 쓰이는 장면

대규모 소셜 그래프에서는 커뮤니티 탐지 이전의 거리 기반 시드 전파, 영향 관계 추정, 추천 그래프 확장에 BFS를 적용할 수 있다. 수십억 간선 규모에서도 실시간 또는 배치 준실시간 분석을 달성할 수 있다.

보안과 위협 헌팅에서는 공격 그래프의 다중 시작점 BFS로 도달성과 최단 침투 경로를 찾는다. 병렬화를 적용하면 윈도우 수초 내 네트워크 단위 스캔이 가능하다.

도로망과 항로망에서는 장애 전파를 시뮬레이션하거나 근사 최단 경로를 초기화하는 데 활용한다. GPU 가속은 이벤트 발생 후 재계산 지연을 줄이는 데 쓰인다. 그래프 데이터베이스의 경우 패턴 매칭 이전에 BFS와 방문 필터링으로 후보를 축소하고, OLAP 워크로드에서 TEPS(Traversed Edges Per Second)를 높여 집계 지연 감소를 기대할 수 있다.

처리량과 확장성의 조건

단일 GPU 기준 BFS 처리량은 1~10+ GTEPS 범위로 보고된 사례가 있다(최신 정보 확인 필요). 10억 간선 그래프를 5 GTEPS로 가정하면 탐색 시간은 ≈ 0.2초이며, 같은 그래프를 CPU 0.1 GTEPS로 가정하면 ≈ 10초가 소요되는 것으로 추정한다.

멀티 GPU로 확장하면 선형에 근접한 확장성을 얻을 수 있다. 다만 파티션 cut율과 통신 은닉 비율에 따라 0.7~0.9 배율 기대값을 산정할 수 있다.

대화형 분석 가능성을 확보하고 운영 의사결정 지연을 줄일 수 있으며, 단일 노드 자원 대비 전력당 성능 효율도 개선된다. 정렬과 중복 제거 같은 결정성·재현성 옵션은 실험 과학 및 금융 규제 환경에서 신뢰성을 높이는 데 사용된다.

접근	성능(TEPS)	확장성	일관성(결과 결정성)	안정성(메모리/에러)	운영 편의
CPU 멀티코어(OpenMP)	낮음~중간	소켓/NUMA 한계 내 선형 근사	높음(순차성 유지 용이)	높음(디버깅 용이)	환경 단순
GPU CUDA 단일	중간~높음	단일 GPU 자원 내 우수	중간(원자/비결정 가능)	중간(OOM/다이버전스 이슈)	드라이버/툴 의존
하이브리드(CPU+GPU)	높음	멀티 GPU+CPU 파이프라인 우수	중간	중간(복잡성↑)	오케스트레이션 필요

성능과 재현성 사이의 운영 선택

CSR 정렬과 정점 재배치(Reverse Cuthill–McKee, degree order)는 캐시와 메모리 결합을 개선할 수 있다. 대신 정렬 비용과 로딩 지연을 함께 감수해야 한다.

push와 pull의 전환 시점은 그래프와 레벨별 활성 간선 비율을 기준으로 동적으로 결정한다. 전환을 과도하게 반복하면 분기 비용이 증가한다.

idempotent 삽입 뒤 compaction을 수행하면 원자 연산을 줄일 수 있지만, 추가 메모리와 compaction 비용이 발생한다. 결과 재현성이 필요한 경우 정렬, 중복 제거, 고정 시드 스케줄을 적용하고 TEPS 저하를 감수할지 정책으로 정해야 한다.

대규모 그래프에서는 OOM 위험이 계속 존재한다. Unified Memory는 코드를 단순하게 만들 수 있지만 예측하기 어려운 페이지 폴트 비용이 생길 수 있다.

검증과 장애 대응을 운영 경로에 넣기

정확성은 소규모 그래프의 기준 결과와 GPU 결과를 비교하고, 방문 수와 거리의 단조 증가성을 검사해 확인한다.

성능 측정에서는 TEPS, 전송-커널 오버랩률, 워프 점유율, 메모리 결합률을 수집한다. 방향 전환 임계값은 이 지표를 바탕으로 튜닝 루프에서 조정한다.

커널 실패 코드를 검사하고 OOM이 발생하면 frontier 샤딩이나 스트리밍 모드로 전환한다. GPU를 사용할 수 없는 경우에는 CPU 폴백 경로를 유지한다. 드라이버와 런타임 버전은 고정하고, Nsight Systems/Compute로 병목을 관측하며 데이터 스키마·버전과 실행 아티팩트(커밋/빌드)를 추적한다.