11장: 스칼라화 및 어파인 실행에 대한 GPU 연구 방향
2장에서 설명한 바와 같이, CUDA와 OpenCL과 같은 GPU 컴퓨팅 API는 프로그래머가 GPU에 많은 수의 스칼라 스레드를 실행할 수 있는 MIMD 유사 프로그래밍 모델을 제공합니다. 이러한 각각의 스칼라 스레드는 고유한 실행 경로를 따르고 임의의 메모리 위치에 액세스할 수 있지만, 일반적인 경우에는 모두 소수의 실행 경로를 따르고 유사한 작업을 수행합니다.
대부분의 현대 GPU에서는 SIMT 실행 모델을 통해 GPU 스레드 간의 수렴 제어 흐름을 활용합니다. 여기서 스칼라 스레드는 SIMD 하드웨어에서 실행되는 워프로 그룹화됩니다(3.1.1절 참조). 이 장에서는 스칼라화와 어파인 실행을 통해 이러한 스칼라 스레드의 유사성을 더 잘 활용하는 일련의 연구를 요약합니다.
이 연구의 핵심 통찰은 동일한 계산 커널을 실행하는 스레드 간의 값 구조[Kim et al., 2013]를 관찰한 것입니다. 예제 11.1의 계산 커널에 나타난 두 가지 유형의 값 구조, 균일 및 어파인은 다음과 같습니다.
균일 변수
모든 스레드에서 동일한 상수 값을 가지는 변수. 알고리즘 11.1에서 변수 a, THRESHOLD 및 Y_MAX_VALUE는 모두 계산 커널의 모든 스레드에서 균일한 값을 가집니다. 균일 변수는 단일 스칼라 레지스터에 저장되어 계산 커널의 모든 스레드에서 재사용될 수 있습니다.
어파인 변수
계산 커널의 모든 스레드에서 스레드 ID의 선형 함수인 값을 가지는 변수. 알고리즘 11.1에서 y[idx] 변수의 메모리 주소는 스레드 ID threadIdx.x의 어파인 변환으로 표현될 수 있습니다:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;이러한 어파인 표현은 베이스와 스트라이드의 두 개의 스칼라 값으로 저장될 수 있으며, 완전히 확장된 벡터보다 훨씬 더 compact합니다.
__global__ void vsadd( int y[], int a ) {
// 코드 내용은 번역하지 않습니다.
}
```여기는 한국어 번역본입니다:
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
알고리즘 11.1: GPU 컴퓨팅 커널에서의 스칼라 및 어파인 연산의 예시 (Kim et al., 2013에서 발췌).
GPU에서 균일 또는 어파인 변수를 탐지하고 활용하는 다양한 연구 제안들이 있습니다. 이 장에서는 이러한 제안들을 두 가지 측면에서 요약합니다.
## 균일 또는 어파인 변수의 탐지
GPU 컴퓨팅 커널에서 균일 또는 어파인 변수의 존재를 탐지하는 두 가지 주요 접근법은 컴파일러 기반 탐지와 하드웨어 기반 탐지입니다.
### 컴파일러 기반 탐지
GPU 컴퓨팅 커널에서 균일 또는 어파인 변수의 존재를 탐지하는 한 가지 방법은 특별한 컴파일러 분석을 통해 수행하는 것입니다. 이것이 가능한 이유는 기존의 GPU 프로그래밍 모델인 CUDA와 OpenCL이 이미 프로그래머가 변수를 전체 컴퓨팅 커널에 걸쳐 상수로 선언할 수 있는 수단을 제공하고, 스레드 ID를 위한 특별한 변수도 제공하기 때문입니다. 컴파일러는 제어 종속성 분석을 수행하여 상수와 스레드 ID에만 의존하는 변수를 탐지하고, 이를 균일/어파인 변수로 표시할 수 있습니다. 이렇게 탐지된 균일/어파인 변수에 대한 연산은 스칼라화의 대상이 됩니다.
AMD GCN [AMD, 2012]은 컴파일러가 균일 변수와 이에 대한 스칼라 연산을 탐지하여 전용 스칼라 프로세서에서 저장 및 처리하도록 의존합니다.
Asanovic et al. [2013]은 수렴 및 변이 분석을 결합하여 임의의 컴퓨팅 커널에서 스칼라화 및/또는 어파인 변환이 가능한 연산을 컴파일러가 결정할 수 있게 합니다. 컴퓨팅 커널의 수렴 영역 내의 명령어들은 스칼라/어파인 명령어로 변환될 수 있습니다. 수렴 영역에서 발산 영역으로 전환될 때마다, 컴파일러는 두 영역 간의 레지스터 종속성을 처리하기 위해 `syncwarp` 명령어를 삽입합니다. Asanovic et al. [2013]은 이러한 접근법을 채택했습니다.다음은 제공된 마크다운 파일의 한국어 번역입니다. 코드의 경우 코드 자체는 번역하지 않고 주석만 번역했습니다.
Temporal-SIMT 아키텍처 [Keckler et al., 2011, Krashinsky, 2011]를 위한 스칼라 연산을 생성하는 분석을 제공합니다.
Decoupled Affine Computation (DAC) [Wang and Lin, 2017]은 별도의 워프로 분리될 스칼라 및 어파인 후보를 추출하기 위해 유사한 컴파일러 분석에 의존합니다. Wang and Lin [2017]은 분기 어파인 분석을 통해 이 프로세스를 보완하여 명령어 스트랜드를 추출하는 것을 목표로 합니다.