[CUDA] ch3. CUDA Thread Programming

 

 

CUDA는 hierarchical thread architecture를 가지고 있어, thread를 group으로 묶어 실행할 수 있다.

 

 

thread가 GPU에서 어떻게 parallel하게 돌아가는지를 이해하게 된다면 parallel programming code를 write할 수 있고

더 나은 performance를 achive할 수 있을 것이다.

 

 

 

CUDA thread가 GPU에서 작동하는 방식

- parallel and concurrent thread execution

- warp execution

- memory bandwidth issue

- contel overhead

- SIMD operation

 

 

 

 

CUDA threads, blocks, and the GPU

 

CUDA programming의 working unit은 CUDA thread이다.

CUDA thread execution model의 가장 기본은 SIMT방식이다.

따라서 kernel function의 body는 동일한 action을 취하면서 multiple thread로 돌아가게 된다.

multiple thread는 group을 이루어 parallel하게 돌아간다.

 

• thread = working unit
• thread block = group of multiple threads = CTA
• grid = group of thread blocks

 

thread -> thread block -> grid

 

 

kernel에 launch할 때 1개 혹은 그 이상의 thread block들이 SM에서 execute한다.

SM은 multiple thread block을 resource availability에 의존하여 실행할 수 있다.

 

thread block 안에 포함된 thread의 갯수는 다양하고, grid 안에 포함된 thread block 갯수 역시 다양하다.

SM에 할당된 thread block의 모습

 

SM에 할당된 multiple thread block들은 concurrent하게 작동한다.

parallel하게 돌아갈 수 있는 thread block의 갯수는 블록이 요구하는 GPU 리소스의 양과 실제 GPU가 이용할 수 잇는 리소스의 양에 따라 달라진다. 

GPU model에 따라서도 SM의 갯수는 다르다. (Tesla V100의 경우 80개, RTX 2080의 경우 48개)

SM은 스레드를 32개로 묶어 컨트롤한다. 

 

 

• thread = working unit
• warp = group of 32 threads -> basic control unit of GPU
• thread block = group of multiple threads = CTA = configure of 1 or multiple warps
• grid = group of thread blocks

 

SM에 할당된 thread block 내부 warp로 묶여진 thread들

 

warp 내에서 multiple thread는 parallel하게 돌아간다.

therad와 block의 operation은 concurrent하다.

 

 

Exploiting a CUDA block and warp

- CUDA thread scheduling

- implicit synchronization

 

__global__ void index_print_kernel() {
 	int idx = blockIdx.x * blockDim.x + threadIdx.x;
 	int warp_idx = threadIdx.x / warpSize;
 	int lane_idx = threadIdx.x & (warpSize - 1);
    
 	if ((lane_idx & (warpSize/2 - 1)) == 0)
 	//thread, block, warp, lane
 	printf(" %5d\t%5d\t %2d\t%2d\n", idx, blockIdx.x, warp_idx, lane_idx);
}


int main() {
 	int gridDim = 4, blockDim = 128;
 	puts("thread, block, warp, lane");
 	index_print_kernel<<< gridDim, blockDim >>>();
 	cudaDeviceSynchronize();
}

 

위 커널을 실행하면 결과는 다음과 같다.

 

이를 통해 알 수 있는 사실:

thread는 warp size로 launch되고 warp의 실행 순서는 결정되지 않았다.

 

1. Out-of-Order block execution : 스레드 블록의 실행순서는 Out-of-Order이다.

2. Out-of-Order warp index with a thread block : 스레드 블록 내에서 워프의 실행순서는 Out-of-Order이다.

3. Grouped therads executed in a warp 

 

 

Understanding CUDA Occupancy

 

Occupancy

각 SM에 할당된 warp들 중 concurrent하게 실행할 수 있는 최대 warp 갯수, 즉 active한 warp의 비율을 말한다.

 

Occupancy trade-off

- Occupancy가 높다면 GPU utilization이 더 높아질 것이다. - stall된 warp의 latency를 hide할 수 있기 때문

- thread 간 resource contention으로 인해 performance가 저하될 것이다.

 

 

Occupancy의 optimal한 지점을 찾는 이유

- 주어진 GPU 리소스로 warp instruction을 효율적으로 issue하기 위해서이다.

- GPU는 SM에서 multiple한 scheduler를 사용해서 multiple warp를 스케줄링한다,

- multiple warp가 효율적으로 scheduling된다면 GPU는 instruction latency 또는 memory latency를 hide할 수 있다.

 

 

core는 multiple warp로 부터 issue된 instruction들을 continuous하게 실행할 수 있다. 

스케줄링 되지 않은 워프들은 다음 instruction이 issue될 때까지 기다려야 한다.

 

• Theoretical occupancy :  커널 리소스 usage와 GPU SM으로부터 이론적으로 커널의 occupancy를 계산 -> maximum upper-bound occupancy로 볼 수 있다.
• Achieved occupancy : SM에서 실제로 concurrent하게 돌아가는 warp의 갯수 그리고 available한 warp의 최대 갯수이다. 프로파일링을 통해 측정될 수 있다.

 

 

Occupancy tunig - bounding register usage

 

커널 알고리즘이 복잡해지거나 datatype이 double precision을 사용하면 register usage 또한 증가한다.

이 경우 active warp size가 제한되어 occupancy가 떨어진다.

따라서 register usage를 제한함으로써 theoretical ocupancy를 증가시킨다.

 

GPU 리소스 usage를 tuning하는 방법은 __launch_boud__를 kernel function에 qualifier로 사용하는 것이다.

이는 NVCC가 최대 블록 사이즈로 SM 당 최소 스레드 블록 갯수를 갖도록 한다.