Post List

2020년 6월 1일 월요일

GPU에서의 희소행렬 곱셈(spGEMM)(3) 수정중

희소행렬이 무엇인지, 희소행렬 곱셈이 어떻게 이루어지는지는 사실 생각보다 간단하다. 하지만 실제 데이터에 대해 연산을 한다면 이런저런 문제가 발생한다.

GPU architecture

일단 GPU를 이용해 연산한다는 가정을 깔고 있기 때문에 간단하게 복기하자면 GPU의 가장 큰 특징은 코어가 여러개란 것이다. 따라서 병렬처리에 용이하지만 만약 자원을 효율적으로 사용하지 못한다면 CPU를 이용해 계산하는 것만 못할 수 있다. CUDA terminology를 빌려 쓰자면 Thread들의 집합인 Thread block 이 GPU의 SM(Streaming Multiprocessor)에 스케줄링 되고, 블록내의 Thread들은 SP(Streaming Processor)에 의해 수행된다. 이때 32개의 Thread 단위로 같은 instruction을 수행하는데 이를 Warp이라고 부르며 GPU의 기본 수행단위이다. GPU 메모리는 크게 Global memory, L2, L1 cache, shared memory, register file로 나뉘어지는데 최대한 Global memory access를 줄이는 것이 유리하다.

Load balancing problem

희소행렬 곱셈의 수행시간은 사실 인풋행렬의 형태에 의해 크게 좌우된다. 만약 행렬의 banded matrix를 띄거나 랜덤분포를 이루면 큰 문제가 없겠지만 어떤 행렬들은 power-law degree distribution이라는 특성을 갖는다. 간단하게 말하면 행렬의 특정 row나 column에 0이 아닌 원소들이 몰려있는 분포이다. 이럴 경우 어떤 문제가 발생할까?
Row-wise product-based spGEMM의 경우 각 행마다 원소의 개수가 크게 차이나게 되면 Thread 마다 workload가 차이나게 된다. 물론 이를 해결하기 위해 따로 Thread들의 작업량을 균등 분배하기 위해 전처리를 가할 수 있지만 "전처리 시간 + 추가 memory access cost"를 고려한다면 효과가 미미하거나 오히려 손해를 볼 가능성이 있다.
Outer product-based spGEMM의 경우엔 column의 각 element에 곱해지는 B의 벡터가 같기 때문에 Thread간 load imbalance 문제는 발생하지 않는다. 하지만 power-law 특성이 더 강해지면 Outer product에선 문제가 더 심화된다.
Power-law 특성이 강해지면 Row-wise나 Outer product 방식이나 두 방식 모두 load balancing 문제가 발생하게 되지만 Outer-product의 경우 column의 원소가 모두 동일한 벡터와 곱해진다. 따라서 column, row 모두 0이 아닌 원소가 많은 밀도 높은 벡터일 경우 workload amplification이 일어난다. 즉 Thread 간 workload는 동일하지만 Thread block 간 심각한 load imbalance가 발생하는 것이다. 보통의 경우 Thread block이 충분히 많다면 Round-Robin 스케줄링으로 격차가 메워지지만 Power-law가 심한 경우엔 심각한 성능저하가 발생한다. 이를 방지하기 위해 벡터를 쪼개는 방식을 이용한다.

Too Sparse Matrix

인풋 행렬이 Power-law 특성을 가진다는 것은 몇몇 노드들은 0이 아닌 원소가 굉장히 많다는 것을 의미하면서 동시에 대부분의 노드들은 0이 아닌 원소가 거의 없다는 것을 뜻한다. GPU에서 가장 작은 수행단위가 32개의 Thread들의 집합인 warp이기 때문에 특별한 전처리를 하지 않는다면 각 연산마다 적어도 32개의 0이 아닌 원소가 필요하다. 하지만 데이터를 분석해보면 대부분의 노드의 차수(0이 아닌 원소의 개수)가 32도 안되는 경우가 많다. 이 경우 GPU의 자원이 낭비된다. Thread 가 instruction을 수행해야 하지만 issue 가능한 instruction이 없기 때문에 첫째로는 sp 코어가 놀게된다. 또 GPU는 많은 양의 register file로 각 Thread가 필요한 data들을 load 해놓고 stall이 필요할 때 context-switching을 이용하여 latency-hiding을 하지만 이 경우엔 latency-hiding을 할 수 없게된다. 이 문제를 해결하는 방법은 여러가지겠지만 Thread coarsening 이라는 직관적이고 간단한 방법으로 전처리 cost가 거의 없이 어느정도 완화할 수 있다.

Output management

희소행렬 곱셈을 수행할 때 생기는 문제들은 이것저것 있지만 사실 가장 큰 문제는 너무 큰 아웃풋 사이즈, 그리고 아웃풋을 관리하는 것이다. 첫째로 아웃풋의 원소가 몇개생길지 계산해보기 전까진 알 방법이 없다. 따라서 미리 아웃풋의 메모리를 할당받아야 하는 GPU의 경우 최대 얼만큼의 메모리가 필요한지 계산하는 과정이 필요하다. 이 과정을 거쳐야 2번째 포스팅 코드의 offset을 구할 수 있다. 두번째로 곱셈계산을 끝내면 위치가 중복된 원소들이 존재한다. Row-wise 방식의 경우 Output Matrix C의 각 행마다 겹치는 원소가 존재하고 Outer-product 방식의 경우 판단위로 겹치는 원소가 존재하는데 결국 이 원소들을 전부 합쳐서 1개의 원소로 만들어야 곱셈연산이 완벽하게 끝난 것이다. 그럼 두가지 선택을 할 수 있는데 1) 모든 원소가 계산된 후 Merge Process를 수행한다. 와 2) 계산 중 겹치는 원소들을 Merge한다. 가 있다. 개인적으로 Row-wise product의 경우 2), Outer product의 경우 1)로 노선을 잡는게 맞는 것 같다. Row-wise product의 경우는 행단위로 merge process가 이루어지기 때문에 겹치는 원소들이 이미 같은 SM에 존재하지만 Outer product의 경우 판단위 merge process가 필요하기 때문에 겹칠 원소들을 미리 모아놓지 않는다면 SM간 communication cost가 필요하고 굉장히 알고리즘적으로 복잡할 것이기 때문이다. 아웃풋을 어떻게 관리할지는 아직 연구가 필요한 부분이라고 생각되며 뚜렷한 해결책이 생각나지 않는다. 사실 merge process는 결국 memory transaction이 굉장히 많이 일어나기 때문에 최적화의 여지가 별로 없다. merge 될 원소와 merge가 필요 없는 원소를 따로 분리하고 memory 접근을 줄이는 방향으로 생각해봤지만 사실 어떤 원소가 merge될지 안될지는 찍기나 다름없다.

이외에도 matrix내의 locality를 증가시킨다거나 하는 접근법이 있긴하지만 1) 전처리 시간이 오래걸린다는 점과 random 분포에선 사실상 NP problem이기 때문에 이는 고려하지 않겠다. 희소행렬 곱셈에 대해 연구하다보면 이런 문제들이 존재하고 이를 해결할 수 있는 접근법이 필요해진다.

2020년 5월 27일 수요일

GPU에서의 희소행렬 곱셈(spGEMM)(2)

저번 포스팅에서는 간단하게 희소행렬을 저장할 때 사용되는 포맷에 대해 올렸다. 간단히 복기하자면 메모리 효율성을 위해 COO(Coordinate list ;Triplet) 포맷이나 CSR, CSC(Compressed Sparse Row/Column) 포맷을 이용한다는 것이다.

희소행렬 곱셈 알고리즘

보통 '행렬곱셈' 하면 내적을 이용한 곱셈을 떠올리지만 희소행렬 곱셈에서는 내적방식이 잘 쓰이지 않는다. 만약 행렬이 2D 그 이유는 내적연산을 위해선 두 벡터의 "인덱스가 같은 원소"를 곱해줘야 하지만 희소행렬은 각 행값, 또는 열값이 idx array에 저장되어있다. 따라서 idx array를 읽고 두 index가 같은지 체크해줘야하는 비효율이 생긴다.

Row-wise product

희소행렬 곱셈을 수행할 때 row-wise(row-row) product이 널리 쓰인다. Nvidia에서 제공하는 cuSPARSE, CUSP 역시 이 방식으로 곱셈을 진행한다. 구체적인 알고리즘은 위와 같이 A의 행의 0이 아닌 원소와 그 원소의 열값에 해당하는 B의 행 전체를 곱하고 결과로 C의 행 1개가 생성된다. A, B 모두 원소를 row-wise로 접근하기 때문에 두 행렬 모두 CSR 포맷을 이용한다. 간단히 Naive 한 GPU 코드를 작성해보면 다음과 같다.

// row-wise product based spGEMM
__global__ void spGEMM_expansion(
    int* a_ptr, int* a_idx, double* a_val,
    int* b_ptr, int* b_idx, double* b_val,
    int* c_ptr, int* c_idx, double* c_val)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    
    for(int a_iterator = a_ptr[bid]; a_iterator < a_ptr[bid + 1]; ++a_iterator)
    {
        int a_col = a_iter[a_iterator];
        
        for(int b_iterator = b_ptr[a_col]+tid; b_iterator < b_ptr[a_col + 1]; b_iterator+=blockDim.x)
        {
            // Merge?
            c_idx[offset] = b_idx[b_iterator];
            c_val[offset] = a_val[a_iterator]*b_val[b_iterator];   
        }
    }
}

간단하게 1개의 Thread block이 C의 행 1개를 생성하는 코드이다. 알고리즘 자체는 굉장히 간단하지만 몇가지 문제들이 발생한다. 이슈는 이후에 언급하도록 하겠다.

Outer product

row-wise(row-row) product 방식 이외에는 outer-product(column-row) 방식이 쓰인다. 구체적인 알고리즘은 위와 같이 A의 열과 B의 행의 곱셈으로 C의 일부가 생성된다. row-wise product 방식과 차이점은 연산의 결과가 전체 행렬 한판이라는 것이다. 이는 꽤나 중요한 차이지만 역시 이후에 언급. A의 access pattern은 column-wise, B는 row-wise이기 때문에 A는 CSC 포맷, B는 CSR 포맷을 이용한다. 간단히 Naive 한 GPU 코드를 작성해보면 다음과 같다.

// outer product based spGEMM
__global__ void spGEMM_expansion(
    int* a_ptr, int* a_idx, double* a_val,
    int* b_ptr, int* b_idx, double* b_val,
    int* c_ptr, int* c_idx, double* c_val)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    
    for(int a_iterator = a_ptr[bid]; a_iterator < a_ptr[bid + 1]; ++a_iterator)
    {
        int a_col = a_iter[a_iterator];
        
        for(int b_iterator = b_ptr[bid]+tid; b_iterator < b_ptr[bid + 1]; b_iterator+=blockDim.x)
        {
            // Merge?
            c_idx[offset] = b_idx[b_iterator];
            c_val[offset] = a_val[a_iterator]*b_val[b_iterator];   
        }
    }
}

간단하게 1개의 Thread block이 C행렬의 일부가 되는 C 한 판을 생성하게 된다.

2020년 5월 21일 목요일

GPU에서의 희소행렬 곱셈(spGEMM)(1)

컴퓨터좀 만져본 사람중 행렬 곱셈 코드를 안짜본 사람은 없다.
나도 이제 깨우치는 중이지만 상당히 많은 데이터가 행렬의 형태를 띄고, 많은 문제들이 수학적으로 모델링 될 때 행렬곱셈 연산이 포함되기 때문에 행렬 곱셈은 컴퓨터공학도에게는 꽤 중요한 연산중 하나다.

희소행렬

최근 데이터가 커지고 네트워크가 희소해지면서 희소행렬에 대한 연산이 중요해지고 있는데 희소행렬은 일반적인 행렬과 달리 2D array의 형태로 저장되지 않는다.
이유는 메모리 효율성 때문이다. 행렬이 N x N 이라면 2D array의 형태로 저장하기 위해선 N2의 메모리공간이 필요하다.

Space complexity for 2D array : \(O(N^{2})\)

따라서 스케일이 크고 희소한 행렬은 크게 두가지 포맷을 이용한다.

COO format
각각의 원소는 (row, col, value)로 표현된다. 위 행렬을 coo 포맷으로 저장하면 다음과 같다.


typedef struct _coo {
    int row;
    int col;
    double val;
}coo;

typedef struct _spmat{
    coo* elem;
}spmat;

int main(){
    spmat A;
    A.elem = (double*)malloc(sizeof(double)*4);
    A.elem[0].row = 1; A.elem[0].col = 0; A.elem[0].val = 5;
    A.elem[1].row = 1; A.elem[1].col = 1; A.elem[1].val = 8;
    A.elem[2].row = 2; A.elem[2].col = 2; A.elem[2].val = 3;
    A.elem[3].row = 3; A.elem[3].col = 1; A.elem[3].val = 6;
    return 0;
}
희소행렬을 COO format으로 저장한다면 행렬원소의 개수가 E일 경우 필요한 메모리 공간은 대략 3E이다.

Space complexity for COO : \(O(3E)\)

COO 포맷은 굉장히 직관적이고 어느정도의 메모리 효율성을 보장하지만 실제 연산에서는 자주 이용되지 않는다. 이유는 각 행이나 열에 direct access 할 수가 없기 때문이다.
실제 연산에서는 CSR(Compressed Sparse Row), CSC(Compressed Sparse Column) 포맷을 이용한다.
Compressed sparse format 는 3개의 array(ptr, idx , val)로 이루어져 있다.
CSR 포맷을 예로 들어보면

ptr array는 각 행의 첫 원소의 위치를 저장하고
idx array는 각 원소의 col을 저장하고
val array는 각 원소의 value를 저장한다.

COO 포맷과 다른 것은 ptr array 뿐이다.

#define ROW_MAJOR 0
#define COL_MAJOR 1
typedef struct _cs {
    int number_of_rows;
    int number_of_cols;
    int number_of_elems;
    int order;
    int* ptr;
    int* idx;
    double* val;
}cs;

typedef struct _spmat{
    coo* elem;
    cs csr;
    cs csc;
}spmat;

int main(){
    spmat A;
    // CSR representation
    A.csr.order = ROW_MAJOR;
    A.csr.number_of_rows = 4;   
    A.csr.number_of_cols = 4;
    A.csr.number_of_elems = 4;
    A.csr.ptr = (int*)malloc(sizeof(int)*(number_of_rows+1));
    A.csr.idx = (int*)malloc(sizeof(int)*(number_of_elems));
    A.csr.val = (double*)malloc(sizeof(double)*(number_of_elems));
    
    A.csr.ptr[0] = 0; // ptr[0] is always zero
    A.csr.ptr[1] = 0; // because there's no nonzero element in row 0
    A.csr.ptr[2] = 2; // two nonzero elements in row 1
    A.csr.ptr[3] = 3; // one nonzero elements in row 2
    A.csr.ptr[4] = 4; // one nonzero elements in row 3
    
    A.csr.idx[0] = 0;
    A.csr.idx[1] = 1;
    A.csr.idx[2] = 2;
    A.csr.idx[3] = 1;

    A.csr.val[0] = 5;
    A.csr.val[1] = 8;
    A.csr.val[2] = 3;
    A.csr.val[3] = 6
    return 0;
}
만약 위 행렬에서 row 1 에 대한 탐색을 하고싶다면 A.csr.ptr[1]을 참조하고 A.csr.ptr[2]-A.csr.ptr[1] 개의 원소를 탐색하면 된다 CSC 포맷은 CSR 포맷과 구조는 같지만 column-major order로 저장되고 ptr는 각 열의 첫번째 원소의 위치, idx array는 행 값을 저장한다.
희소행렬을 CSR/CSC format으로 저장한다면 행렬원소의 개수가 E, 행렬의 크기가 NxN일 경우 필요한 메모리 공간은 N+2E이다.

Space complexity for CSC,CSR : \(O(N+2E)\)

희소행렬 데이터 소스 : https://sparse.tamu.edu/