SpMM - 3 : CSR 변환 GPU 병렬 알고리즘

Compressed Sparse Row Format

SpMM - 2 : Sparse Matrix Representation

저번 포스팅을 통해 sparse matrix를 저장하기 위한 다양한 형식을 배웠으며, 오늘은 CSR, Compressed Sparse Row format을 어떻게 구현하는 지에 대하여 정리할 것이다. 특히 [1]를 참고하여, GPU 디바이스를 통한 CSR 변환 가속화 코드를 구현할 것이다.

 

 

 

CSR 변환의 Sequential Code & 병렬화 전략

CSR은 rowPtr, col, value 배열만을 통해서 sparse matrix를 표현할 수 있다.

일반 행렬을 CSR 형식으로 변환하는 가장 쉬운 방법은 행렬의 첫 번째 요소부터 마지막 요소까지 Sequentially search하는 것일 것이다. Sequential search는 다음과 같은 two-pass로 진행될 것이다.

1. 첫 번째 요소부터 마지막 요소까지 search하여 총 non-zero의 개수를 샌다 -> col, value 배열을 동적 할당한다.
2. 다시, 첫 번째 요소부터 마지막 요소까지 search하며, non-zero가 나타나면 rowPtr, col, value 배열을 채운다.

우리의 목표는 위 sequential algorithm을 GPU로 병렬화하는 것이다.

(1) pass의 경우 parallel reduction을 통해 구할 수 있다. 이때 각 row를 하나의 CUDA warp만을 이용해 reduction할 것이다. 이 경우 warp primitive operation을 통해 매우 빠르게 reduction할 수 있다. 최종적으로 각 row별 non-zero 개수를 알 수 있다. row의 개수만큼의 숫자의 합을 구하면 되기에 CPU를 통해서도 빠르게 계산할 수 있다. 이때 exclusive scan을 하면 한번에 rowPtr도 구할 수 있다. n-번째 rowPtr, 즉 rowPtr[n]은 rowPtr[0] + ... + rowPtr[n-1]과 같기 때문이며, 행렬 A의 row 개수를 M이라 한다면 rowPtr[M] = rowPtr[0] + ... + rowPtr[M-1]이므로 rowPtr[M]은 총 non-zero의 개수를 나타낸다. (유의할 점은 row는 0-번째부터 (M-1)-번째까지만 있다. row의 개수가 M개이기 때문이다. 따라서 rowPtr[0]~rowPtr[M-1]까지만 실제 쓰인다. rowPtr[M]은 총 non-zero를 exclusive scan 한번으로 구하기 위해서 임의로 추가한 것이다.)

 

(2) pass의 경우 마찬가지로 각 row를 하나의 warp가 처리하도록 병렬화한다. 우리는 이미 rowPtr을 알기 때문에 각 row는 독립적으로 처리할 수 있다. Fig 1처럼 rowPtr이 주어지면 각 row에는 몇 개의 non-zero가 있는 지를 미리 알 수 있으며, 각 warp는 다른 warp와 독립적으로 자기가 채워 놓아야 할 col과 value의 index를 미리 알 수 있다.

 

Fig 1. rowPtr이 주어지면 각 warp가 해야할 일의 총 개수를 미리 알 수 있다.

 

다만 row level에서는 non-zero 끼리는 dependence가 있기에, 실제 쓰기를 할 때는 warp의 하나의 thread가 쓰기를 진행한다.

 

 

 

CSR 변환 병렬 알고리즘 정리

GPU를 이용한 CSR 변환 병렬 알고리즘은 rowPtr, col, value 배열을 생성하기 위해 다음과 같은 3가지 단계로 구성된다.

1. 행렬 A[M,N] 의 모든 row에 대하여, 각 row의 non-zero 개수 구하기 (GPU에서 실행) 
   • 하나의 warp(32 threads)가 한 row를 맡음
   • 32 threads가 coalesced memory access를 통해 행렬(글로벌 메모리)을 빠르게 읽어옴
   • __ballot_sync()(warp level primitive)를 통해 32 threads 중 몇개가 non-zero인지를 한 operation으로 알아냄
   • 하나의 threads가 그 결과를 rowPtr에 저장함
   • 그 결과, (M+1)개의 원소를 가지는 rowPtr 배열 생성. 아직까지의 rowPtr은 partial result (완벽한 정답이 아님)
2. rowPtr 배열에 대하여 exclusive scan (CPU에서 실행)
   • rowPtr 배열의 0~(M-1)번 원소는 해당 row 이전까지의 총 non-zero의 개수를 의미함
   • rowPtr 배열의 마지막(M번) 원소는 총 non-zero의 개수를 의미함
   • 총 non-zero 개수를 통해 col, value 배열을 동적 할당함
3. column 및 value 정보 저장 (GPU에서 실행)
   
   • 하나의 warp가 한 row를 맡아서 처리함. n-번 row는 rowPtr[n] 이상 rowPtr[n+1] 미만 index를 가지는 non-zero 처리를 맡음
   • 32 threads가 coalesced memory access를 하여 행렬(글로벌 메모리)을 빠르게 읽어 옴
   • 그 결과를 shared memory에 저장함
   • 하나의 threads가 shared memory에 저장된 값들 중 non-zero를 순서대로 search하여 col, value에 저장

 

 

 

구현 코드

GitHub - jhson989/SpMM: Parallel Sparse Matrix Multiplication via SYCL

 

 

 

Reference

[1] Julien Demouth, Sparse Matrix-Matrix Multiplication on the GPU, NVIDIA, GTC 2012