[CUDA] 프로그래밍 모델( kernel, thread )

​

 

 

 

1. Kernel

CUDA에서는 kernel이라 불리는 함수를 정의할 수 있습니다.

일반 C 함수와 kernel의 차이점은 호출시 N개의 CUDA 쓰레드에서 병렬로 실행된다는 것입니다.

kernel은 __global__이라는 specifier와 호출 시 실행될 횟수를 정의하는 <<<...>>> execution configuration syntax로 정의 됩니다.

각 kernel을 실행하는 쓰레드는 고유한 thread ID를 가지며 threadIdx라는 built-in로 접근할 수 있습니다.

​

다음 코드는 N개의 A와 B 백터를 더하는 샘플 코드입니다.

여기서 N개의 쓰레드가 VecAdd를 실행하여 각각 한번씩 더하기를 실행하게 됩니다.

전체 코드는 CUDA sample의 vectorAdd에서 볼 수 있습니다.

2. 쓰레드 구조

threadIdx는 3개의 component를 가지는 vector여서 1,2,3차원의 thread index를 설정할 수 있습니다.

이 쓰레드들의 묶음은 thread block이라 부르며 마찬가지로 1,2,3 차원 형태를 가질 수 있습니다.

위의 thread index와 thread block기능은 vector, matrix, volume을 위한 연산에 매우 잘 어울립니다.

​

thread index와 thread ID는 매우 직관적인 관계에 있습니다.

1차원: threadID = x

2차원:threadID = x + y * Dx

3차원:threadID = x + y * Dx + z * Dx * Dy

x,y,z = thread Index

(Dx,Dy,Dz) = thread block size

​

아래 코드는 N x N 크기의 매트릭스 A와 B를 합해 C에 저장하는 코드입니다.

CUDA는 쓰레드 블럭당 최대 쓰레드 개수가 제한되어 있으며, 요즘 각 블럭당 최대 1024개 정도로 제한되어 있습니다.

이유는 한 블럭의 모든 쓰레드는 하나의 streaming multiprocessor(SM)에서 실행되며 제한된 메모리와 리소스를 공유하기 때문입니다.

*CUDA GPU는 여러개의 SM들을 가지고 있으며, 각 SM은 수십~수백개의 쓰레드를 병렬로 처리합니다.

각 블럭당 쓰레드 개수 제한은 작지만 kernel은 여러개의 블럭을 동시에 실행할 수 있어 "블럭 개수 * 블럭 당 쓰레드  수" 의 쓰레드 들을 동시에 실행할 수 있습니다.

​

쓰레드 블럭들은 1,2,3 차원 형태로 묶을 수 있으며, 이는 grid(그리드)라 합니다.

아래는 그리드와 블럭에 대한 구조를 나타내는 그립니다.

block은 thread의 묶음이며, grid는 block의 묶음 입니다.

grid당 block수 및 block당 thread수는 <<<...>>> syntax로 설정됩니다.

 

ex) Func<<<numBlocks, threadsPerBlock>>>()

 

다음 MatAdd()예제는 다수의 블럭을 다루는 예제입니다.

각 블럭의 크기는 임의로 16x16(256 쓰레드)으로 설정했지만, 이는 자주 사용되는 크기입니다.

그리드는 쓰레드당 matrix의 element하나를 처리할 수 있도록 블록 개수를 설정했습니다.

쓰레드당 element하나를 처리하기 위해서는 NxN 개수의 쓰레드가 필요합니다.

위에서 블럭당 쓰레드수가 16 x 16이므로, (N/16, N/16)개수의 블럭이 필요합니다.

 

각 쓰레드 블럭은 독립적으로 실행 될 수 있도록 구현되어야 합니다.

이는 각 블록들을 병렬로, 순서 및 시간 코어 개수에 관계없이 스케줄링 하기 위해 필요한 조건입니다.

​

블럭 내의 쓰레드들은 shared memory와 synchronization기능 들을 이용해 데이터를 공유 할 수 있습니다.

예를들어 __syncthread() 기능을 이용해 kernel의 동기화 지점을 설정하면, 블럭 내의 모든 쓰레드들은 다른 쓰레드들이 동기화 지점에 도달하기 전까지 다음 연산을 실행하지 않고 대기할 것입니다. 추가로 Cooperative Groups API를 사용하면 더 풍부한 동기화 기능을 사용할 수 있습니다.

블럭내의 효율적인 협력을 위해 shared memory는 L1캐시와 같이 낮은 지연시간의을 가지며 __syncthreads()는 lightweight합니다.

 

thread -> block -> grid

 

 

 

CUDA 이미지 필터(boxBlur)