그래프 탐색

15.1 배경: 그래프, 그리고 BFS의 위상

그래프(graph)는 정점(vertex)과 간선(edge)으로 이루어진 자료 구조다. 한 줄로 정의가 끝나지만, 그 위에서 다룰 수 있는 문제의 폭은 어마어마하다. 소셜 네트워크의 친구 추천, 길찾기 내비게이션, 검색 엔진의 PageRank, 코드 의존성 분석, 회로 시뮬레이션, 그리고 최근의 그래프 신경망(GNN)까지. 그래프 알고리즘은 컴퓨팅의 한 축이다.

가장 기초적이면서 가장 많이 쓰이는 연산이 그래프 탐색(graph traversal)이다. 시작 정점에서 출발해 도달 가능한 정점들을 체계적으로 방문하는 절차다. 두 가지 대표 방식이 있다. 너비 우선 탐색(BFS)은 거리(레벨)별로 동심원을 그리며 퍼져 나가고, 깊이 우선 탐색(DFS)은 한 갈래를 끝까지 따라 들어갔다 되돌아온다. GPU에서는 BFS가 압도적으로 자주 다뤄진다. DFS는 본질적으로 직렬에 가까워 병렬화가 어렵기 때문이다.

그래프를 메모리에 어떻게 담을지부터 정해야 한다. 인접 행렬(adjacency matrix)은 N × N 비트맵이라 N이 100만만 되어도 1 TB라 비현실적이다. 다행히 14장에서 봤던 CSR 형식이 그래프에도 그대로 들어맞는다. 사실 그래프의 인접 리스트(adjacency list)는 본질적으로 0/1 값을 갖는 희소 행렬의 CSR과 같다.

• row_ptr[u]: 정점 u의 이웃 리스트 시작 위치.
• col[]: 모든 간선의 목적지 정점.

// 4-노드 그래프: 0-1, 0-2, 1-2, 2-3
int row_ptr[5] = {0, 2, 4, 7, 8};
int col[8]     = {1,2,  0,2,  0,1,3,  2};
// 정점 2의 이웃: col[4..7) = {0, 1, 3}

BFS 자체는 친숙하다. 시작 정점을 큐에 넣고, 큐에서 하나 꺼내 미방문 이웃을 모두 큐에 넣는다. 큐가 빌 때까지 반복. 하지만 GPU에서 이걸 돌리려면 직렬 큐가 아닌 레벨 단위(level-synchronous) 방식으로 재구성해야 한다. 같은 거리(level)에 있는 정점들은 서로 독립적이므로 한꺼번에 병렬로 펼칠 수 있다는 관찰을 이용한다.

15.2 BFS의 레벨 단위 병렬화

레벨 단위 BFS는 다음과 같이 돈다. 거리 배열 dist[]를 모두 무한대로 초기화하고 시작점만 0으로 둔다. 매 반복(레벨) k마다, 거리가 k인 정점을 모두 찾아 그 이웃들의 거리를 검사한다. 미방문(거리 = 무한대)인 이웃에 거리 k+1을 적는다. 새로 거리가 갱신된 정점이 한 개라도 있으면 다음 레벨로 넘어간다. 새 갱신이 없으면 종료.

그림 15.1 — 레벨 단위 BFS. 각 레벨이 한 번의 커널 호출에 대응한다.

이제 한 레벨을 어떻게 병렬화할지가 핵심 질문이다. 크게 세 가지 전략이 있다. vertex-centric, edge-centric, 그리고 frontier 기반이다. 셋 다 흥미로운 trade-off를 가지고 있어 그래프 분포에 따라 우열이 갈린다.

15.3 vertex-centric 병렬화: push와 pull

가장 자연스러운 방식은 한 정점 = 한 스레드다. 매 레벨마다 N개의 스레드를 띄워 각자 자기 정점을 본다. 두 가지 방향이 가능하다.

Push 모드: 거리가 현재 레벨 k인 정점이 자기 이웃을 보면서, 미방문 이웃에 k+1을 써 준다. "내가 활성이니 너희도 깨워 줄게" 식이다.

__global__ void bfsPush(int N, int level,
                         const int* row_ptr,
                         const int* col,
                         int* dist,
                         int* changed) {
  int u = blockIdx.x * blockDim.x + threadIdx.x;
  if (u >= N) return;
  if (dist[u] != level) return;     // 활성 정점만 일한다
  int s = row_ptr[u], e = row_ptr[u + 1];
  for (int j = s; j < e; ++j) {
    int v = col[j];
    if (dist[v] == INT_MAX) {
      dist[v] = level + 1;          // race 가능, 같은 값이라 OK
      *changed = 1;
    }
  }
}

Pull 모드: 모든 정점이 자기 이웃을 보면서, 활성(거리 = k)인 이웃이 있으면 자기 거리를 k+1로 갱신한다. "내가 미방문이고 활성 친구가 있으면 깨어날게" 식이다.

두 방식의 차이를 본다. Push는 활성 정점이 적을 때 효율적이다. 일을 시작하는 스레드가 적기 때문. 반대로 모든 정점이 활성에 가까운 후반 레벨에서는 push가 노느니만 못하다 — 이미 방문된 이웃을 다시 검사한다. Pull은 활성 정점이 많을 때, 즉 BFS 중반에 더 효율적이다. 미방문 정점이 줄어드니 검사할 정점도 줄어든다.

15.4 edge-centric 병렬화: 간선 단위로 일하기

vertex-centric의 한 약점은 부하 불균형이다. 정점마다 차수(degree)가 다르니 한 워프 내에서도 어떤 스레드는 1000개의 이웃을, 어떤 스레드는 0개의 이웃을 다룬다. 차수 분포가 멱법칙(power-law)을 따르는 그래프(소셜, 웹)에서는 이게 치명적이다.

대안은 한 간선 = 한 스레드다. 모든 간선을 한 줄로 펼쳐 놓고 스레드를 균일하게 배정한다. 한 간선을 처리하는 스레드는 양 끝 정점의 거리를 본 뒤 갱신 조건이 맞으면 도착 정점의 거리를 갱신한다.

__global__ void bfsEdge(int nnz, int level,
                         const int* edge_src,    // 길이 nnz
                         const int* edge_dst,
                         int* dist,
                         int* changed) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= nnz) return;
  int u = edge_src[i], v = edge_dst[i];
  if (dist[u] == level && dist[v] == INT_MAX) {
    dist[v] = level + 1;
    *changed = 1;
  }
}

장점은 부하 균형이 거의 완벽하다는 점이다. 모든 스레드가 정확히 같은 일을 한다. 그러나 단점도 분명하다. 매 레벨마다 모든 간선을 검사한다. 활성 정점이 한 개여도 m개의 간선을 모두 훑는다. 멱법칙 그래프에서는 이게 누적되어 vertex-centric보다 느려질 수 있다. 또한 edge_src가 활성이 아닐 가능성이 높아 대부분의 검사가 헛수고가 된다(redundant work).

15.5 frontier로 work-efficient하게

두 방식의 약점을 한 방에 잡는 아이디어가 frontier 기반이다. 매 레벨에서 활성 정점만 작은 큐(frontier)에 모아 두고, 다음 레벨에서는 그 큐의 정점들만 펼친다. 이렇게 하면 work-efficient — 즉 알고리즘이 실제로 봐야 할 일만 본다.

그림 15.2 — frontier 기반 BFS. 활성 정점만 들고 다닌다.

구현은 다음과 같다. 두 개의 큐 frontier_in, frontier_out을 둔다. 한 레벨의 커널은 frontier_in의 각 정점 u에 대해 이웃을 검사하고 미방문이면 거리를 갱신하면서 frontier_out에 추가한다. 다음 레벨에서는 두 큐를 swap.

__global__ void bfsFrontier(
    int level,
    const int* frontier_in, int in_size,
    int* frontier_out, int* out_size,
    const int* row_ptr, const int* col,
    int* dist) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= in_size) return;
  int u = frontier_in[i];
  int s = row_ptr[u], e = row_ptr[u + 1];
  for (int j = s; j < e; ++j) {
    int v = col[j];
    // 거리 갱신을 atomicCAS로: 미방문일 때만 성공
    if (atomicCAS(&dist[v], INT_MAX, level + 1) == INT_MAX) {
      // frontier_out에 추가
      int pos = atomicAdd(out_size, 1);
      frontier_out[pos] = v;
    }
  }
}

이론적 work는 더할 나위 없이 효율적이다. 그러나 atomicAdd(out_size, ...)가 한 변수에 모든 스레드의 갱신을 몰아넣는다. 이게 contention의 원흉이다. 큰 프런티어를 만들 때 한 카운터에 수만 번의 atomic이 직렬화된다. SM마다 L2를 거쳐서 싱글 변수에 도달하는 길이 좁아져 처리량이 곤두박질친다.

15.6 사유화로 contention 줄이기

해법은 흔한 패턴 — 사유화(privatization)다. 블록마다 자체 로컬 큐를 셰어드 메모리에 두고, 블록 안의 스레드들은 atomic을 셰어드 메모리에서만 친다. 블록 끝에서 한 번만 글로벌 큐로 머지한다. 글로벌 atomic 횟수가 (스레드 수)에서 (블록 수)로 줄어든다. 1024 스레드 × 64 블록이라면, 65536번의 글로벌 atomic이 64번으로 줄어 1000배 감소.

__global__ void bfsFrontierPriv(
    int level,
    const int* frontier_in, int in_size,
    int* frontier_out, int* out_size,
    const int* row_ptr, const int* col,
    int* dist) {
  __shared__ int sQ[1024];
  __shared__ int sQSize;
  if (threadIdx.x == 0) sQSize = 0;
  __syncthreads();

  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < in_size) {
    int u = frontier_in[i];
    int s = row_ptr[u], e = row_ptr[u + 1];
    for (int j = s; j < e; ++j) {
      int v = col[j];
      if (atomicCAS(&dist[v], INT_MAX, level + 1) == INT_MAX) {
        int p = atomicAdd(&sQSize, 1);   // 셰어드 atomic
        if (p < 1024) sQ[p] = v;
      }
    }
  }
  __syncthreads();

  // 블록당 한 번 글로벌 큐에 머지
  __shared__ int blockBase;
  if (threadIdx.x == 0) {
    blockBase = atomicAdd(out_size, sQSize);
  }
  __syncthreads();
  for (int k = threadIdx.x; k < sQSize; k += blockDim.x) {
    frontier_out[blockBase + k] = sQ[k];
  }
}

로컬 큐가 가득 차면 어떻게 하느냐는 디테일이 따라붙는다. 보통 두 가지 안전장치가 들어간다. 첫째, 로컬 큐가 임계점에 닿으면 즉시 글로벌로 비우고 카운터를 리셋. 둘째, 로컬 큐 크기를 충분히 잡고 워프 단위 prefix sum으로 더 작은 단위에서 미리 정렬해 한 워프가 한 chunk를 글로벌에 코얼레싱으로 쓴다.

15.7 추가 최적화: direction-optimized와 부하 균형

실제 그래프 라이브러리(Gunrock, cuGraph, NVGRAPH)는 위 토대 위에 여러 층의 최적화를 쌓는다.

direction-optimized BFS: 매 레벨 시작 시 frontier 크기 |F|와 미방문 정점 수 |U|를 비교한다. |F| · 평균차수 > |U|이면 push 대신 pull로 전환한다. 활성 비율이 높을 때 pull이 압도적으로 유리하기 때문. 이 한 줄짜리 휴리스틱이 멱법칙 그래프에서 5~10배 가속을 준다.

부하 균형: 한 frontier 정점의 차수가 100만, 옆 정점의 차수가 1이면 한 워프 안에서도 다이버전스가 심하다. 해결책은 정점을 차수 구간별로 나누어 다른 커널 또는 다른 매핑 방식으로 처리하는 것이다(Gunrock의 LB(Load Balanced) 매핑). 작은 차수는 한 스레드가, 중간 차수는 한 워프가, 거대 차수는 한 블록이 협력해서 처리.

커널 호출 오버헤드: 레벨이 매우 깊은 그래프(체인 그래프 등)는 매 레벨마다 커널을 새로 띄워야 하는 부담이 누적된다. 한 호출의 launch latency가 5~10 마이크로초라면 1만 레벨에서 50~100ms의 오버헤드. 이를 줄이려면 CUDA Graph를 캡처하거나 cooperative groups의 grid 동기화로 한 커널 안에서 여러 레벨을 처리하는 기법을 쓴다.

15.8 정리

그래프 탐색은 GPU에 친절하지 않은 워크로드의 대표 사례다. 메모리 액세스가 무작위이고, 차수 분포가 편향되며, 레벨마다 일의 양이 변한다. 그래도 vertex-centric의 push/pull, edge-centric의 균일성, frontier의 work-efficiency, 그리고 사유화로 contention을 누르는 패턴을 조합하면 CPU 대비 5~30배의 가속을 안정적으로 얻는다. 이번 장의 패턴은 BFS만이 아니라 SSSP, connected components, PageRank, Betweenness Centrality 등 거의 모든 그래프 알고리즘에 그대로 응용된다.