WARP Execution

References

• Professional CUDA C Programming
• https://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/issueefficiency.htm

Contents

• Warps 이해하기
• Warp Divergence
• Resource Partitioning
• Latency Hiding
• Occupancy

---

지금까지 CUDA 프로그래밍과 GPU 아키텍처에 대해서 쭉 알아왔는데, Warp에 대해서 자세하게 살펴볼 기회가 없어서 이번에 Warp에 대해서 조금 더 공부하면서 이번 포스팅을 준비했습니다.

 

 

커널이 실행될 때, 소프트웨어의 관점에서는 커널의 모든 스레드들이 병렬로 실행되는 것으로 보입니다. 하지만 하드웨어의 관점에서 살펴보면 모든 스레드들이 물리적으로 동시에 병렬로 실행되는 것은 아닙니다. 이번 포스팅에서는 하드웨어 관점에서의 Warp 실행을 자세히 살펴보겠습니다.

 

---

Warps and Thread Blocks

Warp는 SM(Streaming Multi-processor)의 기본 실행 단위(unit of execution) 입니다. 스레드 블록의 그리드를 실행하면, 그리드의 스레드 블록들은 SM들로 분배됩니다. 스레드 블록이 SM에 스케쥴링되면 스레드 블록의 스레드들은 warp로 파티셔닝됩니다. 32개의 연속된 스레드들로 구성된 하나의 warp는 SIMT(Single Instruction Multiple Thread) 방식으로 실행됩니다. 즉, 모든 스레드는 동일한 명령어를 실행하고, 각 스레드는 할당된 private data에 대해 작업을 수행합니다.

아래 그림은 스레드 블록의 논리적 관점과 하드웨어 관점 간의 관계를 보여줍니다.

스레드 블록은 1,2,3차원으로 구성될 수 있습니다. 하지만, 하드웨어 관점에서 살펴보면, 모든 스레드는 1차원으로 정렬됩니다. 각 스레드는 블록에서 unique ID를 가지고 있습니다. 1차원 스레드 블록에서, 스레드의 unique ID는 CUDA에 내장된 변수인 threadIdx.x에 저장되고, 연속된 threadIdx.x를 가진 스레드들이 warp로 그룹화됩니다.

 

예를 들어, 128개의 스레드로 구성된 1차원 스레드 블록은 다음과 같이 4개의 warp로 파티셔닝됩니다.

 

2차원 또는 3차원 스레드 블록의 논리적 레이아웃은 x 차원을 가장 안쪽의 차원, y 차원을 두 번째 차원, z 차원을 가장 바깥쪽의 차원으로 사용하여 1차원 물리적 레이아웃으로 변환할 수 있습니다.

 

예를 들어, 2D의 스레드 블록이 주어진 경우, 블록의 각 스레드에 대한 고유한 식별자는 내장된 threadIdx와 blockDim 변수를 통해 계산할 수 있습니다.

3D 스레드 블록도 같은 방법으로 아래와 같이 계산할 수 있습니다.

 

따라서, 스레드 블록의 warp 개수는 다음과 같은 공식으로 계산됩니다.

 

하드웨어는 항상 스레드 블록을 warp로 할당합니다. warp는 절대 다른 스레드 블록 사이에서 분리되지 않습니다. 만약 스레드 블록의 크기가 warp 사이즈의 배수가 아니라면, 마지막 warp의 일부 스레드는 비활성 상태가 됩니다.

 

아래 그림은 x 차원에 40개의 스레드, y 차원의 2개의 스레드로 구성된 2차원 스레드 블록을 보여주고 있습니다. SW 관점에서 2차원 그리드에서 80개의 스레드가 레이아웃되어 있습니다.

하드웨어는 이 스레드 블록을 3개의 warp로 할당합니다. 그 결과 총 96개의 하드웨어 스레드가 80개의 소프트웨어 스레드를 지원합니다. 마지막 warp의 절반은 비활성 상태가 됩니다. 비록 이 스레드들은 사용되지 않지만, 이 스레드를 위한 SM 리소스(ex, 레지스터)는 여전히 사용됩니다.

 

SW 관점에서 스레드 블록은 1,2,3차원의 레이아웃의 스레드들로 구성될 수 있습니다.
HW 관점에서 스레드 블록은 1차원의 warp입니다. 스레드 블록의 스레드들은 1D 레이아웃으로 구성되고, 연속된 32개의 스레드들이 각 warp를 형성합니다.

 

---

Warp Divergence

고수준의 프로그래밍 언어에서 Control Flow는 기본적인 구성 중의 하나입니다. GPU는 if..then..else나 for, while과 같은 전통적인 C 스타일의 flow control을 지원합니다.

CPU는 분기(branch) 예측을 수행하기 위한 복잡한 하드웨어를 포함합니다. 만약 예측이 맞다면(true라면), CPU에서 분기는 약간의 성능 저하를 일으킵니다. false라면 instruction pipeline이 flush될 때 CPU가 몇 사이클동안 중지될 수 있습니다. (저도 자세하게는 잘 모르지만, 분기를 처리하는데 약간의 성능 저하가 발생할 수 있다는 정도로만 이해하면 될 것 같습니다.)

 

반면, GPU는 복잡한 분기 예측 메커니즘이 없는 비교적 단순한 디바이스입니다. warp의 모든 스레드들이 반드시 동일한 명령을 같은 사이클에서 수행해야하기 때문에, 만약 하나의 스레드가 명령을 수행하면, warp의 다른 스레드들도 이 명령어를 수행해야 합니다. 따라서, 만약 같은 warp의 스레드들이 다른 분기로 빠져서 명령어를 실행하는 것은 문제가 될 수 있습니다.

한 warp의 16개의 스레드가 cond가 true인 코드를 실행하고, 다른 16개의 스레드는 cond가 false인 코드를 실행한다고 가정해보겠습니다. warp의 절반은 if block의 명령어를 수행할 것이고, 다른 절반의 스레드는 else block의 명령어를 수행하게 될 것 입니다. 이렇게 같은 warp의 스레드들이 다른 명령어를 수행하는 것을 warp divergence라고 합니다. 이미 위에서 warp안의 모든 스레드들은 각 사이클마다 동일한 명령을 실행한다고 했으니, warp divergence는 역설적으로 보입니다.

 

만약 warp 내의 스레드들이 발산(diverge)한다면, warp는 각 branch path를 순차적으로 실행합니다. 해당 path가 아닌 스레드들은 disable됩니다. Warp divergence는 상당한 성능 저하를 야기합니다. 이전 예제에서 warp 내의 병렬 수행은 절반으로 감소합니다. 오직 16개의 스레드만 수행되고 나머지 16개의 스레드는 disable됩니다. 더 많은 조건 분기가 있다면 이러한 loss는 더욱 커집니다.

 

branch divergence는 오직 한 warp 내에서만 발생합니다. 따라서, 다른 warp에서의 다른 조건은 warp divergence를 발생시키지 않습니다.

 

아래 그림은 warp divergence를 보여줍니다. 만약 스레드에서 조건이 true라면 이는 if clause을 실행하지만, 다른 스레드들은 이 실행이 완료될 때까지 중단됩니다.

 

최상의 성능을 얻기 위해서는 같은 warp에서 다른 실행 경로를 피해야합니다. 스레드 블록의 warp 할당은 블록에 의해서 결정됩니다. 따라서, 같은 warp에 있는 모든 스레드들이 동일한 control flow를 갖도록 데이터를 분할하는 것이 가능할 수 있습니다.

 

예를 들어, 아래의 간단한 산술 연산 커널처럼 두 개의 분기가 있다고 가정해보겠습니다. 데이터를 짝수/홀수 스레드 접근(thread approach)으로 파티셔닝할 수 있는데, 이는 warp divergence를 발생시킵니다. (tid %2 == 0) 조건은 짝수 ID의 스레드는 if clause를 실행하고 홀수 ID의 스레드는 else clause를 실행하도록 합니다.

__global__ void mathKernel1(float* c)
{
	int tid = blockDim.x * blockIdx.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;
    
    if (tid % 2 == 0) {
    	a = 100.0f;
    }
    else {
    	b = 200.0f;
    }
    c[tid] = a + b;
}

 

만약 데이터를 warp approach로 파티셔닝한다면, 이러한 warp divergence를 피하고 디바이스의 성능을 극대화할 수 있습니다. ((tid / warpSize) % 2 == 0) 조건은 warp 크기만큼의 스레드들이 한번에 짝수/홀수로 체크되도록 합니다. 따라서, 짝수 warp의 모든 스레드들은 if clause를 실행하고 짝수 warp의 모든 스레드들은 else clause를 실행합니다. 이 커널은 위와 동일한 결과를 생성하지만, 순서가 다릅니다.

__global__ void mathKernel2(float* c)
{
	int tid = blockDim.x * blockIdx.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;
    
    if ((tid / warpSize) % 2 == 0) {
    	a = 100.0f;
    }
    else {
    	b = 200.0f;
    }
    c[tid] = a + b;
}

 

위에서 살펴본 두 개의 커널의 성능을 직접 측정해보겠습니다.

실행 코드는 아래 링크에서 확인하실 수 있습니다.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/simpleDivergence/simpleDivergence.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

nvcc로 컴파일할 때, 최적화 옵션을 주면 컴파일러 최적화가 발생하여 정확하게 커널을 프로파일링할 수 없으므로, '-g'와 '-G' 옵션을 주고 컴파일하면 됩니다.

 

측정은 윈도우에서 Nsight Compute를 사용했습니다. NVIDIA에서 제공하는 이 프로그램을 사용하면, 각 커널에서의 수행 시간과 branch efficiency 등을 측정할 수 있습니다.

만약 CLI로 결과를 간단하게 보고자 한다면, 다음의 커맨드로 프로파일링을 할 수 있습니다.

ncu.bat --metrics smsp__sass_average_branch_targets_threads_uniform.pct,gpu__time_duration.sum ./simpleDivergence.exe

저의 경우 윈도우 환경에서 실행했고, ncu.bat은 Nsight compute가 설치된 폴더에 존재하는 파일입니다.

실행 결과는 다음과 같습니다.

mathKernel1에서는 warp divergence가 발생하기 때문에 성능 하락이 존재하고, 따라서 mathKernel2보다 조금 더 시간이 걸리는 것을 확인할 수 있었습니다.

 

그리고 각 커널에서의 branch_efficiency는 80%, 100% 입니다.

branch efficiency는 다음과 같이 계산할 수 있습니다.

Nsight compute로 branch의 수와 divergent branch의 수를 구할 수 있는데, 이는 다음의 커맨드로 확인하실 수 있습니다.

ncu.bat --metrics smsp__sass_branch_targets.sum,smsp__sass_branch_targets_threads_divergent.sum ./simpleDivergence.exe

mathKernel1의 branch efficiency는 \(\frac{10-2}{10} = 0.8\)로 계산되고, mathKernel2는 \(\frac{9-0}{9}=1.0\)으로 계산됩니다.

 

branch의 수가 이상하다고 생각이 될텐데, 이는 CUDA의 nvcc 컴파일러가 여전히 제한된 최적화를 수행하고 있다는 것을 보여줍니다. mathKernel1의 코드만 보면, 브랜치 효율이 50%가 나와야하는데, 80%로 측정되고 있습니다.

 

---

Resource Partitioning

Warp의 local execution context는 주로 다음의 리소스들로 구성됩니다.

• Program counters
• Registers
• Shared memory

SM에 의해서 실행되는 각 warp의 execution context는 warp의 전체 수명동안 on-chip으로 유지됩니다. 따라서 하나의 execution context에서 다른 context로 전환하는 데 cost가 전혀 없습니다.

 

각 SM은 레지스터 파일(register file)에 저장된 32-비트 레지스터들의 집합들을 가지고 있으며, 스레드 블록들 사이에서 분할된 고정된 크기의 공유 메모리를 가지고 있습니다. 주어진 커널에서 SM에 동시에 상주할 수 있는 스레드 블록과 워프의 수는 SM에서 사용할 수 있거나 커널이 요구하는 레지스터의 수와 공유 메모리 크기에 따라 다릅니다.

 

아래 그림은 각 스레드가 더 많은 레지스터를 소모할 때, 더 적은 수의 warp가 SM에 위치하는 것을 보여줍니다.

만약 커널이 사용하는 레지스터의 수를 감소시킨다면, 더 많은 warp들이 동시에 수행될 수 있습니다.

 

아래 그림은 스레드 블록이 더 많은 공유 메모리를 사용할 때, 더 적은 스레드 블록들이 SM에 의해서 동시에 수행되는 것을 보여줍니다.

마찬가지로 각 스레드 블록에서 사용하는 공유 메모리의 양을 줄일 수 있다면, 더 많은 스레드 블록들이 동시에 수행될 수 있습니다.

 

사용 가능한 리소스는 일반적으로 SM당 상주하는 스레드 블록의 수에 따라 제한됩니다. SM당 레지스터의 수와 공유 메모리의 크기는 디바이스마다 다릅니다. 만약 각 SM에 적어도 하나의 블록을 처리하기에 충분한 레지스터나 공유 메모리가 없다면, 커널 launch는 실패할 것입니다.

조금 오래된 장치들이지만, 아래 표에서 각 디바이스에서 제한되는 리소스들을 보여줍니다.

Resource Limits Pertaining to Compute Capability

본인의 디바이스의 리소스를 확인하고 싶다면, Device Query를 통해 확인하실 수 있습니다. 필요하시다면 아래 포스팅을 참조해주세요.

리소스 동적 분할 및 제한 사항 (+ device query)

리소스 동적 분할 및 제한 사항 (+ device query)

 

레지스터나 공유 메모리와 같은 연산 리소스가 스레드 블록에 할당되었을 때, 이 스레드 블록을 active block이라고 합니다. 이 스레드 블록에 포함된 warp는 active warp라고 합니다.

Active warps는 다음의 3가지 유형으로 분류할 수 있습니다.

• Selected warp
• Stalled warp
• Eligible warp

SM의 warp 스케쥴러는 모든 사이클 때마다 active warps를 선택하여 execution unit에 이들을 dispatch 합니다. 이렇게 실행 중인 워프를 selected warp라고 합니다. active warp가 실행할 준비는 되었지만 현재 실행 중이지 않다면 eligible warp라고 합니다. 만약 워프가 실행할 준비도 되지 않았다면, 이는 stalled warp입니다.

 

워프는 다음의 두 가지 조건을 만족하면 실행할 준비가 완료되어 eligible warp가 됩니다.

• 32개의 CUDA core를 실행할 수 있을 때
• 현재 instruction에 대한 모든 arguments가 준비되었을 때

예를 들어, Kepler 아키텍처의 SM에서의 active warp의 수는 언제든지 아키텍처의 제한인 64개보다 적거나 같아야 합니다. Kepler 아키텍처에서의 selected warp의 수는 어느 사이클에서나 4개보다 적거나 같습니다. 만약 warp가 정지(stalled)되면, warp 스케쥴러는 eligible warp를 골라서 실행시킵니다. 연산 리소스는 워프들 사이에서 파티셔닝되고 워프의 전체 lifetime동안 on-chip에서 해당 리소스들이 유지되기 때문에 warp context를 전환하는 속도는 매우 빠릅니다.

 

CUDA 프로그래밍에서 연산 리소스 분할은 주의깊게 살펴봐야 합니다. 연산 리소스는 active warp의 수를 제한합니다. 그러므로, 하드웨어에 의한 제약과 구현한 커널에 의해서 사용되는 리소스를 잘 아는 것이 중요합니다. GPU를 최대한으로 활용하려면, active warp의 수를 최대한 많이 유지해야합니다.

 

---

Latency Hiding

SM의 기능 단위의 활용을 최대화하기 위해서 스레드 수준의 병렬화에 의존합니다. 따라서, 이러한 활용은 상주하는 warp의 수와 직접적으로 관련이 있습니다. 명령이 실행되고 완료되기까지의 클럭 사이클의 수는 instruction latency로 정의됩니다. Full compute resource 활용은 모든 warp 스케쥴러가 매 클럭 사이클마다 eligible warp를 가지고 있을 때 달성할 수 있습니다. 이는 상주하는 다른 warp에 다른 명령어를 실행시켜서 각 명령어의 latency를 숨길 수 있도록 해줍니다.

 

CPU의 C 프로그래밍과 비교해서, latency hiding은 CUDA 프로그래밍에서 특히 더 중요합니다. CPU 코어는 한 번에 하나 또는 두 개의 스레드에 대한 latency를 최소화하도록 설계되었지만, GPU는 처리량을 극대화하기 위해 많은 수의 스레드를 처리하도록 설계되었습니다. GPU 명령어의 latency는 다른 warp의 연산에 의해서 숨겨집니다.

 

명령어 latency를 고려할 때, 명령어는 두 가지의 기본적인 유형으로 분류될 수 있습니다.

• Arithmetic instructions (산술 명령어)
• Memory instructions (메모리 명령어)

산술 명령어 latency는 산술 명령의 시작과 결과가 생성될 때까지의 시간입니다. 메모리 명령어 latency는 load 또는 store 명령이 실행되고 data가 destination까지 도착할 때까지의 시간입니다.

각 케이스에 대응되는 latency는 대략 다음과 같습니다.

• 10~20 cycles for arithmetic operations
• 400~800 cycles for global memory accesses

 

아래 그림은 warp 0이 stall된 실행 파이프라인의 간단한 예를 보여줍니다. warp 스케줄러는 다른 warp를 실행하도록 선택하고, warp 0이 eligible 되면 warp 0을 실행합니다.

이러한 latency를 감추기 위해 필요한 active warp의 수를 측정하는 방법은 Little's Law를 통해 근사적으로 알 수 있습니다.

 

아래 그림은 Little's Law를 시각적으로 보여줍니다. 커널의 명령에 대한 평균 latency가 5 사이클이라고 가정해봅시다. 매 사이클당 6 warps의 처리량을 유지하기 위해서는 적어도 30개의 warp가 필요합니다.

 

 

산술 연산에서 필요한 병렬(parallelism)은 산술 latency를 숨기기 위해 필요한 연산의 수로 표현할 수 있습니다. 아래 표는 Fermi와 Kepler 디바이스에서 필요한 연산의 수를 보여줍니다. 예시로 사용되는 산술 연산은 32비트 부동소수점의 곱셈-덧셈 (a + b x c)로, 이는 각 SM에서의 clock cycle 당 연산의 수로 표현됩니다. 이 처리량은 산술 명령어마다 다릅니다.

SM Parallelism Required to Maintain Full Arithmetic Utilization

처리량(throughput)은 각 SM당 사이클 당 연산의 수로 명시되어 있고, 하나의 명령을 실행하는 1개의 warp는 32개의 명령에 대응됩니다. 따라서 최대 연산 리소스 활용을 유지하기 위해서 SM당 필요한 warp의 수는 Fermi GPU에서는 640 / 32 = 20 warps로 계산할 수 있습니다. 

즉, 산술 연산에 필요한 병렬은 명령의 수 또는 warp의 수로 표현할 수 있습니다. 이러한 간단한 단위 변환은 병렬성을 증가시키는 두 가지 방법이 있다는 것을 알려줍니다.

• Instruction-level parallelism (ILP): 스레드 내에서 더 많은 independent instructions
• Thread-level parallelism (TLP): 동시에 더 많은 eligible threads

 

메모리 명령에서 필요한 parallelism은 메모리 latency를 숨기기 위해 필요한 사이클당 바이트 수로 표현할 수 있습니다. 아래 표는 Fermi와 Kepler 아키텍처에서의 측정분석 결과를 보여줍니다.

Device Parallelism Required to Maintain Full Memory Utilization

메모리 처리량은 보통 초당 기가바이트로 표현되기 때문에, 먼저 메모리 frequency를 사이클당 기가바이트로 변환해야합니다. 디바이스의 메모리 frequency는 다음의 커맨드로 확인할 수 있습니다.

nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory"

Fermi 아키텍처인 Tesla C2070에서 측정한 메모리 frequency는 1.566GHz 입니다. 1Hz는 초당 1 사이클로 정의되므로, 초당 기가바이트로 변환하면 다음과 같습니다.

결과값에 메모리 latency를 곱하면 Fermi 메모리 연산에 요구되는 parallelism의 값인 74KB에 근사한 값을 유도할 수 있습니다. 메모리 bandwidth는 전체 디바이스에 대해서 주어지므로, 이 값은 SM이 아닌 디바이스 전체에서의 값입니다.

이 값을 warp 또는 thread에 연결짓는 것은 어플리케이션에 따라 다릅니다. 각 스레드가 연산을 위해 Global Memory에서 SM으로 1 float의 데이터(4 bytes)를 전달한다고 가정하면, Fermi GPU의 모든 메모리 latency를 hiding하려면 18,500개의 스레드 또는 579개의 warp가 필요합니다.

Fermi 아키텍처는 16개의 SM을 가지고 있습니다. 그러므로 모든 memory latency를 숨기려면 SM당 579 warps / 16 SMs = 36 warps가 필요합니다. 만약 각 스레드가 4-byte load 이상을 수행한다면, 더 적은 스레드가 요구될 것 입니다.

 

명령어 latency와 비슷하게, 각 thread/warp 내에서 더 독립적인 메모리 작업을 생성하거나 동시에 더 많은 active thread/warps를 생성하여 사용 가능한 병렬 처리를 늘릴 수 있습니다.

Latency hiding은 SM당 active warps의 수에 따라 달라지며, 이는 execution configuration(<<<...>>>)과 리소스 제약(커널 내 사용되는 레지스터와 공유메모리)에 의해서 결정됩니다. 최적의 execution configuration을 선택하는 것은 latency hiding과 resource utilization 사이의 밸런스를 조절하는 문제입니다.

 

바로 이어서 이 문제에 대해서 자세히 살펴보겠습니다.

---

Occupancy

각 CUDA core 내에서 명령어는 순사적으로 실행됩니다. 하나의 warp가 stall되었을 때, SM은 다른 eligible warps를 실행하도록 switch합니다. 이상적으로는 디바이스의 Core를 계속 사용할 수 있을 만큼 충분한 warp를 가지는 것이 좋습니다. 이때, Occupancy(점유율)은 SM 당 active warps와 최대 warps 수의 비율입니다.

다른 포스팅(리소스 동적 분할 및 제한 사항 (+ device query)에서 살펴봤지만, 간단한 device 쿼리 코드를 다시 작성하여, 필요한 정보만을 살펴보겠습니다.

코드는 아래 링크를 참조바랍니다.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/deviceQuery/simpleDeviceQuery.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

SM당 최대 warp 수는 cudaDeviceProp 구조체의 maxThreadsPerMultiProcessor의 값에 warpSize를 나누면 구할 수 있습니다. 제 GPU 디바이스에서는 SM에서 최대 스레드의 수는 1024개이며, warp 사이즈는 32이기 때문에 SM당 최대 warp의 수는 32로 계산됩니다.

 

CUDA Toolkit은 CUDA Occupancy Calculator라는 엑셀파일을 포함하고 있습니다. 이는 CUDA가 설치된 폴더의 tools 폴더에서 확인하실 수 있습니다.

위에서 먼저 본인의 GPU 디바이스의 Compute Capability 정보를 입력하고, 커널의 리소스 사용량을 입력합니다.

compute capability를 지정한 후에, physical limits 섹션의 데이터가 자동으로 채워집니다. 다음으로 아래의 커널 리소스 정보를 입력합니다.

• Threads per block (execution configuration)
• Registers per thread (resource usage)
• Shared memory per block (resource usage)

스레드 당 레지스터와 블록 당 공유 메모리 리소스 사용량은 nvcc로 컴파일할 때 아래의 compile flag를 추가하여 얻을 수 있습니다.

--ptxas-options=-v

위에서 살펴본 simpleDivergence.cu를 위의 옵션을 추가하여 컴파일하면 다음와 출력 결과를 확인할 수 있습니다.

어떻게 이러한 결과로 출력되는지는 아직 잘 모르겠습니다 ㅠ.ㅠ

 

Occupancy를 향상시키기 위해서 스레드 블록 configuration의 크기를 조정하거나 리소스 사용량을 조정하여 더 많은 active warp를 사용하고, 연산 리소스의 활용을 향상시켜야할 필요가 있습니다. 스레드 블록을 극단적으로 조절하면 리로스 활용도가 제한될 수 있습니다.

• Small thread blocks: 블록당 너무 적은 스레드는 모든 리소스를 완전히 활용하기 전에 하드웨어가 SM당 warps의 수를 제한합니다.
• Large thread blocks: 블록당 너무 많은 스레드는 각 스레드가 이용가능한 하드웨어 리소스를 줄입니다.

Guidelines for grid and block size
- 블록당 스레드의 수는 warp size(32)의 배수로 유지
- 작은 block size는 피한다. 블록 당 스레드 수는 적어도 128이나 256부터 시작
- 커널의 리소스 사용량에 맞게 block 크기를 조절한다.
- 디바이스에 충분한 병렬 처리를 제공하기 위해서 SM 수보다 블록 수를 더 많이 유지한다.
- 실험을 통해 최적의 execution configuration과 resource usage를 찾는다.