3D Client Programmer

이 함수를 호출 함으로써 DirectX와 CUDA 양쪽에서 ID3D11Resource를 사용할 것이라고 CUDA 런타임에게 명시할 수 있다. CUDA 런타임은 변수 resource를 통해 버퍼를 가리키는 CUDA 전용 핸들을 하나 반환한다. 이 핸들은 차후에 CUDA 런타임의 API를 호출시 ID3D11Resource를 참조하기 위해 사용될 것이다.

이 함수는 아래 사이트에서 확인하자.

https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__D3D11_gf0545f2dd459ba49cfd6bcf2741c5ebd.html

cudaMallocPitch()는 CUDA에서 2차원 배열을 생성하는 함수이다. C처럼 2차원 배열을 만들 수 없기 때문에 이렇게 따로 함수가 있다. 이것에 대해서는 다음에 다시 글을 작성하도록 하겠다.

cudaMemset()은 말 그대로 값 세팅을 한다.

cudaGraphicsMapResources()는 앞에서 ppResource를 CUDA에 주소를 할당하고, cudaGraphicsUnmapResources()은 ppResource의 주소 할당을 해제한다.

앞의 함수들을 호출했다면, 이제 DirectX를 초기화 해야한다.

앞에서 g_pCudaCapableAdapter를 구했었다. 이 변수를 이용해서 D3D11의 Device를 생성해야한다.

이때 주의해야 할 점은 D3D11CreateDevice()를 사용할 때 첫 번째 인자로 Adapter를 지정했다면, 두 번째 인자는 D3D_DRIVER_TYPE_UNKNOWN으로 반드시 지정해야 한다.

그리고 CUDA Adapter를 이용하여 Device를 생성했다면, 이제 CUDA에도 D3D Device를 지정해줘야 한다.

이 함수는 ::IDXGIFactory::EnumAdapters에서 찾은 Adapter에 해당하는 CUDA 호환 장치를 *device에 넣어 반환한다.

이 함수는 Adapter가 CUDA를 지원해야만 성공적으로 cudaSuccess 리턴을 한다.

다음 함수는 Cuda를 지원하는 Device를 찾아 IDXGIAdapter의 변수 g_pCudaCapableAdapter에 저장하는 함수이다.

이 g_pCudaCapableAdapter는 이후 D3D11CreateDeviceAndSwapChain()를 할 때 사용된다,

CUDA와 DirectX의 상호운용 방법에 대해서 살펴보기 전에, 몇가지 구조체들과 함수에 대해 알아보자.

이 함수는 디바이스 장치의 개수를 획득하는 함수이다.

parameter로 count의 포인터를 넘겨 받기 때문에 당연히 count에는 Device의 개수가 리턴되어서 나오게 된다.

이 구조체는 CUDA Device의 정보들을 저장할 수 있는 구조체이다. 

이 구조체의 변수는 매우 많기 때문에 다음 링크를 참조하도록 하자.

http://horacio9573.no-ip.org/cuda/structcudaDeviceProp.html

이 함수는 CUDA Device에서 정보들을 가져오는 함수이다. 이 함수를 이용하여 각 Device들의 정보를 가져와서 알맞게 사용할 수 있다.

위의 함수들을 이용하여 Device를 읽어오는 코드를 보자.

Device를 찾을 때 또 다른 유용한 함수도 있다.

cudaDeviceProp 구조체의 값을 채운 후 이 함수를 호출하면 이 정보에 맞는 Device를 찾아서 Device 번호를 첫 번째 매개변수에 넣어준다.

CUDA에서는 많은 스레드를 사용하게 된다. 앞에서 메모리의 계층구조를 보면, 공유 메모리라는 것도 있다. 

많은 스레드들이 공유 메모리에서 서로 데이터를 '쓰기, 읽기' 를 하기 위해서는 당연히 동기화를 해야만 한다.

이 동기화를 위해 C언어에서는 굉장히 머리 아픈 상황이 많이 발생하지만, CUDA C에서는 다행히 동기화 함수를 제공하고 있다. 

__syncthreads() 이 함수는 블록 내의 다른 모든 스레드가 __syncthreads()를 호출해야만 다음 명령어로 넘어가게 되어 있다.

그림 - __syncthreads() 사용 예

스레드 동기화에서 주의 할 점은 일단 동기화 자체가 작업이 먼저 끝난 스레드가 아직 작업이 끝나지 않은 스레드를 기다리는 것이기 때문에 성능에 영향이 갈 것이다. 꼭 사용해야 되는 부분에서만 사용하도록 하자.

그림 - CUDA 메모리 계층 구조

그림 - CUDA 메모리 특성

액세스 속도 (느림 -> 빠름)

Global -> Local -> Shared -> Register

메모리 용량 (작음 -> 큼)

Register -> Shared -> Local -> Global

액세스 범위 (좁음 -> 넓음)

Register -> Local -> Shared -> Global

Register Memory는 온 칩 프로세서에 있는 메모리로 직접 연산을 수행하는 가장 빠른 메모리이다. CUDA C를 이용해 구현할 때도 몇 개의 레지스터를 사용하는 지와 레지스터가 어떻게 재사용 되는지 알 수 없다.

다만 커널 함수 내에서 로컬 변수를 많이 사용할수록 레지스터 사용량이 증가하고, GPU마다 레지스터의 수는 차이가 있다.

sin(), cos() 등 수학 함수를 사용할 때도 레지스터를 사용한다.

로컬 메모리(Local Memory)는 커널 함수 내에서 너무 많은 로컬 변수를 사용하거나, 배열형 변수로 큰 용량을 사용하면 프로세서 밖에 있는 DRAM에 메모리가 할당된다. 로컬 변수가 레지스터로 사용될지 로컬 메모리에 할당 될지 명확하지 않다.

Local Memory로 할당되는 경우는 다음과 같다.

1. 너무 많은 레지스터 변수를 사용했을 때

2. 너무 많은 로컬 변수를 사용했을 때

3. 로컬 변수로 배열을 사용했을 때

쉽게 말하면, 각 Thread는 자신만이 사용하기 위한 메모리이다.

공유 메모리(Shared Memory)는 CUDA의 큰 장점 중 하나이다. CUDA C 컴파일러는 공유 메모리의 변수들을 일반 변수들과는 다르게 취급한다. CUDA C 컴파일러는 GPU에서 실행되는 각 블록마다 변수의 복사본을 하나씩 생성한다. 블록 내의 각 스레드는 메모리를 서로 공유하지만, 다른 블록 내에서 보고 있는 변수의 복사본을 보거나 수정할 수는 없다. 하지만 공유 메모리는 블록 내의 스레드들이 계산하는 데 통신과 협력을 할 수 있게 함으로써 훌륭한 수단을 제공한다. 더욱이 공유메모리 버퍼는 오프칩(off-chip)인 DRAM에 상주하는 것과 달리 물리적으로 GPU 상에 상주한다. 이 때문에 공유 메모리의 접근 지연 시간은 블록당 공유 메모리를 효율적으로 생성하는 일반 버퍼, 소프트웨어로 관리되는 캐시 또는 스크래치패드(Scratchpad)보다 훨씬 짧은 경향이 있다.

GPU의 SM 안에 공유 메모리가 있고, SM은 블록 단위로 프로그램을 처리 하기 때문에 공유 메모리의 데이터는 동일한 블록 안에 있는 스레드 사이에서만 공유할 수 있다. 공유 메모리를 할당하는 방법은 정작할당, 동적할당이 있다.

쉽게 말하면, 각 Thread 간에 data 공유를 위한 메모리이다.

__shared__ int a; 

전역 메모리(Global Memory)는 비디오 카드에 장착된 DRAM 메모리를 의미한다. 전역 메모리는 GPU 칩 외부에 있기 때문에 액세스 속도는 레지스터나 공유 메모리에 비해 많이 느리지만, CPU의 메모리 보다는 빠르다.

쉽게 말하면, Block들의 집합인 Grid간에 data 공유를 위한 메모리이다.

상수 메모리(Constant Memory)는 DRAM에 있는 데이터를 읽기 전용으로 사용하며 캐시를 지원한다. NVIDIA의 그래픽 카드는 Constant Memory의 크기가 64KB이다. 최초 데이터는 DRAM에서 가져오기 때문에 속도가 좀 느리지만, 한번 캐시에 올라온 값을 반복하여 재 사용하기 때문에 속도가 빠르다. Host에서는 값을 쓰고, Device 영역에서는 값을 읽기만 할 수 있다.

쉽게 말하면, Half - Warp내의 모든 스레드가 동일한 주소의 데이터를 상수 메모리로 부터 요청하면 GPU는 단 한번의 읽기만 요청한 후 모든 스레드에게 그 데이터를 알려주는 메모리이다.

__constant__ int a;

텍스처 메모리(Texture Memory)는 상수 메모리처럼 칩에 캐싱된다. 따라서 어떤 경우엔 텍스처 메모리는 오프칩인 DRAM에 대한 메모리 요청을 줄임으로써 매우 효과적인 대역폭을 제공할 수 있다. 텍스처 캐시는 공간 구역성(spatial locality)를 자주 드러내는 메모리 접근 패턴을 가진 그래픽스 어플리케이션들을 위해 설계되었다. 

texture<자료형> a;

출처 : http://blog.naver.com/riverrun17/220420579990

CUDA 메모리에 대해 공부하다가 어떤 분이 잘 정리해 놓으셔서 긁어왔습니다....

이번에는 앞전에 했던 블록을 이용한 벡터의 합을 스레드를 이용한 벡터의 합으로 구현해보려고 한다. 

사실 코드는 몇줄 바뀌지 않는다. 단순하게 블록에서 하던 일을 스레드에서 하게 바꾸는 작업만 하면 되는 일이니까...

앞에서 작성했던 코드와 비교해서 바뀐 부분은 빨간색 네모를 확인하면 된다.

그림 - add()에서 바뀐 부분

그림 - main()에서 바뀐 부분

앞에서 하드웨어의 제약으로 인해 한 번에 실행 가능한 블록의 수는 65,535개를 초과할 수 없다고 했었다. 동일하게, 하나의 커널을 실행하는 데 있어서 블록당 스레드 수 역시 하드웨어의 제약을 받는다. 구체적으로, 이 수는 우리가 3장에서 보았던 디바이스 속성 구조체의 'maxThreadsPerBlock' 멤버의 값을 초과해서는 안된다. 현재 이용가능한 많은 그래픽스 프로세서들에서 블록당 스레드의 최대 개수는 512개이다.

그렇다면 배열의 개수가 512개보다 더 클 때 어떻게 스레드 기반으로 두 벡터의 합을 수행할 수 있는지 생각해보자. 뭐 방법은 간단하다. 스레드와 블록을 같이 사용하는 방법이다.

그림 - 스레드와 블록을 같이 사용

여기서 'blockDim'이라는 내장 변수를 사용하는데 앞에서 설명했다시피, 각 블록당 스레드의 개수 이다. 여기서는 1차원의 블록을 사용하고 있기 때문에 그냥 'blockDim.x'만 사용했고, 2차원이라면 'y'도 사용 해야 할 것이다.

또 다른 변경은 커널 실행 자체에 있다. 실행할 N개의병렬 스레드들이 여전히 필요하지만, 우리에게 부과된 512개의 스레드 제한 수를 초과하지 않도록 다수의 블록들을 통해 스레드들을 발동시키길 원한다. 하나의 해결책은 임의로 블록 크기를 고정된 스레드의 개수로 지정하는 것이다. 예를 들어, 블록당 128개의 스레드들을 사용한다고 고정시켜보자.

그리고 블록을 생성하는 개수를 정할 차례인데, 여기서 묘안은 N/128이 정수의 나눗셈이라는 것이다. 만약 N이 127이라면 N/128은 정수로 0이 된다. 즉 0개의 블록을 생성하게 되므로 스레드도 실행되지 않게 된다. 하지만, 이렇게 구현해버리면, N이 정확히 128의 배수여야만 올바른 블록의 개수가 실행될 것이다. 이것은 좋지 못하다. 이를 해결하기 위한 일반적인 요령은 다음과 같다.

그림 - 커널의 다른 실행 방법

여기서 N이 10개 이므로, 블록은 10개 스레드는 128개가 된다. 즉 1280개의 스레드가 생성되어 돌아가기 때문에 잘못된 계산이 나올수 있으므로 예외 처리를 해줘야 한다. 하지만, add() 커널 함수에서 이미 예외 처리를 해놨었다.

그림 - 예외 처리가 된 부분

tid가 위에서 계산되었을 때 N보다 클 때에는 if문에서 걸러지기 때문에 올바르게 작동하게 된다.

하지만 이렇게 구현하게 되면, 하드웨어 제약사항으로 인한 블록의 개수(65,535) * 스레드의 개수(512) 개를 초과하게 되면, 병렬 처리가 정상적으로 처리가 되지 않게 된다. 그러므로 더 수정을 해야 한다.

그림 - 더 업그레이드 된 add() 커널 함수

위와 같이 각 스레드가 현재 인덱스에서의 작업을 끝내면 그리드 내에서 작동 중인 스레드의 총 수만큼 각 인덱스를 증가시킬 필요가 있다. 즉, while() 문에서 tid 값을 위와 같이 증가 시키면 각 스레드는 유니크한 인덱스를 구할 수 있을 것이다.

그림 - 고정된 개수의 블록과 스레드를 실행하는 add() 커널 함수

위와 같이 블록의 개수(128) * 스레드의 개수(128) = 16,384개 이지만, 3만개든 10만개든 덧셈이 동시에 가능해진다.

이전에 커널 함수에서 Kernel<<<Block, Thread>>> 이라고 잠깐 설명을 했었다.

여기서 Grid, Block, Thread에 대해 개념을 어느 정도 잡고 있어야 할 것 같아서 정리한다.

우선 한줄 요약을 하면 다음과 같다. "Thread가 모여서 Block이 되고, Block이 모여서 Grid가 된다."

즉, 'Thread -> Warp -> Block -> Grid' 라고 생각하면 된다.

그림 - 동작 단위

커널 함수에서 Kernel<<<N, 1>>> 은 N개의 블록이 각각 1개의 스레드를 가지고 있다는 뜻이다.

Kernel<<<1, N>>> 은 1개의 블록이 N개의 스레드를 가지고 있다는 뜻이다.

앞에서 blockIdx로 블록의 인덱스를 이용해서 데이터를 처리했었다. 

마찬가지로 threadIdx로 스레드의 인데스를 이용해서 데이터를 처리할 수 있다.

blockDim은 블록당 스레드의 수를 나타내고, gridDim은 그리드당 블록의 수를 나타낸다.

위에 있는 Warp에 대해 설명한다. CUDA 아키텍처에서의 워프는 "서로 엮여 있으면서" 정확히 동일한 방식으로 발맞추어 실행되는 32개의 스레드들의 집단과 관계가 있다. 프로그램 코드의 매 라인마다 하나의 워프 내의 각 스레드는 다른 데이터를 가지고 동일한 명령어를 수행한다.