Real-Time GPU Fluid Dynamics 번역 (1)

2. Introduction

Here is explained how a CPU fluid simulation can have its performance increased by using CUDA, a GPU solution. Before presenting any implementation details we will review some of the CUDA keypoints. After that, the author explains how to deploy a thread layout for fluid parallel computing. Furthermore, the reader can find GPU optimized alternatives to CPU libraries. Finally, we will compare the CUDA implementation performance to the CPU implementation.

CPU 유체 시뮬레이션이 GPU 솔루션인 CUDA를 사용하여 성능을 높일 수 있는 방법을 설명한다. 구현 세부 정보를 제시하기 전에 몇 가지 CUDA 요점을 검토할 것이다. 그 후, 저자는 유체 병렬 컴퓨팅을 위한 스레드 레이아웃을 효율적으로 사용하는 방법을 설명한다. 게다가, 독자는 CPU 라이브러리의 GPU에 최적화된 대안을 찾을 수 있다. 마지막으로, 우리는 CUDA 구현 성능을 CPU 구현과 비교할 것이다.

3. CUDA Basics

CUDA is a general purpose computing toolkit. Using the collection of functions at the CUDA API, a set of specifiers, and the CUDA compiler, the programmer can parallelize a C function (kernel) for N threads. Since the fluid simulation is based on a grid of vector velocities, we have to pay special attention to the optimization of two-imensional arrays memory accesses.

CUDA는 범용 컴퓨팅 툴킷이다. CUDA API, 지정자 집합, CUDA 컴파일러의 함수 모음을 사용하여 프로그래머는 N개의 스레드에 대한 C 함수(커널)를 병렬화할 수 있다. 유체 시뮬레이션은 벡터 속도의 그리드를 기반으로 하기 때문에 2차원 배열 메모리 액세스의 최적화에 특히 주의해야 한다.

3.1 Memory and Thread Hierarchy

Threads are grouped into blocks that share the same streaming multiprocessor (SM). Thus, they share the same shared memory and the same L1 cache (introduced in Fermi architectures) for that SM. For this reason, it’s of extreme importance that different blocks are as independent from one another as possible, this reduces the penalty associated with synchronization to access the content stored in external shared memory as well as the penalty associated with accessing a higher level of cache. For convenience, each kernel spawns N threads that can be indexed in one dimension (each thread is identified by an unique index, linear layout), two dimensions (each thread is identified by two indexes, grid layout), or even three dimensions. This is really valuable for the programmer. Both memory and thread hierarchy explanations are well depicted in fig. 2 where the reader can see, for two kernels, their respective two grid memory layouts (grid 0 and grid 1). In addition, the grid-0 thread layout and indexing are also illustrated at the left.

스레드는 동일한 스트리밍 멀티프로세서(SM)를 공유하는 블록으로 그룹화된다. 이러한 이유로, 서로 다른 블록이 가능한 한 서로 독립적이라는 것이 매우 중요하다. 이렇게 하면 외부 공유 메모리에 저장된 컨텐츠에 액세스하기 위한 동기화 관련 패널티뿐만 아니라 더 높은 수준의 캐시에 액세스하기 위한 패널티도 줄어든다. 편의를 위해 각 커널은 1차원(각 스레드는 고유 인덱스, 선형 레이아웃으로 식별됨), 2차원(각 스레드는 2개의 인덱스, 그리드 레이아웃으로 식별됨), 심지어 3차원까지 인덱싱할 수 있는 N개의 스레드를 생성한다. 이것은 프로그래머에게 정말 가치가 있다. 메모리와 스레드 계층 설명은 그림 2에 잘 나타나 있는데, 여기서 독자는 두 커널에 대해 각각의 두 그리드 메모리 레이아웃(그리드 0과 그리드 1)을 볼 수 있다. 또한 왼쪽에는 grid-0 스레드 레이아웃과 인덱싱도 나와 있다.

그림 2: 스레드 블록 그리드 (왼쪽), 메모리 계층 구조(오른쪽, CUDA Toolkit 설명서 참조)

When the blocks are totally independent from each other the block abstraction benefits are more clear: blocks can be distributed across several GPUs, at different machines.

One final note, because a block is restrained to a SM, the maximum number of threads per block is limited to the maximum number of threads the SM supports (Up to 1024 threads in modern GPUs).

The numbers of cores per SM is called the warp size, which is 32 in all (See Table 13. Technical Specifications per Compute Capability at CUDA Toolkit Documentation) NVidia multiprocessors.

블록이 서로 완전히 독립적일 때 블록 추상화 이점은 더 명확하다. 블록은 여러 GPU에 분산되어 서로 다른 기계에서 사용할 수 있다. 마지막으로, 블록이 SM으로 제한되기 때문에 블록당 최대 스레드 수는 SM이 지원하는 최대 스레드 수(현대 GPU에서는 최대 1024개 스레드)로 제한된다. SM당 코어 수를 워프 크기라고 하는데, 이는 모든 NVidia 멀티프로세서에서 32이다.(Technical Specifications per Compute Capability at CUDA Toolkit Documentation 테이블 13 참조)

3.1.1 Two-Dimensional Arrays Memory Access

We commonly find memory access patterns were we access a grid with two indexes, x and y.

일반적으로 메모리 액세스 패턴은 x와 y라는 두 개의 인덱스로 그리드에 액세스하는 경우를 찾을 수 있다.

// host-equivalent to: memPtr = malloc(sizeof(cData)*width*height)
cudaMalloc(&memPtr, sizeof(cData)*width*height);
...
array_index = width*y + x

Because of inherent design characteristics of memory and cache lines, the width of the thread block, as well as the 2D array width, must be a multiple 5of the warp size in order to ensure a maximum performance, as is described in the CUDA guide of best practices. This, for example, will ensure that all warps (SM) can access 128 bytes chunks of data in only one L1 cache request and do not need more requests as happens in fig. 3.

메모리와 캐시 라인의 고유한 설계 특성 때문에, 최고의 성능을 보장하기 위해 스레드 블록의 폭과 2D 어레이 폭은 워프 크기의 배수여야 한다. 예를 들어, 모든 워프(SM)가 하나의 L1 캐시 요청에서 128바이트의 데이터 청크에 액세스할 수 있으며 그림 3과 같이 더 많은 요청이 필요하지 않다.

그림 3: 2개의 128바이트 L1 캐시 라인에 맞는 정렬되지 않은 순차 주소( CUDA Toolkit Documentation 참조)

From the programmer point of view, all he has to do is to call cudaMallocPitch() when he needs to allocate a 2D array of memory: this function will make sure the memory is properly aligned for best performance.

프로그래머의 관점에서 볼 때, 해야 할 일은 2D 메모리 배열을 할당해야 할 때 cudaMallocPitch()를 호출하는 것이다.

그림 4: 패딩이 추가되어 블록 크기가 워프 크기의 배수일 때 스레드 블록의 메모리 액세스가 L1 캐시에 맞춰진다.

Because of this memory alignment, a padding is introduced, and we no longer index memory positions with the data width: we have to use the pitch value (provided by cudaMallocPitch) to access memory. This is shown in fig. 4 (As the reader can notice, what we are calling pitch is what we were calling pdx when we had to iterate though an in-place fast Fourier transform array. Please, refer to the Real-Time CPU Fluid Dynamics[1] paper). The code we explained at the beginning of this section transforms to:

이러한 메모리 정렬로 인해 패딩이 도입되고 더 이상 메모리 위치를 데이터 크기로 인덱싱하지 않고 메모리에 접근하려면 피치 값(cudaMallocPitch에서 제공)을 사용해야 한다. 이것은 그림 4에 나와 있다(독자가 알아차릴 수 있듯이, 우리가 pitch라고 부르는 것은 내부 고속 푸리에 변환 배열을 통해 반복해야 할 때 pdx라고 부르는 것이다). Real-Time CPU Fluid Dynamics[1] 문서를 참조하십시오. 따라서 이 섹션의 시작 부분에서 설명한 코드는 다음과 같이 변환된다.

cudaMallocPitch(&memPtr, &pitch, sizeof(cData)*width, height);
...
array_index = pitch*y + x

Anyway, the programmer still have to make sure the block size is multiple of 32.

그래도, 프로그래머는 여전히 블록 크기가 32의 배수인지 확인해야 한다.

3.2 Kernels

A kernel is defined adding the global specifier before a function definition. When launching the kernel, the programmer can specify the number of blocks (NB) and threads per block (TPB) with <<NB,TPB>> in the function call after the function name and before the parenthesis for arguments. In order to use a thread (and block) grid indexing layout, dim2 syntax (see code sample below) need to be used. The function is executed NB · T P B times in parallel. blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y are the variables which individually identify each thread in the grid of blocks and threads. These identifiers will help us to access the matching array index position for each thread (in case there is 1 thread for each array position).

커널은 함수 정의 앞에 전역 지정자를 추가하는 것으로 정의된다. 프로그래머는 커널 함수 호출에서 함수 이름과 파라미터에 대한 괄호 사이에 <<<NB,TPB>>>로 블록의 수(NB)와 블록당 스레드 수(TPB)를 지정할 수 있다. 스레드(및 블록) 그리드 인덱싱 레이아웃을 사용하려면 dim2 구문(아래 코드 샘플 참조)을 사용해야 한다. 커널 함수는 NB · TPB를 병렬로 실행한다. blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y는 블록과 스레드 그리드에서 각 스레드를 개별적으로 식별하는 변수이다. 이러한 식별자는 각 스레드에 대해 일치하는 배열 인덱스 위치에 액세스하는 데 도움이 된다(각 배열 위치에 대해 스레드가 1개인 경우).

__global__ void MatAdd(float A[NN], float B[NN], float C[NN]) {
    // If there is only one block: blockIdx.x and blockIdx.y are always 1
    int array_index = N*threadIdx.y + threadIdx.x
    C[array_index] = A[array_index] + B[array_index];
}
int main() {
    ...
    dim2 numBlocks(1, 1); // in this case only 1 block: (1,1)
    dim2 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}