-- ChangWonLee - 2011-01-11

CUDA Programming

CUDA

• CUDA - Computer Unified Device Architecture
  • 일반 고급 언어를 이용하여 CPU에서 작업하는 일반적인 연산을 GPU를 사용해 수행함
  • GPU의 많은 수의 코어를 이용하여 컴퓨팅 성능을 개선시켜주는 NVIDIA의 병렬 컴퓨팅 아키텍처
  • 고급언어를 이용하여 프로그래밍 가능하다는 점에서 GPGPU Programming과 다름

• 지원 제품
  • 지원 제품 리스트

• 설치
  1. OS에 맞는 Driver install
  2. Toolkit install
  3. GPU Computing SDK code samples install
  • [참고] - nvidia cuda 설치에서 getting started guide 참고

• kernel
  • cuda 디바이스를 여러 thread를 통해 동시에 실행하기 위한, 호스트에서 호출하는 함수
    즉, cuda 함수
  • kernel_name<<<gridDim, blockDim>>>(param...)
  • kernel은 총 thread의 개수만큼 호출됨

OpenMP Programming

OpenMP

• schedule - schedule([static, dynamic], size) => size만큼 잘라서 thread에 배분함. nested parallel에선 맨 뒷부분의 schedule만이 적용됨. 그 외에 private은 모두 적용됨.
• schedule을 사용하지 않고 thread를 나누지 않은 경우에는 1부터 시작하여 설정된 thread개수와 나눈다. 그 후 몫이 thread의 개수보다 작거나 같을 때 그 몫을 thread의 개수로 함. (단, 소수점은 올림!)

Experiments

Matrix Multiplication

Algorithm

degree <- n;
matrix1, matrix2, result <- n by n matrix;
for i=0 to degree-1{
    for j=0 to degree-1{
        sum <- 0;
        for k=0 to degree-1 {
            sub <- sub + matrix1[i*degree + k] * matrix2[k*degree + j];
        }
        result[i*degree+j] = sum;
    }
}

Performance

degree	openMP	CUDA	CUDA op time	Serial
256	0.026	1.434	0.001749	0.110
512	0.172	1.509	0.012932	0.937
1024	1.37	1.509	0.103234	11.715
2048	21.316	2.367	0.886828	187.312
4096	312	9.071	7.431047	1533.741
8192	1905.377	68.106s	65.742883	12884.293

• 그래프를 보면 cuda operation부분만 Exponential하게 증가하고 나머지는 polynomial하다는 것을 알 수 있다.
• Serial한 Matrix Product와 openMP, CUDA의 수행시간을 비교해 보면 openMP는 최대 8.7배, CUDA는 189배 빠른 속도로 수행하였다.
• 이 외에 branch나 dependency한 프로그램에 해대서는 그 내부에 independent한 부분을 병렬화 시키거나 프로그램 자체를 수정하여 독립적으로 만들어야 parallel programming한 효율이 높게 나타난다.
• independent한 부분의 loop가 그리 많지 않다면 parallel하게 할 경우 thread를 생성하여 값을 복사하거나 GPU memory로 copy하는 과정에 의해 더 많은 시간이 소요될 수 있다.
• Matrix Product에서 본다면, openMP에서는 적어도 256 by 256인 square matrix. 즉, 256*256*256(2^24)정도는 되어야 4배 정도의 효율을 보였다.
• CUDA의 경우에는 1024 by 1024 square matrix부터 7.76배의 빠른 수행시간을 보였다.
• 또한 CUDA는 Memory size에 제약이 많기 때문에 매우 큰 size의 데이터를 어떻게 independent하게 나눠야 하는가에 대한 부분이 중요하다.
• openMP와 CUDA를 통해서 높은 효율을 기대하기 위해서는 기존의 프로그램을 parallel program으로 재설계하는 설계방식에 대한 공부가 필수적일듯 하다.
• 앞으로의 목표는 openCL을 이용하여 CPU, GPU를 사용하는 Heterogeneous computing환경을 효율적으로 구성하는 방향으로 알아볼 것이다. 이번 UWURF기간에 코코링크사에서 오셔서 진행했던 세미나에서 OpenCL과 CUDA C를 사용한 프로그램의 성능이 5%밖에 차이나지 않는다고 하니, 5%를 감수하고 확장성 높은 방향으로 진행하는것이 유리할듯 하다. 게다가 CPU와 GPU에 대해 모두 수행 가능하기 때문에 긴 loop부분은 GPU로 처리하고 그 외에 짧은 loop부분은 CPU를 이용하여 더 높은 성능향상을 기대할 수 있을듯 하다.
• 사실 그 외에도 intel의 CPU와 GPU를 통합한 프로세서인 Larrabee가 무산되기는 했지만, 결국에는 이기종의 computing unit에 공유 메모리가 탑재되는 방향으로 개발이 이뤄질것 같기 때문에 Nvidia의 생각보다 제한이 많은 CUDA C는 무조건적으로 수행속도를 빠르게 내야하는 경우를 제외하고는 그리 많이 사용될것 같지 않다고 생각한다.

performance analysis

• cores, threads

	openMP	CUDA
cores	64	240
threads	128	960

• NVIDIA GeForce GTX 275는 CUDA core가 240개이고 Intel Xeon Processor E5530의 경우 4개의 core를 갖는다. 여기서 /proc/cpuinfo를 보면 총 16개의 processor가 존재하므로 물리적인 core의 개수는 총 64개이다.
• CUDA는 30개의 SM(Streaming Multiprocessor)에 총 8개의 SP(Streaming Processor)가 존재한다. 여기서 하나의 실행 단위는 warp라 하는데 32개의 thread가 모여서 동시에 실행된다. 하나의 warp는 8개의 SP에 4개의 thread씩 할당되므로, SM당 32개의 thread가 존재한다. 따라서 960개의 thread가 존재한다.

GPU와 CPU는 동일한 clock, 한 cycle당 동일한 명령어를 처리한다고 가정한다. GPU는 960개의 thread로 동일한 작업을 나누어 수행하고, openMP를 사용한 프로그램에서는 128개로 나눠 수행한다. 하지만, CPU는 일반적으로 명세서에 나온 성능보다 1/10의 효율을 낸다. 그리고 simple scalar에서 cache의 크기를 matrix의 크기와 같게한 후 simulation한 결과 cpu utilization이 0.68이 나왔으므로(simple scalr의 Experiments - Matrix Multiplication 참조) 128개의 thread에 1/10의 성능인 0.1과 100%성능을 위한 1.5를 곱해주면 19.2가 된다. 즉, GPU를 이용하는 쪽이 960/19.2 = 50배 빠를것이라 추측할 수 있다.

degree	openMP/CUDA
256	8.66
512	34.4
1024	105.38
2048	507.52
4096	1950

위의 표를 보면 degree가 2배(matrix의 size가 4배)씩 늘어날 때마다 GPU와 CPU사이의 속도차이는 점점 심해짐을 알 수 있다.