CUDA Thread 동기화 및 스케쥴링 / 리소스 할당

References

• Programming Massively Parallel Processors

Contents

• 동기화(Synchronization)
• Transparent Scalability
• 리소스 할당(Resource Assignment)
• Thread Scheduling (Warp)

---

Synchronization & Transparent Scalability

CUDA는 같은 블록 내에 있는 스레드들이 barrier synchronization(배리어 동기화) 함수인 __syncthreads()를 사용하여 스레드들의 작업을 조율할 수 있게 합니다. __syncthreads()를 호출한 스레드는 호출한 위치에서 멈추게 되는데, 블록 내의 모든 스레드들이 해당 위치에 도달할 때까지 멈추게 됩니다. 이렇게 하면 블록의 모든 스레드들은 커널 내에서 다음 단계의 작업을 진행하기 전에 현재 단계의 작업을 마무리할 수 있게 해줍니다.

 

배리어 동기화는 병렬 작업을 조정하기 위해서 간단하면서 널리 사용되는 방법입니다. 실생활에서도 여러 사람에 의해 병렬적으로 수행되는 작업을 조정하기 위해서 배리어 동기화를 사용하는데, 예를 들어, 4명의 친구가 한 차를 타고 다같이 백화점에 간다고 가정해봅시다. 친구들은 각자 다른 가게에 가서 옷을 살 수 있는데, 이를 병렬 작업이라고 할 수 있습니다. 그리고 쇼핑이 끝난 후 백화점을 떠나기 위해서는 배리어 동기화가 필요합니다. 백화점을 떠나기 전에 4명의 친구들이 모두 차에 돌아올 때까지 기다려야 합니다.

 

아래 이미지는 배리어 동기화가 어떻게 수행되는지 보여주고 있습니다.

barrier synchronization

블록에 N개의 스레드가 있고, 시간은 왼쪽에서 오른쪽으로 흐르고 있습니다. 몇몇의 스레드는 배리어 동기화 구문에 일찍 도착하고 몇몇은 훨씬 더 늦게 도착합니다. 배리어에 일찍 도착한 스레드들은 늦게 도착하는 스레드들을 기다리게 됩니다. 그리고, 마지막 스레드가 배리어에 도달하면, 모든 스레드들이 이제 다음 작업을 시작하게 됩니다.

 

CUDA에서 __syncthreads() 구문이 존재하면, 반드시 블록 안에 있는 모든 스레드들은 이 구문을 실행해야합니다. 만약 __syncthreads() 구문이 if문 내부에 있다면, 모든 스레드들은 if문을 포함하는 경로를 수행하거나 아니면 모두 다른 경로를 택해야합니다. 만약 if-else 구문에서 두 경로에 모두 __syncthreads()를 가지고 있다면, if로 분기하는 스레드와 else로 분기하는 스레드가 서로 다른 __syncthreads()를 수행하게 됩니다. 이 두 개의 __syncthreads()는 서로 다른 배리어 동기화 지점이므로, 블록의 한 스레드는 if 경로를 수행하고, 다른 스레드가 else 경로를 실행하는 경우 두 개의 스레드는 서로 다른 동기화 지점에서 기다리게 되므로 서로를 영원히 기다리게 됩니다. (아래에서 몇 가지 예제로 조금 더 살펴보도록 하겠습니다.)

 

이러한 동기화 기능은 블록 내의 스레드 실행에 제약을 가하게 됩니다. 스레드들은 거의 동일한 시점에 수행되어야 오래 기다리는 것을 피할 수 있습니다. 실제로, 배리어 동기화에 관여하는 모든 스레드들이 배리어에 도달하는 데 필요한 리소스에 액세스할 수 있는지 확인해야 합니다. 그렇지 않다면, 배리어 동기화 지점에 도달하지 못하는 스레드로 인하여 다른 모든 스레드들이 영원히 대기상태로 빠질 수 있습니다. CUDA 런타임 시스템은 모든 스레드들을 블록 단위로 실행 리소스(execution resources)에 할당함으로 이러한 제약을 만족시킵니다. 런타임 시스템이 수행을 완료하기 위한 블록의 모든 스레드에 필요한 리소스를 확보하고 있을 때만 블록의 실행을 완료할 수 있습니다. 블록에 있는 어떤 스레드가 실행 리소스에 할당된다면 동일한 블록 내의 모든 스레드들도 동일한 리소스를 할당받습니다. 이렇게 함으로써 블록 내의 모든 스레드들은 거의 동일한 시점에 실행되며 배리어 동기화로 인하여 발생할 수 있는 긴 대기(지연) 시간을 피할 수 있습니다.

 

이는 CUDA 배리어 동기화의 설계에 있어서 중요한 tradeoff를 보여줍니다. 서로 다른 블록의 스레드 사이에는 서로 배리어 동기화를 하지 않도록 하면 그 스레드들은 서로를 기다리지 않으므로 CUDA 런타임 시스템은 블록을 임의의 순서대로 실행할 수 있습니다. 이러한 유연함(flexibility)은 아래 이미지처럼 확장 가능한 구현을 가능하게 합니다.

왼쪽의 device처럼 실행 리소스가 크지 않은 저가의 시스템에서는 한 번에 2개의 블록씩을 실행할 수 있고, 오른쪽의 고가의 device처럼 실행 리소스가 큰 시스템에서는 한 번에 4개의 블록을 동시에 실행할 수 있습니다. 

이처럼 동일한 프로그램의 코드가 다양한 리소스에 맞춰서 실행될 수 있기 때문에 특정 산업에서 요구하는 가격과 전력, 성능에 따라서 구현을 다양화할 수 있습니다. 예를 들어, 모바일 프로세서에서 어플리케이션은 느리게 실행되지만 매우 낮은 전력 소비로 실행할 수 있고, 데스크탑 컴퓨터에서는 동일한 어플리케이션을 더 많은 전력을 소비하면서 더 빠른 속도로 실행할 수 있습니다. 둘 다 코드 변경없이 완전히 동일한 어플리케이션을 실행합니다. 이렇게 리소스가 다른 하드웨어에서 동일한 어플리케이션 코드를 실행하는 기능을 transparent scalability라고 합니다. 이러한 특성은 개발자의 부담을 감소시켜주고, 어플리케이션의 사용성을 향상시켜줍니다.

---

간단한 예제 코드로 동기화에 대해 살펴보도록 하겠습니다.

myKernel 커널 함수는 함수가 수행되는 스레드의 위치를 출력하도록 작성되어 있습니다.

#include <stdio.h>
#include <cuda_runtime.h>

__global__
void myKernel(void)
{
    printf("THREAD %d in BLOCK (%d, %d)\n", threadIdx.x*blockDim.y + threadIdx.y, blockIdx.x, blockIdx.y);
}

int main(void)
{
    dim3 dimBlock(2, 2);
    dim3 dimGrid(2, 2);

    printf("CUDA kernel launch with (%d x %d) blocks of (%d x %d) threads\n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y);
    myKernel<<<dimGrid, dimBlock>>>();

    return 0;
}

그리고, main 함수에서는 각 블록이 2x2 스레드를 갖고, 총 2x2 블록의 그리드를 생성하여 커널을 수행하도록 했습니다. 이 코드를 컴파일 후 실행하면 어떻게 될까요?

main 함수에서의 출력만 나타나고, 커널 함수의 출력은 나타나지 않았습니다. 이를 보면, host code에서 호출된 커널은 device에서 수행되는데, host code가 device와는 별도로 계속 진행한다는 것을 보여주고 있습니다. 따라서, host code에서 커널이 실행됬지만, 커널 함수가 실행되기 전에 main문의 return을 만나서 프로그램이 종료했습니다.

커널이 정상적으로 출력하기 위해서 우리는 host 코드에서 호출된 커널이 모두 완료되기 전까지 return에 도달하지 못하도록 대기시켜야 합니다. CUDA 런타임 API에서는 host code에서 device와 동기화를 시켜주기 위한 cudaDeviceSynchronize()를 제공합니다. main 함수를 아래와 같이 수정하고 다시 컴파일 후 실행해보겠습니다.

int main(void)
{
    dim3 dimBlock(2, 2);
    dim3 dimGrid(2, 2);

    printf("CUDA kernel launch with (%d x %d) blocks of (%d x %d) threads\n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y);
    myKernel<<<dimGrid, dimBlock>>>();
    cudaDeviceSynchronize();

    return 0;
}

이제 커널이 시작된 후에 host code에서는 cudaDeviceSynchronize가 호출된 위치에서 대기하다가 각 스레드의 커널 함수가 모두 수행된 후에 프로그램을 종료합니다. 정확하게는 함수 호출 이전에 시작된 모든 연산이 종료될 때까지 host code를 대기시킵니다.

 

다시 __syncthreads() 를 살펴보기 위해서, 이번에는 배열을 뒤집어주는 reverse 커널을 한 번 작성해보겠습니다.

#define N 256

__global__
void reverse(int *arr)
{
    __shared__ int temp[N];
    int idx = threadIdx.x;
    int idx_inv = N - idx - 1;
    temp[idx] = arr[idx];

    arr[idx] = temp[idx_inv];
}

굳이 이렇게 해야되나 싶지만... __syncthreads()의 효과를 살펴보기 위해서 위 코드를 사용했습니다. 이 커널 함수에서는 device memory에 배열 하나만을 사용하여, 내부에서 shared memory에 임시 배열을 하나 만들어주고, 이 임시 배열에 원래 배열의 값을 저장한 뒤에 원래 배열에 임시 배열의 값을 역순으로 넣어줍니다.

main문을 다음과 같이 작성하고 컴파일 후 실행해보도록 하겠습니다. main 함수에서 커널은 N개의 스레드를 갖는 1개의 블록의 그리드를 생성합니다.

int main(void)
{
    int h_arr[N], result_arr[N], comp_arr[N];

    for (int i = 0; i < N; i++)
    {
        h_arr[i] = i;
        comp_arr[i] = N - i - 1;
    }

    int *d_arr;
    cudaMalloc((void**)&d_arr, N * sizeof(int));

    // run version with static shared memory
    cudaMemcpy(d_arr, h_arr, N * sizeof(int), cudaMemcpyHostToDevice);
    reverse<<<1, N>>>(d_arr);
    cudaMemcpy(result_arr, d_arr, N * sizeof(int), cudaMemcpyDeviceToHost);
    for (int i = 0; i < N; i++) {
        if (result_arr[i] != comp_arr[i]) {
            printf("Error: result_arr[%d]!=com_arr[%d] (%d, %d)\n", i, i, result_arr[i], comp_arr[i]);
        }
    }

    cudaFree(d_arr);

    return 0;
}

실행해보면, 위와 같은 결과를 확인하실 수 있습니다. reverse 커널 내에서 공유 메모리인 임시 배열에 값을 넣어주는 작업이 각 스레드에서 수행되고, 마지막에 원본 배열의 값에 뒤집어진 값을 넣어줍니다. 즉 아래 코드에서 line 9에서 원래 배열 arr에 temp 배열의 값을 넣어주기 위해서는 이 구문이 수행되기 전에 temp 배열은 arr 배열의 값을 모두 담고 있어야합니다. 하지만, 스레드마다 미세하게 수행 시점이 다르기 때문에 temp에 배열 값이 저장되지 않았는데, line 9에서 저장되지 않은 temp값을 참조하여 원본 배열에 값을 써주면서 위와 같은 에러가 발생하는 것입니다. 

__global__
void reverse(int *arr)
{
    __shared__ int temp[N];
    int idx = threadIdx.x;
    int idx_inv = N - idx - 1;
    temp[idx] = arr[idx];
    
    arr[idx] = temp[idx_inv];
}

따라서, temp에 arr 값을 넣어주고 난 후에 스레드는 다른 모든 스레드들이 line 7의 작업을 완료할 때까지 대기해주어야 합니다. 이때, __syncthread()를 사용합니다.

__global__
void reverse(int *arr)
{
    __shared__ int temp[N];
    int idx = threadIdx.x;
    int idx_inv = N - idx - 1;
    temp[idx] = arr[idx];
    
    __syncthreads();

    arr[idx] = temp[idx_inv];
}

수정된 reverse 커널 함수입니다. 이제 다시 컴파일 후, 실행해보도록 하겠습니다.

에러없이 프로그램이 정상적으로 종료된 것을 확인할 수 있습니다.

 

---

Resource Assignment

커널이 실행되면, CUDA 런타임 시스템은 해당되는 스레드들의 그리드를 생성합니다. 위에서 언급했듯이, 이 스레드들은 블록 단위로 실행 리소스에 할당됩니다. GPU 하드웨어는 Streaming Multiprocessors(SMs)로 구성됩니다. 아래 이미지는 각 SM에 할당되는 다중 스레드 블록을 보여주고 있습니다.

Thread block assignment to Streaming Multiprocessors(SMs)

GPU 하드웨어마다 SM당 할당할 수 있는 블록의 수는 제한되어 있습니다.

예를 들어, 위 이미지처럼 하나의 SM에 최대 8개의 블록을 할당할 수 있는 CUDA device가 있다고 가정해보겠습니다. CUDA 런타임은 만약 8개의 블록을 동시에 실행하는데 필요한 리소스가 부족하면, 제한된 리소스에 맞춰서 각 SM에 할당되는 블록 수를 자동으로 줄입니다. SMs의 수가 제한되고 각 SM에 할당되는 블록의 수가 제한되면, CUDA 장치에서 실행할 수 있는 블록의 수도 제한됩니다(대부분의 그리드는 8개의 블록보다는 더 많은 블록을 포함합니다.). 런타임 시스템은 실행해야하는 블록의 목록을 관리하고 이전에 할당된 블록이 완료되면 SM에 새로운 블록은 할당합니다.

 

위 이미지는 각 SM에 3개의 블록이 할당된 예시를 보여줍니다. SM 리소스 제한 중의 하나는 동시에 추적되고 스케쥴링될 수 있는 스레드의 개수입니다. SM은 스레드와 블록 인덱스를 관리하고 이들의 실행 상태를 추적하기 위해서 하드웨어 리소스(built-in registers)를 가지고 있습니다. 따라서, 각 세대의 하드웨어에서 하나의 SM에 할당할 수 있는 블록과 스레드의 수가 제한되어 있습니다. 예를 들어, Fermi architecture에서는 각 SM에 8개의 블록과 1536개의 스레드까지 할당할 수 있습니다. 따라서, 각 SM에 256 스레드를 가지는 6개의 블록이나, 512 스레드를 갖는 3개의 블록을 할당할 수 있습니다. 각 SM은 8개의 블록까지 허용되므로 128 스레드를 갖는 12개의 블록을 할당하는 것은 불가능합니다. 

만약 CUDA device가 30개의 SM을 갖고 있고 각 SM에 1536개의 스레드를 할당할 수 있다면, 이 device는 최대 46,080개의 스레드를 동시에 수행할 수 있습니다.

 

---

Thread Scheduling and Latency Tolerance

스레드 스케쥴링은 엄밀히 말하면 구현 컨셉(implementation concept)입니다. 따라서, 스레드 스케쥴링은 특정 하드웨어 구현의 맥락에서 논의되어야 합니다. 지금까지 나온 하드웨어의 대부분의 경우, SM에 할당된 블록은 warps(워프)라고 불리는 32개의 스레드 유닛으로 더 나뉘어집니다. 워프의 사이즈는 구현에 따라 다른데, 사실 워프는 CUDA 명세에 속하지 않습니다. 하지만 워프를 이해하면 특정 세대의 CUDA device에서의 어플리케이션 성능을 이해하고 최적화하는데 큰 도움이 될 것입니다. 

Blocks are partitioned into warps for thread scheduling.

SM에서 워프는 스레드 스케쥴링의 단위입니다. 위 이미지는 구현에 블록들이 워프로 나누어진 것을 보여줍니다. 각 워프는 연속된 threadIdx 값을 갖는 32개의 스레드로 구성되어 있습니다. 첫 번째 워프에는 0~31 스레드가 속하고, 두 번째 워프에는 32 ~ 63 스레드가 속합니다. 위 이미지에서는 3개의 블록이 SM에 할당되는 것을 보여줍니다. 각 블록은 스케쥴링의 목적으로 워프 단위로 더 나누어집니다.

SM에 할당된 블록의 수와 블록의 크기가 주어진다면 SM에 있는 워프의 수를 계산할 수 있습니다. 예를 들어, 각 블록에 256개의 스레드가 있다면, 256/32를 계산하여 각 블록에는 8개의 워프가 존재합니다. SM에 3개의 블록이 있으므로, 각 SM은 24개의 워프가 있습니다.

 

SM은 다중 데이터(SIMD) 모델이며, 워프에 속한 모든 스레드들을 Single Instruction으로 실행할 수 있도록 설계되었습니다. 그래서 각 스레드에 할당된 데이터에 동일한 명령어를 적용할 수 있습니다. 결과적으로 워프의 모든 스레드들은 항상 동일한 실행 타이밍을 갖게됩니다.

 

위 이미지는 실제로 명령을 실행하는 여러 Streaming Processors(SPs)를 보여주고 있습니다. NVIDIA에서는 이 SP를 CUDA Core라고 부릅니다. 일반적으로 각 SM에 할당된 수보다 SP의 수가 훨씬 적습니다. 따라서, 각 SM은 어느 시점에서든 SM에 할당된 모든 스레드의 부분 집합(일부)만 동시에 실행할 수 있는 정도의 하드웨어만 가지고 있습니다. 초기 GPU 설계에서는 각 SM은 주어진 순간에 하나의 워프에 대해 하나의 명령어만 실행할 수 있었는데, 최근에는 각 SM은 언제든지 적은 수의 워프에 대해 명령을 수행할 수 있습니다. 이는 어떤 하드웨어든지 SM에 있는 모든 워프의 들의 하위 집합에 명령을 실행할 수 있다는 것을 의미합니다.

 

그렇다면, 한 SM에 이렇게 많은 워프를 가져야할 필요가 있을까요?

이에 대한 답은 global memory에 액세스하는 것과 같이 지연시간이 오래 걸리는 연산을 CUDA 프로세서가 어떻게 효율적으로 실행시킬 수 있는가에 대한 답과 같습니다.

워프에 의해 수행되어야할 명령어가 이전에 시작된 긴 지연시간을 갖는 연산의 결과를 기다려야할 때, 그 워프는 실행하도록 선택되지 않습니다. 대신 더 이상 결과를 기다리지 않는 또 다른 워프를 실행하도록 선택합니다. 만약 둘 이상의 워프가 실행을 하기 위한 준비가 된 경우, 우선순위 메커니즘을 사용하여 실행할 워프를 선택합니다. 이렇게 연산의 긴 지연시간을 다른 스레드의 작업으로 채우는 이러한 메커니즘을 'latency tolerance' 또는 'latency hiding(지연시간 은닉)'이라고 합니다.

 

워프 스케쥴링은 파이프라인되어 수행되는 부동소수점(floating-point) 연산이나 분기 명령과 같은 다른 종류의 지연시간이 긴 연산을 감내하기 위해서도 사용될 수 있습니다. 충분한 수의 워프가 주어진다면, 하드웨어는 언제든지 실행할 수 있는 워프를 찾을 수 있고, 따라서 긴 지연시간이 발생하더라도 하드웨어를 충분히 활용할 수 있습니다. 실행이 가능한 워프를 선택하는 것은 실행 시간을 전혀 지연시키지 않는데, 이를 'zero-overhead thread scheduling'이라고 합니다. 워프 스케쥴링에 의하여 워프 명령의 긴 대기 시간은 다른 워프들의 명령을 실행함으로써 숨길 수 있습니다. 이렇게 긴 대기 시간을 감내하는 기능 때문에 GPU는 보통의 CPU가 캐시 메모리와 분기 예측을 위해 사용하는 칩의 면적만큼 필요하지 않습니다. 그렇기 때문에 GPU는 칩 면적의 대부분을 부동소수점 실행 리소스에 할당할 수 있습니다.

 

---

 

저에게는 조금 익숙하지 않은 하드웨어 개념들이 많아서, 이해하는데 조금 어려웠고 지금도 사실 완벽하게 이해하고 있는 것 같지는 않습니다... ㅠ

다음 포스트부터 CUDA 예제를 조금 살펴볼 예정인데, 예제를 살펴보며 조금 더 하드웨어 리소스 관련하여 연관지어서 생각하는 시간을 가져보도록 해야겠습니다.