ISSN 1598-17 (Print) ISSN 2287-1136 (Online) http://www.jksii.or.kr GP-GPU 의캐시메모리를활용하기위한병렬블록 LU 분해프로그램의구현 Implementation of parallel blocked LU decomposition program for utilizing cache memory on GP-GPUs 김영태 1* 김두한 1 유명한 1 Youngtae Kim Doo-han Kim Myoung-han Yu 요 약 GP-GPU 는그래픽처리를위한 GPU 의다중쓰레드를일반수치계산에활용하여초고속으로계산하는장치이다. GP-GPU 에서는 CPU 의캐시메모리와는달리다중쓰레드가공유하는공유메모리의형태로캐시메모리를제공하며, 공유메모리는사용자프로그램에서직접제어할수있다. 본연구에서는 GP-GPU 의캐시메모리를사용하여계산성능을향상시키기위한블록구조의병렬 LU 분해프로그램을구현하였다. Nvidia CUDA C 로구현된병렬블록 LU 분해프로그램은동일한 GP-GPU 상에서일반 LU 분해프로그램에비교하여 7~8 배이상의속도개선을보였다. 주제어 : LU 분해프로그램, GP-GPU, Nvidia CUDA, 병렬프로그램 ABSTRACT GP-GPUs are general purposed GPUs for numerical computation based on multiple threads which are originally for graphic processing. GP-GPUs provide cache memory in a form of shared memory which user programs can access directly, unlikely typical cache memory. In this research, we implemented the parallel block LU decomposition program to utilize cache memory in GP-GPUs. The parallel blocked LU decomposition program designed with Nvidia CUDA C run 7~8 times faster than nun-blocked LU decomposition program in the same GP-GPU computation environment. keyword : LU decomposition, GP-GPU, Nvidia CUDA, Parallel program 1. 서론 GP-GPU(General Purposed Graphic Processing Unit) 는그래픽계산장치를일반계산에사용하는장치로서, 격자구조로배열된다중프로세서를사용하여행렬을병렬로계산할수있는저전력, 고성능의초고속계산장치이다 [1][2][3][4]. 캐시메모리는메인메모리보다작지만속도면에서현저히빠른메모리로서메인메모리에상주하는데이터의일부를저장하여프로세서가메인메모리를접근하지않고직접데이터를접근할수있는저장장치이다. 한편 GP-GPU 는수백개의다중쓰레드를사용하는병렬계산의특성상시스템을통하여캐시메모리를자동으로 1 Department of Computer Science, Gangneung-Wonju National University, Wonju, 22-711, Korea. * Corresponding author (ykim@gwnu.ac.kr) [Received 1 August 213, Reviewed 5 August 213, Accepted 14 October 213] 사용하기가어렵다. 따라서일반적인 CPU와는달리응용프로그램에서직접사용할수있는 ( 다중쓰레드들이공유하는 ) 공유메모리의형태로캐시메모리를제공한다. 캐시메모리는메모리처리속도를줄임으로써계산성능을향상시키기때문에프로그램에서캐시메모리를최대한활용하는것이성능에있어서매우중요하다 [5]. 본연구에서는 Nvidia GP-GPU의캐시메모리 ( 이하공유메모리 ) 를사용하여프로그램의계산시간을단축하기위하여병렬블록 LU 분해프로그램을구현하였다. LU 분해프로그램은 2차원행렬을하삼각행렬 (L) 과상삼각행렬 (U) 의곱으로표현하는프로그램으로서연립방정식을계산하기위한가우스소거법, 역행렬및행렬의 determinant 등을계산할수있는선형대수의가장기본적인프로그램중의하나이다 [6]. 구현된병렬블록 LU 분해프로그램은블록을사용하지않는 LU 분해프로그램과동일한복잡도인 O(⅓n 3 ) 와계산결과를갖지만, 행렬을블록으로나누어서삼각행렬의시스템해법계산 Journal of Internet Computing and Services(JICS) 213. Dec: 14(6): 41-47 41 http://dx.doi.org/1.7472/jksii.213.14.6.41
및행렬의곱을사용하기때문에이과정에서나누어진블록을공유메모리에저장하여계산시간을단축할수가있다. 본연구에서의병렬블록 LU 분해프로그램은동일한 GP-GPU 계산환경에서기존 LU 분해프로그램과비교하여괄목할만한속도개선을이루었다. 본논문에서는 2장에서병렬블록 LU 분해프로그램에대해설명하고, 3장에서는병렬블록 LU 분해프로그램의성능을기존의 LU 분해프로그램과비교하여속도개선에대하여알아보며, 4장에서는결론으로맺는다. 2. 병렬블록 LU 분해프로그램 2.1 GP-GPU 와 CUDA GP-GPU 는그래픽을처리하기위한장치로서 수백개의코어들이격자구조로배치되어있으며 이를다중쓰레드를통하여실행하기때문에행렬 의계산에있어서매우효율적인구조이다. 그림 1 은 CPU 와 GP-GPU 구조의차이점을보여준다. (a) 에서와같이멀티코어로이루어진 CPU 는각코어 가별도의캐시메모리와연결되어있지만 (b) 에서 의 GP-GPU 에서는여러개의코어들이사용할수 있는메인메모리인 global 메모리와쓰레드블록이 사용할수있는별도의캐시메모리를제공한다. (a) CPU 의구조 (b) GP-GPU 의구조 ( 그림 1) CPU 와 GP-GPU 의구조비교 (Figure 1) Structure comparison of CPU and GP-GPU 일반적으로 CPU 에서는운영체제를통하여캐시메모 리를자동으로관리하는반면에 GP-GPU 에서는운영체제 에서각각의코어에캐시메모리를자동으로처리하기가 어렵기때문에응용프로그램에서모든코어들이공유하 는공유메모리형태의캐시메모리를직접사용하게된다. 따라서응용프로그램에서공유메모리를효율적으로활 용하는것이전체계산성능에매우중요한역할을한다. 참고로, [1] 에서제안한행렬의곱프로그램을구현하여비교한결과, GP-GPU 의공유메모리를사용한경우가 global 메모리를사용한경우와비교하여최대 9배의계산속도개선의차이를보였다. GP-GPU 를처리할수있는프로그램언어로는 OpenCL, OpenGL, CUDA(Computer United Device Architecture) 등을사용할수있다 [1]. 이중 Nvidia에서개발한 CUDA 모델은고급프로그램언어인 C와 Fortran 등을제공하기때문에일반계산에서가장많이사용하는방식으로본연구에서도 CUDA C 프로그램언어를사용하였다. 2.2 Nvidia CUDA 환경에서의 LU 분해프로그램 LU 분해프로그램은주어진행렬 A에대해 A=LU 형태로하삼각행렬 L과상삼각행렬 U의두개의행렬로나누는프로그램이다. LU 분해프로그램은연립방정식의해를구하는가우스소거법과거의동일한알고리즘을사용하지만분해된 L 및 U 삼각행렬을사용하여연립방정식의해뿐만아니라역행렬및행렬의 determinant 등을계산할수있기때문에선형대수분야에서는가장대표적인프로그램중의하나이다 [6]. 본연구에서는 LU 분해프로그램구현을위해 GP-GPU에서의프로그래밍언어로 CUDA C를사용하였다. CUDA C에서는먼저동시에계산하고자하는쓰레드블록을쓰레드의개수로설정하고, 전체계산도메인을나타내는그리드는쓰레드블록으로나뉘어져해당영역들이순차적으로반복계산한다. 이때그리드를계산하기위한쓰레드블록들의계산순서는임의로결정된다 [1]. 다음은 LU 분해프로그램에서쓰레드블록과그리드를정의하는코드이다. 코드에서그리드의크기를하나의도메인으로정의하는이유는데이터의존도를해결하기위한블록의계산순서를프로그램에서결정하기위한것이며이방식은블록을사용하지않는병렬 LU 분해프로그램을위하여 [7] 에서제안한바있으며병렬블록 LU 분해프로그램에서도동일한방식을사용한다. dim3 dimblock(num_thread, num_thread); dim3 dimgrid(1, 1); 그림 2는블록을사용하지않는 LU 분해프로그램의병렬성을보여준다. 그림에서와같이프로그램의대각선의원소들을기준으로해당열을수정하고이를사용하 42 213. 12
여화살표가표현하는바와같이부분도메인의원소들 을수정하게되며반복이진행되면서계산하는행렬의 크기는작아지게된다 [7][8]. 형태를보여준다. 반복문의수는블록의수이며반복문 내에서의 LU 분해프로그램은블록을사용하지않은 LU 프로그램을말하며부분피벗팅을포함한다. for(b = ; B < nblocks; B++){ 블록을사용하지않는 LU 분해 ; 삼각행렬시스템의해 ; 행렬의곱 ; ( 그림 2) LU 분해프로그램 (Figure 2) LU decomposition program 다음코드는그림 2 의방식을구현한 LU 분해프로그 램이다. 코드에서 1 은열을수정하고, 2 에서는부분도 메인을수정한다. CUDA C 프로그램에서는행렬의각원 소에대한반복문이묵시적으로이루어져서반복문이없 지만, 이코드에서 i 와 j 에대한반복문이존재하는이유 는앞에서설명한바와같이데이터의의존성를해결하 기위하여명시적반복문을통하여부분행렬을순서대 로계산하기때문이다. 그림 3은병렬블록 LU 분해프로그램에서의 LU 분해와삼각행혈시스템해법의알고리즘을보여준다. 그림에서볼수있듯이그림 2의 LU 분해프로그램과는달리각원소를계산하는것이아니라블록단위로계산하게된다. 그림에서 1은블록을사용하지않는 LU 프로그램의행렬이며 2는삼각행렬의시스템해에서의행을보여준다. 이때계산에사용되는행렬의일부분을공유메모리에저장하여사용하므로성능을향상시킨다. for(k=; k<n; k++) { for(i=k+1; i<n; i++) { A(i,k) = A(i,k)/A(k,k);...1 for(j = k+1;j<n; j++) { A(i,j) = A(i,j) - A(k,j)*A(i,k);...2 2.3 병렬블록 LU 분해프로그램 블록 LU 분해프로그램은주어진행렬을균등한블록 으로나누어서블록단위로삼각행렬의시스템해법및 행렬의곱프로그램을사용하여계산하는방식이다 [8][1]. 이때블록의크기는임의로결정할수있는데본 연구에서는공유메모리를최적으로사용하기위하여블 록의크기를 GP-GPU 에서의쓰레드블록의크기로정하 였다. 다음의알고리즘은블록 LU 분해프로그램의기본 ( 그림 3) LU 분해와삼각행렬시스템해법 (Figure 3) LU decomposition and triangular system solver 다음 CUDA C 프로그램은그림 3 에서 2 의삼각행렬 시스템해프로그램이다. 프로그램에서 scol 로표현된 1 차원공유메모리의행렬에해당행을 1 과같이저장하 고이를쓰레드블록이공유하여다음계산인 2 에서사 용할수있다. 코드에서 i 와 j 에대한반복문이없는이유 는 CUDA C 에서는반복문을생략하기때문이다 [1]. i = B*blockDim.x + threadidx.x; for(kk=; kk<blockdim.x; kk++) { k = B*blockDim.x + kk; scol(threadidx.x) = A(i,k);...1 한국인터넷정보학회 (14 권 6 호 ) 43
for(bi=b+1; bi<nblocks; bi++) { j = bi*blockdim.y + threadidx.y; if (j > k) { if (i > k) { A(i,j) = A(i,j)-A(k,j)*scol(threadIdx.x);..2 j = bj*blockdim.y + threadidx.y; s_a(threadidx.x,threadidx.y) = A(B*blockDim.x+threadIdx.x,j);...2 for(k=; k<blockdim.x; k++) { A(i,j) -= s_a(threadidx.x,k) * s_b(k,threadidx.y);...3 그림 4는블록 LU 분해프로그램에서의행렬의곱프로그램을보여준다. 3에서의화살표가표시하는데로각블록은행렬의곱을통하여수정이되며, 이때계산에사용되는부분행렬은행렬을수정할때반복하여사용이되기때문에해당블록을 2차원의공유메모리에저장하여성능을향상시킬수있다. 2.4 부분피벗팅부분피벗팅은 LU 분해프로그램의반복에서해당열에서절대값이가장큰행을축이되는행과교환함으로써계산결과를좀더안정되게한다 [6]. 블록 LU 분해프로그램에서의부분피벗팅은 [7] 에서의방식과동일한방식을사용하였다. 3. 성능평가 3.1 시스템사양 ( 그림 4) 블록 LU 에서의행렬의곱 (Figure 4) Matrix multiplication in the blocked LU decomposition 다음프로그램은그림 4 에서의알고리즘을구현한 CUDA C 프로그램이다. 프로그램에서 1 과 2 에서와같 이해당블록을 2 차원공유메모리인 s_b 와 s_a 에각각 저장을하고이후에 3 에서의행렬의곱프로그램에서 사용하게된다. for(bi=b+1; bi<nblocks; bi++) { i = bi*blockdim.x + threadidx.x; s_b(threadidx.x,threadidx.y) = A(i,B*blockDim.y+threadIdx.y);...1 for(bj=b+1; bj<nblocks; bj++) { 병렬블록 LU 분해프로그램의성능평가를위해두 개의다른 GP-GPU 를사용하였으며사양은표 1 과같다. Nvidia Tesla C16 의프로세서의수는 3 개이며각프로 세서는 8 개의코어로구성이되어있고최대한 512 쓰레 드를동시에사용할수있다. 반면에 Tesla C275 는 8 개의 코어로구성이된 14 개의프로세서로서최대한 1,24 개의 쓰레드를사용할수있다. 다중쓰레드가사용할수있는 캐시 ( 공유 ) 메모리는 C16 이 16K bytes 이고 C275 가 48K bytes 이다. Tesla C275 가 Tesla C16 보다코어의수는적 지만, 최대병렬쓰레드의수가많고메모리처리속도가 향상이되어계산량이많은경우에보다효율적인성능을 보인다. 본논문에서는성능평가를위하여두 GP-GPU 를 사용함으로써성능의이식성을보이고자한다. ( 표 1) 성능비교에사용된 GP-GPU 의사양 (Table 1) System specification of GP-GPUs used in performance comparison Tesla C16 Tesla C275 Clock frequency 1.3 GHz 1.15 GHz multiprocessors의수 3 개 14 개 cores의수 24 개 112 개 최대쓰레드블록의수 512 개 124 개 ( 블록당 ) 공유메모리의크기 16K bytes 48K byres global memory의크기 4G bytes 4G bytes 44 213. 12
3.2 성능비교병렬블록 LU 분해프로그램의성능평가를위하여블록을사용하기않은병렬 LU 분해프로그램과실행시간을비교하였다. 성능의비교를위하여 124 124, 248 248, 372 372, 496 496 크기의행렬에대해각각쓰레드의수를다르게설정하여실행시간을측정하였다. 계산에는모두부분피벗팅을포함하였다. 그림 5는 Tesla C16에서의병렬블록 LU 분해프로그램의성능을블록을사용하지않는병렬 LU 분해프로그램과비교하였다. 그림에서 (a) 는 64(=8x8) 개의쓰레드를사용한경우이며 (b) 는 256(=16x16) 개의쓰레드를사용한경우이다. 그림에서볼수있듯이쓰레드블록이클수록성능의개선은향상되며, 256개의쓰레드를사용할경우 ( 그림 (b) 참조 ) 에는최대 9배정도의향상된성능을보인다. 그림 6은 Tesla C275를사용하여병렬 LU 분해프로그램들의성능을비교하였다. 그림5에서와유사하게병렬블록 LU 분해프로그램의향상된성능을보였으며최대 8배정도의성능을보인다. Performance Comparison (Thread block: 8x8) Performance Comparison(Thread block: 8x8) (a) 64(=8x8) 쓰레드를사용한계산시간의비교 Performance Comparison (Thread block: 16x16) (a) 64(=8x8) 쓰레드를사용한계산시간의비교 Performance Comparison(Thread block: 16x16) (b) 256(=16x16) 쓰레드를사용한계산시간의비교 ( 그림 5) 블록 LU 분해프로그램과일반 LU 분해프로그램의성능비교 (Tesla C16) (Figure 5) Performance comparison of the blocked LU decomposition and the non-blocked LU decomposition programs(tesla C16) (b) 256(=16x16) 쓰레드를사용한계산시간의비교 ( 그림 6) 블록 LU 분해프로그램과일반 LU 분해프로그램의성능비교 (Tesla C275) (Figure 6) Performance comparison of the blocked LU decomposition and the non-blocked LU decomposition programs(tesla C275) 그림 5 와 6 에서나타나듯이캐시메모리를활용한블록 LU 분해프로그램의성능이기존의 LU 분해프로그램에 비하여뛰어나게향상된것을볼수있으며, 쓰레드블 록의크기가클수록캐시메모리의사용이증가하여성능 의개선은향상되었다. 두 GP-GPU 를비교하면 Tesla C275 가 C16 보다계산속도는빠르지만성능개선의 한국인터넷정보학회 (14 권 6 호 ) 45
효과는비슷하게나타난다. 4. 결론 GP-GPU는저비용및저전력의고성능병렬계산장치로서많은활용이기대되는차세대계산장치이다. 본연구에서는 Nvidia GP-GPU에서캐시메모리를사용하여성능을높이기위하여블록 LU 분해프로그램을구현하였다. 프로그래밍언어는 CUDA C를사용하였으며구현된병렬블록 LU 분해프로그램은동일한계산환경에서블록을사용하지않은 LU 분해프로그램과비교하여 7~8배이상의성능향상을보였다. 일반적으로캐시메모리는응용프로그램에서직접사용하기가어렵기때문에성능의향상을위한프로그램의최적화는제한적이다. 하지만 Nvidia GP-GPU 에서는응용프로그램이직접캐시메모리 ( 공유메모리 ) 를사용하는것이가능하며이는성능향상을위해서는절대적으로필요하다. 따라서본연구에서제안한프로그램방식은뛰어난성능의향상을보였으며, 이방식은향후유사한행렬의수치프로그램에서활용할수있을것으로기대한다. 참고문헌 (Reference) [1] Nvidia, CUDA Programming Guide 4.2. [2] J. Nickolls, "Scalable Parallel Programming with CUDA", ACM Queue, vol. 6, no. 2, pp.4-53 8. [3] E. Lindholm, "NVIDIA Tesla: A Unified Graphics and Computing Architecture", IEEE Micro, vol. 28, no. 2, pp.39-55 8. [4] Nico et al., "LU-GPU: Efficient Algorithms for Solving Dense Linear Systems on Graphics Hardware", SC '5 Proceedings of the 5 ACM/IEEE conference on Supercomputing, pp. 3. [5] John L. Hennessy, David A. Patterson. "Computer Architecture: A Quantitative Approach". 211. [6] Golub, Gene H. Van Loan, Charles F. (1996), Matrix Computations (3rd ed.), Baltimore: Johns Hopkins. [7] Shin, B., Y. Kim, Implementation of high performance parallel LU factorization program for multi-threads on GPGPUs, Journal of Korean Society for Internet Information, Vol. 12, No. 3, pp. 131-137, 211. [8] Kim, Y., Performance Comparison of Two Parallel LU Decomposition Algorithms on MasPar Machines, Journal of IEEE Korea Council, Vol. 2, No. 2, pp. 247-255, 1999. [9] G. Fox, M. Johnson, G. Lyzenga, S. Otto, J. Salmon, and D. Walker, 'Solving Problems on concurrent Processors Vol. 1.', Prentice Hall, Englewood Cliffs, NJ, 1988. [1] Gallivan et al., "Parallel Algorithms for Matrix Computations", SIAM, Philadelphia, 1991. 46 213. 12
저자소개 김영태 (Youngtae Kim) 1986 년연세대학교수학과 ( 이학사 ) 1991 년 Iowa State Unniv. 대학원 Department of Computer Science( 석사 ) 1996 년 Iowa State Unniv. 대학원 Department of Computer Science(Ph.D.) 1998 년 ~ 현재강릉원주대학교컴퓨터공학과교수관심분야 : 병렬처리 E-mail : ykim@gwnu.ac.kr 김두한 (Doo-han Kim) 213 년강릉원주대학교컴퓨터공학과 ( 공학사 ) 213 년 ~ 현재강릉원주대학교컴퓨터공학과대학원생관심분야 : 병렬처리 E-mail : mkonjo@naver.com 유명한 (Myoung-han Yu) 9 년강릉대학교컴퓨터공학과 ( 공학사 ) 211 년강원대학교전자공학과 ( 공학석사 ) 1998 년 ~ 현재강릉원주대학교컴퓨터공학과대학원생관심분야 : 무선통신, 컴퓨터네트워크 E-mail : greatymh@gmail.com 한국인터넷정보학회 (14 권 6 호 ) 47