CUDA 프로그래밍에서의 위험요소
공식문서
기본 튜토리얼
http://www.youtube.com/watch?v=m0nhePeHwFs&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK
NVCC 컴파일 및 실행
용어 및 개념
Dim3 : x,y,z 성분을 갖는 3차원 벡터
SM(Streaming Multiprocessor, 줄여서 Multiprocessor) : 보통 16개의 쓰레드 블럭을 동시에 실행시키는 단위
Warps : 워프는 SM에 할당된 스레드들을 32개 씩 한 사이클에 동시에 수행하도록 스케줄링 한다. (SIMD 같은 방식으로)
여러 가지 Kernel definition
__device__ : gpu가 요청하여, gpu에서 실행되는 함수
__host__ : cpu가 요청하여, cpu에서 실행되는 일반적인 C++함수로 __host__ 선언은 생략 가능하하다.
Kernel execution configuration syntax
- 함수이름<<<쓰레드 블럭 수, 각 블럭당 쓰레드 수>>>(함수 인자 전달)
ex) AddIntsCUDA<<<100,10>>>(a,b); 실행시 해당 글로벌 함수를 각10개의 쓰레드를 갖는 블럭 100개를 만들어 총 1000개의 쓰레드로 실행시킨다.
그러나 이는 엄밀하게 말하면 1차원 thread block과 1차원 grid에 대한 예시이다. 왜냐하면 여기서 M과 T가 int인 경우 1차원을 생성하지만, dim3 object로 지정할 경우 2차원, 3차원의 thread block과 grid 생성이 가능하다.
Heterogeneous Programming
Thread Hierarchy
CUDA Built-in variable
CUDA에서는 프로그래머가 아주 편리하게 thread를 indexing할 수 있도록, 3차원(threadIdx.x / threadIdx.y / threadIdx.z)의 indexing을 제공한다. 만약 thread block을 1차원으로 정의하면 x만 쓰면되고, 2차원이면 x,y 3차원이면 x,y,z를 쓰면된다. 각 thread의 ID는 다음과 같이 계산된다
그래서 1차원 block의 경우 각 쓰레드의 고유 id 는 (쓰레드 블럭*각 블럭당 쓰레드) + 해당 쓰레드의 index 로 구할 수 있다.
ex) int threadID = blockIdx.x * blockDim.x + threadIdx.x;
Thread와 Thread block
그러나 이러한 제약이 없는 쓰레드 블럭은 최대 2^32-1개(거의 무한)까지 만들 수 있고, 역시 그래픽카드 종류에 따라 2^16-1개 일 수도 있다.
보통의 일반적인 GPU에서는 총 쓰레드 개수(블럭X쓰레드수) 가 2000~3000개 정도를 갖도록 설정하는게 일반적이다.
쓰레드 블럭과 각 블럭당 쓰레드의 수는 GPU의 스펙에 맞게 설정하면 더 최적화된 성능을 낼 수도 있다.
이때 Thread block들이 실행하는 각 task들은 서로 어떠한 dependency도 없이 호출되어야만한다. 즉 어떤 순서로 실행되든지 상관이 없어야한다는 의미이다. 그리고 한 block내에서의 thread들은 적절한 synchronization을 한다면 shared memory를 통해서 서로 데이터를 공유할 수 있다. 이 때쓰이는 함수로 __syncthreads()가 있고, 이것은 일종의 barrier(장벽)으로 block내의 모든 thread들이 반드시 기다려야하는 지점이다.
Grid
이 Thread block들은 다시, grid란 단위로 묶이는데 이 역시 1차원, 2차원, 3차원 grid를 형성 할 수 있다. 대게의 경우 한 grid 안에 들어 있는 thread block의 개수는 데이터의 크기인 경우가 많고, 보통의 경우 processor의 개수를 초과한다. 그래서 최종적으로, 2차원 grid와 2차원 block 쓰는 예시를 그림으로 표현하면 다음과 같다.
(출처: Nvidia공식 cuda documentation)
Memory Hierarchy
CUDA 메모리 함수
GPU는 시스템메모리 즉, 컴퓨터에 장착된 DRAM을 엑세스할 수 없다.
따라서 GPU는 GPU에 장착된 메모리(global memory)에 데이터를 복사해와서 읽어야하며, 수행한 작업도 local 메모리에 쓴다음 다시 시스템 메모리로 옮겨야 한다.
이 작업은 굉장한 speed bottle neck이다. 따라서 최적화가 필요하다.
cudaMalloc(포인터, sizeof(타입)*개수) : gpu의 글로벌 메모리에 메모리를 동적할당한다.
cudaFree(포인터) : 위에서 할당한 메모리를 해제한다.
cudaMemcpy(저장할 포인터, 값, sizeof(타입)*개수, 방향) : gpu의 메모리와 cpu의 메모리간의 복사 즉, 메모리 통신을 수행한다.
cudaMemset() :
cudaDeviceReset() : gpu의 모든 메모리와 작업중인 쓰레드를 초기화시킨다.
cudaSuccess : 각 쿠다 api가 성공할 경우 리턴하는 상수
ex) if(cudaMalloc(&d_b, sizeof(int)*count) != cudaSuccess){
cout<<"Error B!"<<endl;
}
Device Memory L2 caching
CUDA C++ api를 import하는 header의 차이점
stackoverflow.com/questions/6302695/difference-between-cuda-h-cuda-runtime-h-cuda-runtime-api-h
- kernel에서 shared memory 할당하기 :
1. 스태틱 셰어드 메모리 사용 -> 커널 함수에서 다음과 같은 형식으로 변수를 셰어드 메모리에 할당한다.
ex) __shared__ int i=1;
2. 다이나믹 셰어드 메모리 사용 -> 커널 호출시 세 번째 파라미터로 블럭당 셰어드 메모리 크기를 지정한다.
ex) AddIntsCUDA<<<100,10,128>>>(a,b);
그다음 커널 함수에서 다음과 같이 선언
extern __shared__ int parray[]; // 이러면 위에서 선언한 128만큼이 자동으로 할당된다.
char * chararray = (char *)&parray[0];
float * floatarray = (float *)&parray[10];
__syncthreads(); 를 호출하면 모든 쓰레드 블럭내의 쓰레드들이 그 지점에 도달 할 때까지 기다리며 동기화를 하게된다.
- CUDA에서 제공하는 기본 Structure
float3 : x,y,z 로 구성된 3차원 공간의 점을 나타내는 구조체
- Memory Architecture of GPU
Global Memory : 그래픽카드에 장착된 램. 읽기 속도가 매우 느림. 대신 자유롭게 어떤 쓰레드든지 어떤 주소라도 엑세스가 가능하고 용량이 큼.
Local Memory : 각 쓰레드 마다 개별로 할당된 DRAM이다. 레지스터에 할당된 메모리가 가득차면 그다음에 Local memory를 사용하게 되어 속도가 매우 느리다.(레지스터가 꽉찼을 때만 사용한다) 각 로컬메모리를 소유한 쓰레드만이 메모리를 엑세스할 수 있다.
Shared Memory (= L1 Cache) : 각 스트리밍 멀티프로세서가 소유하고 있는 매우 빠른 메모리로 레지스터와 같은 속도를 갖는다. 같은 블럭 안의 쓰레드들은 모두 이 셰어드 메모리를 통해 쓰레드간의 통신을 한다.(같은 쓰레드 블럭에 있는 쓰레드끼리만 가능)
그러나 bank conflict가 일어나면 속도가 매우 느려진다.(뱅크 컨플릭트는 아래 글을 참조)
http://newsight.tistory.com/134
capability 2.0부터는 32개의 bank가 각 셰어드 메모리에 존재하며, 순서대로 4byte(dwords)씩 다른 bank에 들어가 있다. 따라서 이 경우 각 쓰레드가 32의 배수 개로 존재해서, 32*4byte 의 단위로 index를 접근하는게 최고 속도로 셰어드 메모리를 읽는 방법이 되겠다.
Register : 각 쓰레드는 아주 작은 레지스터를 갖고있으며 이는 매우 빠른 속도의 메모리이다. kernel에서 선언한 변수는 제일 먼저 레지스터에 할당되어 쓰이며, 레지스터 용량이 가득차면 로컬메모리에 변수를 할당한다.
또한 GPU에는 수천개의 레지스터가 있기 때문에, 충분한 수의 쓰레드가 있어야만 최고 성능을 낼 수 있다.
Caches : 캐시는 L1 cache(Level 1)과 L2 cache,L3 cache 로 구분되는데, L1이 가장 비싸고 가장 빠르며 가장 적은 용량을 갖고, 레벨이 올라갈수록 느려지고 용량은 커진다. 그러나 L1 캐시의 용량을 무한정 늘리는 것은 불가능한데, 그 이유는 L1캐시가 커질 수록 cache hit 는 올라가겠지만, 데이터를 찾는 속도가 느려져서 오히려 성능이 저하된다. (책장이 클수록 책을 찾는게 느려지듯. 또한 노트북같은 모바일 기기는 L3캐시의 크기가 매우 큰데, 그 이유는 램을 엑세스하는 것은 캐시에비해 전력소모가 매우 크기 때문이다.)
실제로 GPU에서는 어떤 변수를 엑세스하면 먼저 L1 캐시에서 찾아보고 없으면 L2 캐시, 없으면 글로벌 메모리에서 찾는다.(L3 캐시가 없는 듯)
L1 캐시는 shared memory와 물리적으로 동일한 메모리로 각 스트리밍 멀티프로세서에 위치하며 속도가 매우 빠르다. 그래서 둘사이의 용량 비율을 선택할 수도 있다.(쓰레드 블럭당 셰어드 메모리와 L1 캐시의 비율)
ex) 커널 실행 전, C 코드에서 다음과 같이 선언할 것
cudaFuncSetCacheConfig(커널 함수 이름, CudaFuncCachePreferNone); -> default 값
cudaFuncSetCacheConfig(커널 함수 이름, CudaFuncCachePreferShared); -> L1 16 Kbytes : Shared 48 Kbytes
cudaFuncSetCacheConfig(커널 함수 이름, CudaFuncCachePreferL1); -> L1 48 kbytes : Shared 16 Kbytes
cudaFuncSetCacheConfig(커널 함수 이름, CudaFuncCachePreferEqual); -> L1 32 kbytes : Shared 32 Kbytes
L2 캐시는 글로벌 메모리처럼 모든 쓰레드에서 접근이 가능하다.
그런데 사실 L1 및 L2캐시는 하드웨어에서 알아서 자주 사용하는 변수들이 캐싱하는 용도이므로 프로그래머가 활용할 수 없다. 대신 셰어드 메모리는 프로그래머가 직접 할당해서 잘 사용해야 한다.(셰어드 메모리로 L1 캐시의 일부를 프로그래밍 하는 것이라고 볼 수도 있다.)
Constant Memory : 원래 GPU에서 그래픽 계산을 위한 파이값이나 뭐 그런걸 저장해놓고 빠르게 읽는 용도인 듯. CPU가 커널을 실행하기전에 변수들을 저장해 놓고, 커널에서는 오직 읽기만이 가능하다. 이 또한 레지스터와 셰어드 메모리 수준의 속도를 갖고 있으며 모든 쓰레드에서 글로벌메모리 처럼 접근 가능하다. 대신 64k의 아주 작은 용량을 갖는다.
(아주 빠른 읽기 전용 글로벌 메모리라고 생각)
Texture Memory : constant memory와 마찬가지로 cpu가 값을 세팅해 놓으면 gpu에서는 읽을 수만 있다. 역시 글로벌 메모리처럼 모든 쓰레드에서 접근이 가능하다.
- 개발환경 설정(인터넷강의 튜토리얼 2번, 8분~15분 참조)
1. Visual Studio 2010 설치(dreamspark 에서 다운로드)
https://www.dreamspark.com/Student/Software-Catalog.aspx
2. CUDA 최신버전 설치
https://developer.nvidia.com/cuda-downloads
3. Visual Studio 에서 C++ empty 프로젝트 생성
4. 프로젝트에서 우클릭, Build Customizing 에서 CUDA 체크
5. 프로젝트 속성에서 VC++ 디렉토리 - include directory에 다음 경로 추가
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\include (본인의 쿠다 설치 경로 참조)
6. 같은 창에서, Library directory 에 다음 경로 추가
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\lib\x64 (본인의 쿠다 설치 경로 참조)
7. 프로젝트 속성에서 Linker - input 의 추가 종속성에 다음 추가
$(CudaToolkitLibdir)\cudart.lib;
(다음과 동일 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\lib\x64\cudart.lib;)
예제 소스 :
위의 예제 소스에서 __global__ 함수를 다음과 같이 바꾸어보자.
__global__ void AddIntsCUDA(int *a, int *b){
for(int i=0; i<10000000000000000;i++){
a[0] += b[0];
}
}
그러면 일시적으로 컴퓨터의 화면이 꺼지면서 프로그램은 GPU를 쓰지않은 결과를 낸다. 그리고 1~2분 정도 지나면 다시 원래대로 돌아온다. 이것은 일반적인 윈도우에서의 GPU작업은 매우 빠른 시간동안에 끝나기 때문에, GPU에서 2초이상의 작업을 수행할 경우 문제가 있다고 판단, GPU를 reset시켜버리기 때문이다.
따라서 이를 막기위해서는 윈도우 레지스트리를 수정해야한다.(다음 사이트 참조)
http://msdn.microsoft.com/en-us/library/windows/hardware/ff569918(v=vs.85).aspx
TdrDelay 를 새로 추가하고, 값을 30정도로(30초를 의미) 지정해주면 적당하다.(레지스트리 수정 후 재부팅)
- CUDA SDK Sample 소스코드 컴파일 하기
1. 아래 주소에서 예제 소스를 다운 받는다.
http://docs.nvidia.com/cuda/cuda-samples/#samples-reference
2. VC++ 프로젝트에 코드를 복사하고, 프로젝트 속성에서 VC++ 디렉토리 - include directory에 다음 경로도 추가한다.
C:\Program Files\NVIDIA GPU Computing Toolkit\NVIDIA_CUDA-6.5_Samples\common\inc
3. 컴파일 한다.
- 내 컴퓨터에 장착된 그래픽카드 성능 분석
1. 예제 소스에서 deviceQuery를 컴파일한다.
내 컴퓨터에 장착된 GTX 650의 사양
CUDA 코어 : 384 개
글로벌 메모리 : 1024 Mbytes
constant memory : 65536 bytes
shared memory per block : 49152 bytes
register per block : 65536
warp size : 32
- 성능 최적화
1. 적절한 쓰레드 블럭과 쓰레드 수 정하기
가장 간단한 방법은 커널을 실행할 때 블럭수와 쓰레드수를 가변적으로 for문을 돌려서 가장 수행시간이 적게 나오는 숫자를 찾는다.
2. SM과 compute 숫자 변경
프로젝트 속성 - configuration - CUDA C/C++ - Device 의 Code Generation 에서 숫자를 변경한다.
3. 32 bit / 64 bit 머신 변경
프로젝트 속성 - configuration - CUDA C/C++ - Common 의 Target Machine Platform 에서 변경한다.
경우에 따라 64bit에서 더 빠를 수도, 더 느릴 수도 있다.
- NVVP를 이용해 개발한 CUDA 프로그램 프로파일링하기
1. Visual Profiler 를 실행시킨다.(시작메뉴 또는 아래 주소에 있음)
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.5\libnvvp
2. File - New Session을 클릭하고, 분석할 쿠다 프로그램(.exe)를 지정한다.
(Visual Studio 프로젝트 폴더에서 debug에 보면 있음)
3, 어느 부분에서 시간이 가장 많이 소모되는지 확인한다.
- 코드에서 수행시간 측정하기
1. cpu에서 다음 코드 이용
long startTime = clock();
GPU 커널 실행
long finishTime = colck();
printf("%d\n", finishTime-startTime);
2. NVVP를 이용한 프로파일링
3.
- 질문
질문 1. CUDA는 MPI인가 open MP인가?
CUDA는 incremental parallelization 이 아닌, 처음부터 병렬화를 고려하는 방식이기 때문에 방법론적으로는 MPI가 맞다. 그러나 셰어드 메모리를 통해 통신한다는 점은 open MP에 해당한다. 그러므로 관점을 어디에 두느냐에 따라 달라진다.
질문 2. kernel에는 for문이 필요없는가? 필요할 수도 있고, 필요 없을 수도 있다.
기본적으로 모든 쓰레드가 각자 똑같은 커널 함수를 실행시킨다. for문을 넣어서 각 쓰레드가 긴 작업을 할 수도 있고, for문 없이 각 쓰레드가 1회만 실행하게 할 수도 있다.
질문 3. kernel에서 printf() 실행이 가능한가?
쿠다 아키텍쳐 2.0이상의 그래픽카드에서는 실행이 가능하다. 호환성을 위해서는 아래와 같이 사용할 수 있다.
#if __CUDA_ARCH__ >= 200 printf("Hi Cuda World"); #endif질문 4. Shared memory를 통한 쓰레드간의 통신은 블럭간에서도 가능한가? 아니면 같은 블럭안에서만 사용 가능한가?
같은 쓰레드 블럭안에서만 공유가 가능하다. 즉 SM은 각 쓰레드 블럭에게 Shared Memory를 조금씩 할당해주는 방식으로 보인다. 그럼 여기에도 Bank 개념이 포함되겠군..아마 ?
질문 5. SM은 1개의 코어를 갖는가, N개의 코어를 갖는가?
SM은 보통 16개의 FPU 코어를 갖는다.
'Development > for Machine Learning' 카테고리의 다른 글
| Practical Theano Tutorial (0) | 2015.03.11 |
|---|---|
| Theano 사용법 (0) | 2015.01.20 |
| CUDA / CuDNN / Theano / TensorFlow 설치하기 (4) | 2014.07.07 |
| Weka (0) | 2014.06.23 |
| MATLAB 문법, API, 단축키, 환경설정 (0) | 2014.06.23 |
댓글