C++ CUDA Warp 발산

CUDA 연산 시 그래픽카드의 warp size 만큼 동일한 병렬 연산이 이루어 진다.

(오해의 소지가 있으니 자세히 말하자면, GPU의 코어 숫자 만큼 병렬 연산이 이루어 진다.

내 그래픽 카드인 RTX2070 super의 코어 수는 2560개 이므로, 최대 2560개의 코어를 동시에 사용할 수 있다.

warp_size는 동일 연산이 일어나야 하는 권장 코어 숫자라고 생각하면 쉽다.

2560 / warp_size = 80 이다. 따라서 동시에 80개의 warp들이 병렬 연산이 가능 한데,

80개 중 각각의 warp 에서는 동일 연산이 이루어져야 되는게 좋다는 뜻이다.)

 

우리가 평소에 잘 사용하는 NVIDIA GTX/RTX nvidia 계열의 그래픽 카드는 warp size가 32다.

만약에 gpu 내부에서 실행되는 함수에서 if else 같은 분기점이 있어서

16개는 더하기 연산이 실행

16개는 빼기 연산이 실행된다면

32개 사이즈가 동시에 실행될 수 있었는데 16개밖에 동시에 실행이 안 되는 것과 마찬가지이므로

성능 손실이 일어난다.

이런 경우를 warp 발산이라고 한다.

(사실 cuda compiler는 이런 경우에 컴파일 시 최적화를 해 주지만, 복잡한 코드 같은 경우에는 최적화가 잘 안 될 수도 있다.)

 

코드>>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdlib.h>
#include <stdio.h>
 
__global__ void code_without_warp_divergence()
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0;
 
    int warp_id = gid / 32;
    
    if (warp_id % 2 == 0)
    {
        a = 100.0;
        b = 50.0;
    }
    else {
        a = 200.0;
        b = 75.0;
    }
}
 
__global__ void code_divergence()
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0;
    
 
    if (gid % 2 == 0)
    {
        a = 100.0;
        b = 50.0;
    }
    else {
        a = 200.0;
        b = 75.0;
    }
}
 
int main()
{
    int size = 1 << 22;
 
    dim3 block_size(128);
    dim3 grid_size((size + block_size.x - 1) / block_size.x);
 
    code_without_warp_divergence << < grid_size, block_size >> > ();
    cudaDeviceSynchronize();
    code_divergence << <grid_size, block_size >> > ();
    cudaDeviceSynchronize();
 
 
    cudaDeviceReset();
    return 0;
}