[CUDA] Thread할당

Cuda kernel 함수를 개시하면, kernel function을 수행하는 thread들이 grid를 형성한다. kernel 함수는 program 실행 중 kernel이 개시될때 생성되는 thread들이 수행할 C문장을 정의한 것이다.  grid내부의 thread들의 구성과 자원할당, 스케줄링을 살펴본다. 

 

CUDA thread의 구성 

 - grid내부 thread들은 모두 동일한 kernel fuction을 실행하기에 서로 구별하고 사용할 data영역을 구별하기 위해 고유한 좌표값(blockIdx, threadIdx)를 사용한다. 

 - blockIdx, threadIdx는 kernel 함수에서 사용할 수 있도록 이미 초기화된 변수로 주어진다. 즉, thread가 kernel 함수를 수행하는 도중 blcokIdx, threadIdx를 참조해 thread의 좌표값을 얻는다. 

 - gridDim, blockDim도 이미 정의된 변수들로 각각 grid, block의 크기를 알려준다. 

[그림1] cuda thread

 그림 1을 기준으로 보면 gird는 N개의 thread blcok으로 구성되어 있으며(0~N-1의 blcokIdx.x) 각 thread block은 M개이 thread로 구성되어 있다. (0~M-1의 threadIdx.x) => 각 grid당 MxN의 thread가 존재한다. 

 

각 code는 threadID=blcokIdx.x * blockDim.x + threadIdx.x의 식으로 계산된 값을 이용해 input data를 읽을 위치와 output data를 쓸 위치를 알 수 있다. 

 

thread 계산의 예)

[그림2] thread 계산의 예

일반적으로 gird는 2차원 배열, block은 3차원 배열의 thread로 구성된다. 이는 kernel이 실행될때, 정해진다. 

[그림3] dimGrid, dimBlock의 선언

다음은 host code에 의해 개시되는 작은 크기의 2차원 grid를 나타내고 있다. 

[그림4] 총 64개의 thread가 있다.

Idx의 용도

" blockIdx와 threadIdx는 동일한 kernel을 실행하는 thread들을 구별하기 위한 수단이며, thread가 작업을 수행할 data영역을 결정하는 것이다. " 

 

만약 blockIdx를 사용하지 않는다면, 각 block은 최대 512개의 thread를 가질 수 있으므로 최대 512개의 Pd원소만 계산할 수 있다. 즉, 16*16의 연산은 가능하지만, 32*32연산은 한번에 처리하지 못한다. => 더큰 matrix를 수용하기 위해 여러개의 thread block을 사용할 필요가 있다. 

 

cuda에서의 행렬 곱 연산 

 - Pd를 정방향 타일로 나눈다. 타일의 모든 Pd원소는 동일한 block의 thread에 의해 계산된다. 

 - thread는(bx*tile_width+tx)를 이용해 Pd의 x index, (by*tile_width+ty)를 이용해 Pd의 y index를 찾는다. 

[그림5] cuda에서의 행렬 곱 연산

각 thread block안에서 수용되는 곱셈의 동작 

[그림6] thread block에서의 연산

- blcok(0,0)의 thread들은 4개의 내적값을 생성. 

- thread(0,0)은 Md의 행0과 Nd의 열0의 내적 연산을 하여 Pd(0,0)을 생성한다. 그림 6은 다수의 block을 사용하도록 수정된 kernel함수이며, blockIdx,threadIdx값을 사용해 계산할 Pd의 row, col을 찾는다. 그림 9은 수정된 kernel을 개시하는 host code이다.

[그림7] 수정되기 전 code

[그림8] block단위의 연산

[그림9] 수정된 host code

Thread의 할당 

kernel이 개시되면 cuda system은 해당되는 thread의 grid를 생성한다. 이 thread들은 block단위로 실행 자원에 할당된다. 

 

 - 각 SM은 block이 필요한 충분한 자원을 가지고 있는 한 최대 8개까지의 block을 할당한다. 예를들어 SM이 30개가 있다면 최대 240개의 block을 동시에 할당할 수 있다. 대부분 grid는 240개보다 더 많은 block을 가지는데 runtime system은 실행해야하는 block list를 관리하면서 이전에 할당된 block이 끝나게 되면 새로운 block을 SM에 할당하는 방식을 사용한다. 

[그림10] thread 할당

그림10을 기준으로 보면, 각 SM에 3개의 thread block을 할당했음을 알 수 있고, 동시에 SM자원 제약 중 하나로 작용한다. SM이 thread와 blcokId의 실행상태 추적을 위한 자원은 필수적이여한다. 

 

Thread 스케줄링 & 지연시간 

thread scheduling은 구현에 속한 개념이기에, 특정 hardware구현에 대해 설명이 필요하다. 하나의 blcok은 하나의 SM에 할당되며, warp라 불리는 32개의 thread unit으로 나눠야한다. SM에서 warp은 thread schdeuling의 단위이다. 

[그림11] warp

그림 11은 하나의 block이 warp로 나누어진것을 볼 수 있다. 각 warp는 연속된 threadIdx값을 갖는 32개의 thread로 구성된다.(0~31 : 1번째 warp, 32~63 : 2번째 warp) 각 block이 256개의 thread를 가지고 있다면 각 block은 256/32 = 8개의 warp를 가지고 있을 것이다. 

 

한 SM에서 많은 warp를 가져야되는 이유는 global memory를 접근하는 것과 같이 지연시간이 긴 연산을 CUDA device가 효율적으로 처리하기 위해서이다. 긴 지연시간이 발생했을때 다른 thread의 작업으로 그 latency를 overlap하기 위해서이다.