CUDA programming_01

FROM http://newsight.tistory.com/135

인터넷 강의 : http://www.youtube.com/watch?v=m0nhePeHwFs&list=PLKK11Ligqititws0ZOoGk3SW-TZCar4dK

Kernel : CPU가 GPU에게 큰 load의 작업을 병렬수행하도록 시키는 코드
이것을 만드는 것이 바로 쿠다 프로그래밍이고 gpgpu이다.

함수 종류 지정(함수의 이름 앞에 다음을 선언하여 실행 방식을 결정한다.)

__global__ : cpu(host)가 요청하여, gpu(device)에서 실행하는 함수
(쿠다 프로그래밍에서 가장 많이 쓰이는 핵심 함수이다.)

__device__ : gpu가 요청하여, gpu에서 실행하는 함수

__host__ : cpu가 요청하여, cpu가 실행. 일반적인 C++함수로 생략 가능하다.

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;
}

kernel 함수 실행시 configuration 설정 : 함수이름<<<쓰레드 블럭 수, 각 블럭당 쓰레드 수>>>(파라미터 전달);
ex) AddIntsCUDA<<<100,10>>>(a,b); 실행시 해당 글로벌 함수를 10개의 쓰레드를 갖는 블럭 100개를 만들어 총 1000개의 쓰레드로 실행시킨다.

이때 각 블럭당 쓰레드는 최대 1024개를 넘을 수 없다.(그래픽카드에 따라 512개 일 수도 있다)
그러나 쓰레드 블럭은 최대 2^32-1 개까지 만들 수 있다.(그래픽카드에 따라 2^16-1개 일 수도 있다)

일반적인 사양의 GPU에서는 총 쓰레드가 2000~3000개 정도를 갖도록 설정하는게 일반적이다.
쓰레드 블럭과 각 블럭당 쓰레드의 수는 GPU의 스펙에 맞게 그때 그때 설정하는 데에 목적이 있다.

쓰레드 블럭 : 쓰레드끼리 서로 내부 통신하는 단위(shared memory를 갖는 단위)

각 쓰레드의 고유 id 는 (쓰레드 블럭*각 블럭당 쓰레드) + 해당 쓰레드의 index 로 구할 수 있다.
ex) int id = blockIdx.x * blockDim.x + threadIdx.x; (각각은 이미 쿠다에 정의된 변수임)

Grid : 쓰레드 블럭의 묶음

Dim3 : x,y,z 성분을 갖는 3차원 벡터

SM(Streaming Multiprocessor, 줄여서 Multiprocessor) : 보통 16개의 쓰레드 블럭을 동시에 실행시키는 단위

Warps : 워프는 SM에 할당된 스레드들을 32개 씩 한 사이클에 동시에 수행하도록 스케줄링 한다. (SIMD 같은 방식으로)

- 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];

- Shared Memory 사용시 주의사항 : 동기화

쓰레드 블럭내의 모든 쓰레드들이 접근하므로 동기화에 주의해야한다. 따라서 필요한 경우 mutex나 semaphore 같은 동기화 기법이 필요하다. 그러나 동기화는 시퀀셜  메모리 엑세스를 발생시키므로 성능이 저하될 위험이 있으므로 최대한 피해야한다.

대신에 쓰레드간의 수행 단계를  동기화하는 방법을 보통 사용한다.
__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;)

예제 소스 : 

 main.cu

위의 예제 소스에서 __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 코어를 갖는다.