스레드 - 워프 - 블록 - 그리드 -SM - 스케줄링

1. 커널 함수 = 설계도(코드)

__global__ void gemm_kernel(...) {
    // 여기에 있는 코드가 모든 스레드가 공유해서 실행하는 ‘설계도’
}

 

2. 스레드 = 커널 코드를 실행하는 로봇(복제본)

• 커널이 launch 되면 GPU 는 스레드를 수천~수십만 개 만든다
• 스레드마다 threadid.x/y/z, blockidx.x/y/z 가 다르다
• 스레드는 자기만의 레지스터 파일, 자기만의 변수, 자기만의 Program Counter 를 가진다

스레드 = 커널 코드 + 다른 인덱스를 가지고 실행되는 작은 CPU

 

3. 워프 = 32 개의 스레드를 묶은 실행 단위

• NVIDIA GPU 에서는 항상 32개의 스레드가 묶여서 동작
• 32 개 묶음을 warp 
• warp 내 스레드들은 동일한 line 의 코드를 동시에 실행한다

 

4. 블록 = 워프 묶음

block 은 논리적 그룹

• shared memory 를 공유
• _syncthreads() 로 동기화
• 서로 데이터를 주고 받을 수 있음
• block 은 같은 SM 에서만 실행됨

blockDim = (8, 32) 같은 형태면 총 256 threads = 8 warps

 

5. 그리드 = 모든 블록의 집합

<<<gridDim, blockDim>>>

gridDim 은 블록의 배치, 각 블록은 독립적

 

6. SM Streaming Multiprocessor = 물리적 실행 엔진

GPU 는 여러 개의 SM 을 가지고 있고,

각 SM 은 동시에 여러 워프를 스케줄링한다.

SM 은 다음을 보유

• 레지스터 파일
• shared memory
• warp scheduler
• ALU, load/store unit
• Tensor Core

즉, SM = 워프들을 스케줄링하고 명령을 실행하는 프로세서

 

7. 스케줄링 : warp 단위로 명령을 실행 

한 SM 은 동시에 많은 warp 를 보유할 수 있음

1. warp scheduler 가 실행 가능한 warp 를 선택
2. 그 warp 의 다음 명령을 issue
3. memory stall 이 생기면 다른 warp 로 전환

 

 

우리는 block 단위로 프로그래밍, 실제 실행은 warp 단위로 돌아감, block 의 구성이 직접적으로 스케줄러의 동작 방식을 바꾸진 않지만, 스케줄링 가능성은 block 구조에 의해 간접적으로 영향을 받음

 

사용자 관점 : block 단위 프로그래밍

• shared memory 는 block 단위로 공유
• _syncthreads 도 block 단위로 작동
• block 은 grid 단위로 반복해서 실행
• block 크기를 우리가 정함 

사용자 시각에는 스레드들은 block 단위로 묶여서 행동

 

 

GPU 현실 : warp 단위 실행 (scheduling)

실제 명령 launch 는 항상 warp 단위, 

 

여러 블록들이 동시에 여러 SM 에서 병렬 실행 가능한 점을 생각하면, block 에 대한 실행 단위 를 정리할 수 있을듯?

 

GPU 에는 여러 개의 Streaming Multiprocessor 가 존재,

각 SM 은 여러 블록을 수용 가능

 

한 SM 안에서조차 여러 블록이 동시에 실행될 수 있음

한 블록의 shared memory 와 레지스터 수가 적다면, SM 은 여러 블록을 동시에 탑재 가능