스레드 블록 아키텍처(Thread Block Architecture) - 그리드 블록 모델 -2

2.3 스레드 블록 생성 가이드라인

-CUDA 프로그래밍에서 1개의 블록은 최대 512개의 스레드를 가질 수 있다.

-프로그램을 구현할떄 1개 블록이 가지는 스레드의 개수는 32 또는 64의 배수로 지정하는 것이 좋다.

(SM이 32배수 단위로 동작하게 되어 있기 때문에 나머지가 발생하지 않는다. NVIDIA의 사양서는 64배수를 권장한다.)

-최대 블록의 크기는 65,535 X 65,535 까지 지정할 수 있으므로 CUDA프로그래밍의 스레드와 블록의 수는 작업을 분할하는 기준으로 많은 수를 지정하는 것이 좋다. 그러면 GPU는 SM단위로 블록을 처리하면서 작업을 진행하게 된다.

| GPU의 블록 단위 작업 처리

| CUDA 프로그램의 자동 확장성

위의 그림은 SM이 프로그램을 블록 단위로 병렬처리하는 과정을 나타내고 있다. 블록과 스레드를 충분히 큰 수로 지정하면 추후 하드웨어의 성능이 발전했을 때, 기존에 구현한 프로그램의 속도가 증가하는 자동 확장성 (Automatic Scalability) 을 기대할 수 있다.

8개의 블록으로 구성된 CUDA 프로그램이 GPU의 SM(스트리밍 멀티프로세서)의 수가 증가할수록 더 많은 수의 블록을 동시에 처리하여 처리 속도가 빨라지는 것을 나타내고 있다. 

동일한 원리로 블록 안의 스레드 수를 충분히 생성하도록 지정하면 SM 안에 코어 역할을 하는 SP(스트리밍 프로세서)의 수가 증가하였을 때 자동으로 블록의 처리 속도가 빨라지게 된다.

2.4 스레드 블록의 개수 설정

스레드와 블록을  생성하는 것은 커널 함수를 호출할 때 설정할 수 있다.

문법

__global__ void kernel<<< Dg, Db, Ns, S>>>( );

| 1번째 인수 : Dg(dimensions of the grid : type dim3)

 

-그리드의 크기로 그리드 안의 블록 개수를 3차원적으로 지정한다. 

-현재는 3차원의 블록지정을 지원하지 않지 때문에 'Dg.z= 1'로 제한하고 있다.

-dim3형의 변수로 지정하며 생략할 수 없다. 

-Dg(x, y, z)로 각 차원을 지정하여 dim3 Dg(3, 2, 1)과 같이 사용한다.

| 2번째 인수 : Db(dimensions of the block : type dim3)

-블록의 크기로 블록 안의 스레드 개수를 3차원적으로 지정한다.

-dim3형의 변수로 지정하며 생략할 수 없다.

-Db(x, y, z)로 각 차원을 지정하여 dim3 Db(4, 4, 4)와 같이 사용한다.

-x는 1부터 512까지, y는 1부터 512까지, z는 1부터 64 까지 사용할 수 있다.

- x * y * z의 값은 512를 넘으면 안 된다.

| 3번째 인수 : Ns(number of bytes shared memory dynamically allocated / block : type size_t)

-각 블록에서 사용하는 공유 메모리의 크기를 지정한다.

-정수형으로 입력해야 하며 생략 가능하다.

-생략했을 때 디폴트 값으로 '0'이 대입된다.

| 4번째 인수 : S(associated cudaStream_t)

-스트림은 비동기 처리를 하려고 사용한다.

-실행 스트림 번호를 '0' 이상의 정수로 지정한다.

-생략 가능라며 생략했을 경우는 스트림 번호 '0'으로 그 커널 함수가 실행된다.

-스트림 번호 '0'은 비동기 처리를 하지않고 동기 구동을 하는 것을 의미한다.

| 스레드의 생성   

1차원 스레드의 생성	2차원 스레드의 생성
dim3 Dg(16, 1, 1); dim3 Db(32, 1, 1); kernel<<< Dg, Db >>>(a, b, c); or kernel<<< 16, 32 >>>(a, b, c);	dim3 Dg(3, 4, 1); dim3 Db(4, 3, 1); kernel<<<Dg, Db>>>(a, b, c);

2.5 블록과 스레드의 2차원 배치

구현해야 할 프로그램에 따라 처리해야 할 데이터가 2차원 이상인 경우가 많이 있다. 스레드와 블록을 작업의 분할 개념으로 접근하면서 2차원 이상의 데이터에 대해 편리하게 처리하기 위해 CUDA는 블록과 스레드의 2차원 이상 생성을 지원한다.

| 그리드 3 x 4, 블록 4 x 3로 구성한 스레드

| code

dim3 Dg (3, 4, 1);

dim3 Db (4, 3, 1);

kernel <<< Dg, Db >>> (a, b, c);

스레드 - 블록 구조를 2차원으로 구성하게 되면 2차원으로 이루어진 배열 또는 작업을 분할하기 편리하다.

| 2차원으로 분할된 작업

위 그림은 진행해야 할 작업을 12 X 12개의 작업으로 분할한 것을 나타내고 있다. 따라서 NX = 12, NY = 12 이고 제일 작은 사각형 하나는 스레드 하나를 나타내고 있다. 스레드가 4 X 3형태로 모여 하나의 블록을 나타내고 있고, 그림에서는 조금 진한 사각형 하나가 하나의 블록을 의미한다. 이러한 블록 하나가 3 X 4 의 형태로 모여 하나의 그리드를 이루고 있다.

작업을 분할하고 커널 함수 내에서 작업을 수행하기 위해 스레드를 수분하는 인덱스 작업이 필요해진다.

| code

//작업을 분할하는 인덱스

int tid, tx, ty;

tx = blockDim.x * blockIdx.x + threadIdx.x;

ty = blockDim.y * blockIdx.y + threadIdx.y;

tid = NX * ty + tx;