Evaluating Application Characteristics for GPU Portability Layer Selection

본 논문은 주요 고에너지 물리학 실험의 대표적인 이기종 애플리케이션을 분석하여 다양한 GPU 이식성 계층 전반의 성능에 영향을 미치는 핵심 특성을 식별함으로써, 개발자가 자신의 특정 요구 사항에 가장 적합한 기술을 선택할 수 있도록 안내하는 연구를 제시한다.

핵심

당신이 거대한 연회를 요리하려는 셰프라고 상상해 보세요. 당신의 주방에는 세 가지 다른 종류의 고성능 오븐이 있습니다. 하나는 NVIDIA가 만들었고, 하나는 AMD, 다른 하나는 Intel이 만들었습니다. 각 오븐은 음식을 다르게 익히며, 서로 다른 조절 노브를 사용하고, 최상의 결과를 내기 위해 각기 다른 레시피를 요구합니다.

만약 당신이 NVIDIA 오븐에 특화된 레시피(CUDA라는 언어)를 작성한다면, 그 레시피를 AMD나 Intel 오븐에 그대로 넣을 수 없습니다. 레시피 전체를 다시 써야 할 것입니다. 이것은 문제입니다. 왜냐하면 내일 당신의 주방에 어떤 오븐이 들어올지 항상 알 수는 없기 때문입니다.

이 문제를 해결하기 위해, 이 논문은 "포터빌리티 레이어(portability layers, 이식성 계층)"에 대해 논의합니다. 이것들은 만능 번역기 또는 스마트 어댑터라고 생각하면 됩니다. 이들은 하나의 마스터 레시피를 작성하면, 번역기가 이를 각 오븐이 이해할 수 있는 특정 언어로 변환해 줍니다. 이 논문은 여러 종류의 번역기(Kokkos, SYCL, OpenMP, Alpaka)를 살펴보고, 어떤 것이 다양한 요리 작업에 가장 적합한지 테스트합니다.

저자들은 고에너지 물리학 실험(예: 아원자 입자를 연구하는 데 사용되는 실험)의 실제 "레시レシピ(레시피)"를 사용하여 이 번역기들을 테스트했을 때 다음과 같은 결과를 얻었습니다.

1. "스타트업 시간(Startup Time)" 문제

GPU(오븐)를 켜는 것은 즉각적이지 않습니다. 작동 준비를 마치는 데 몇 밀리초가 걸립니다.

• 문제점: 어떤 번역기들은 요리 과정을 시작하는 데 시간이 오래 걸립니다. 예를 들어, Kokkos는 AMD 오븐을 사용할 때 상당한 지연을 발생시킬 수 있습니다. 만약 당신의 요리 작업이 매우 짧다면(예: 10초 동안 달걀 삶기), 번able기(번역기)가 가스레인지를 켜는 데만 5초가 걸린다면, 당신은 시간의 절반을 낭비하게 됩니다.
• 교훈: 작업이 아주 작고 빠르다면, 시작 속도를 느리게 만드는 번역기는 피하십시오.

2. "붐비는 주방(Crowded Kitchen)" 문제

실제 물리학 실험실에서 GPU는 혼자 일하지 않습니다. GPU는 많은 사람(스레드)이 동시에 오븐을 사용하려고 하는 더 큰 시스템의 일부입니다.

• 문제점: 어떤 번역기들은 인파를 처리하는 데 서툽니다. 예를 들어, Kokkos는 "오직 한 사람만이 한 번에 오븐에 말을 걸 수 있다"라는 규칙을 가지고 있어, 여러 명의 셰프가 동시에 작업을 실행하려고 하면 교통 체증이 발생합니다. SYCL은 일관성이 부족합니다. 사용하는 번역기의 버전에 따라 모두가 동시에 요리할 수 있게 해주기도 하고, 때로는 줄을 서서 기다리게 만들기도 합니다.
• 교훈: 애플리케이션에 많은 사람이 동시에 작업해야 한다면, 문을 잠가버리지 않고도 붐비는 주방을 관리할 줄 아는 번역기가 필요합니다.

3. "도구 상자 호환성(Toolbox Compatibility)" 문제

물리학 레시피는 종종 수학과 데이터 처리를 돕는 특별한 도구들(ROOT나 Eigen 같은 라이브러리)을 사용합니다.

• 문제점: 일부 도구들은 번역기와 잘 어울리지 않습니다. 예를 들어, 인기 있는 수학 도구인 Eigen은 많은 번역기가 의존하는 NVIDIA 컴파일러와 함께 사용할 때 자주 깨집니다. 또한, CPU용 컴파일러와 GPU용 컴파일러라는 두 가지 서로 다른 컴파일러를 한 프로젝트에서 사용하는 것은 마치 서로 맞지 않는 두 세트의 설계도로 집을 짓는 것과 같아서, 소프트웨어를 구축(빌드)하는 데 악몽이 됩니다.
• 교훈: 번역기를 선택하기 전에, 당신이 즐겨 쓰는 도구들이 그 안에 들어갈 수 있는지 확인하십시오.

4. "가구 배치(Furniture Arrangement)" 문제

GPU는 단순하고 평평한 구조를 좋아합니다. 데이터가 깔끔한 상자 열처럼 배열되어 있는 것을 선호합니다. 하지만 물리학 데이터는 종종 복잡하고 무질서한 형태(예: 크기가 제각각인 여행 가방 더미)로 옵니다.

• 문제점: 번역기들은 이 무질서를 해결하기 위해 데이터를 특별한 컨테이너로 감싸려고 시도합니다. 비록 이것이 코드를 이식 가능하게 만들어 주기는 하지만, "오버헤드(overhead)"를 발생시킵니다. 이는 마치 양말 한 짝만 옮기면 되는데도 모든 물건을 가방에 담아 옮기는 것과 같으며, 속도를 늦춥니다. 또한, 어떤 번역기도 물리학에서 매우 흔하게 나타나는 "들쭉날쭉한(jagged)" 데이터(길이가 서로 다른 행)를 처리하는 데 능숙하지 않습니다 않습니다.
• 교훈: 데이터가 복잡하고 무질서하다면, 번역기가 이를 정리하려다 속도를 늦출 수 있습니다.

5. "특화된 도구(Specialized Tools)" 문제

때때로 난수 생성기(RNG)나 고속 푸리에 변환(FFT)과 같은 특정 도구가 필요합니다.

• 문제점: 각 오븐 제조사는 자신들만의 초고속 특화 도구를 가지고 있습니다. 범용 번역기들은 이러한 특화된 도구들을 포함하지 않거나, 자체적인 더 느린 버전의 도구를 사용합니다. 물론 번역기가 오븐의 네이티브 도구를 강제로 사용하도록 할 수도 있지만, 그렇게 하면 그 도구가 특정 오븐에서만 작동하기 때문에 "이식성(portability)"이 깨지게 됩니다.
• 교훈: 특정 도구에 크게 의존한다면, 속도(오븐의 네이티브 도구 사용)와 이식성(번역기의 일반적인 도구 사용) 사이에서 하나를 선택해야 할 수도 있습니다.

6. "건설 시간(Construction Time)" 및 "이사 날(Moving Day)" 문제

• 레시피 만들기: 어떤 번역기들은 "요리 시간(컴파일 시간)"을 훨씬 길게 만듭니다. 거대한 프로젝트의 경우, 특정 번역기를 사용하면 빌드 과정이 몇 분 대신 몇 시간씩 걸릴 수 있습니다.
• 주방 옮기기: 만약 당신이 특정 오븐(예: NVIDIA V100)을 위해 소프트웨어를 빌드했다면, 그것이 더 새로운 오븐(예: NVIDIA A100)에서는 작동하지 않을 수 있습니다. 어떤 번역기들은 당신이 마주칠 수 있는 모든 유형의 오븐에 대해 별도의 버전을 빌드할 것을 요구합니다. 이는 소프트웨어를 여러 연구소에 배포하는 데 있어 엄청난 물류적 어려움을 초래합니다.

최종 결론

논문은 **"완벽한 번역기는 없다"**고 결론짓습니다.

• Kokkos는 어떤 것들에는 훌륭하지만, 특정 하드웨어에서 동시성(concurrency)과 스타트업 시간 문제로 어려움을 겪습니다.
• SYCL은 강력하지만 컴파일러 버전에 따라 일관성이 없을 수 있습니다.
• OpenMP 등은 메모리 처리와 서로 다른 하드웨어를 다루는 방식에 있어 각자의 강점과 약점을 가지고 있습니다.

핵-심 요약: 단순히 유명하다고 해서 번역기를 선택해서는 안 됩니다. 당신의 구체적인 "레시피(애플리케이션)"를 살펴봐야 합니다. 만약 코드가 짧고 빠르다면, 스타트업 시간이 낮은 번역기를 선택하십시오. 만약 코드가 복잡하고 많은 도구를 사용한다면, 그 도구들과 잘 어울리는 것을 선택하십시오.

저자들은 이러한 기술들이 매년 새로운 모델의 오븐이 출시되는 것처럼 빠르게 진화하고 있다는 점을 언급합니다. 오늘 가장 좋은 것이 내일은 바뀔 수 있으므로, 개발자들은 계속해서 변화하는 환경을 주시해야 합니다. 미래에는 새로운 표준들이 이러한 선택을 더 쉽게 만들어 줄 수도 있지만, 현재로서는 철저한 테스트만이 적합한 것을 찾는 유일한 방법입니다.

기술 요약

기술 요약: GPU 이식성 계층 선택을 위한 애플리케이션 특성 평가

문제 정의

고에너지 물리학(HEP) 컴퓨팅은 성능 향상을 위해 GPU에 점점 더 의존하고 있으나, 하드웨어 지형은 NVIDIA를 넘어 AMD와 Intel을 포함하여 다양해지고 있다. CUDA와 같은 벤더 전용 언어는 높은 성능을 제공하지만, 특정 하드웨어로의 배포를 제한하여 이식성을 떨어뜨린다. 반면, Kokkos, SYCL, HIP, OpenMP, Alpaka, std::par와 같은 이식성 계층은 다양한 CPU 및 GPU 아키텍처에서의 실행을 약속한다. 그러나 이러한 계층들은 기능적으로 동일하지 않으며, 성능, 병렬성 지원, 그리고 기존 코드베이스와의 호환성 측면에서 뚜렷한 특성을 보인다. 핵심적인 문제는 개발자가 애플리케이션의 특정 특성, 커널 구조 및 런타임 환경에 따라 가장 적절한 이식성 계층을 선택할 수 있는 명확한 프레임워크가 부족하다는 점이다.

방법론

고에너지 물리학 계산 우수성 센터(HEP-CCE)는 주요 HEP 실험의 대표적인 이종 컴퓨팅 애플리케이션을 사용하여 비교 연구를 수행하였다:

• ATLAS: FastCaloSim (매개변수화된 칼로리미터 시뮬레이션).
• CMS: Patatrack (40개 이상의 커널, 멀티스레딩, 메모리 풀을 포함하는 이종 픽셀 재구성) 및 P2R (단일 커널 트랙 재구성).
• DUNE: Wire-Cell Toolkit (래스터화, scatter-add 및 FFT 컨볼루션 커널).
• 이벤트 생성기: Sherpa(Pepper) 및 Madgraph(자동 생성된 코드를 활용함)의 GPU 포트.

이 테스트베드들은 코드 특성이 성능과 호환성에 미치는 영향을 식별하기 위해 다양한 이식성 기술으로 포팅되었다. 본 연구는 커널 실행 지연 시간(kernel launch latency), 병렬성 처리, 라이브러리 호환성, 메모리 관리, 컴파일 오버헤드를 포함한 다차원적 측면에서 이 애플리케이션들을 평가하였다.

주요 결과 및 발견

1. 커널 실행 지연 시간 (Kernel Launch Latency)

커널 실행은 즉각적이지 않으며 마이크로초에서 수십 마이크로초 범위에 걸쳐 발생할 수 있다.

• Kokkos: 특히 NVIDIA나 Intel보다 높은 네이티브 실행 지연 시간을 가진 AMD GPU에서 수십 마이크로초의 상당한 오버헤드가 관찰되었다.
• SYCL, HIP, OpenMP, Alpaka: 네이티브 구현과 비교했을 때 큰 추가 페널티를 보이지 않았다.
• AMD 특이사항: AMD GPU에서 rocRAND를 호출할 때 극심한 첫 호출 지연 시간(1초 초과)이 발생하였으며, 이는 이식성 계층이 실행 오버헤드를 악화시킬 경우 짧게 실행되는 커널들이 심각한 성능 저하를 겪을 수 있음을 시사한다.

2. 병렬성 및 스레드 풀 (Concurrency and Thread Pools)

HEP 애플리케이션은 종종 멀티스레딩(예: Intel TBB)을 사용하는 더 큰 프레임워크 내에서 실행된다.

• Kokkos: 외부 스레드 풀과의 심각한 불호환성을 보였다. 이의 시리얼 백엔드는 호출을 직렬화하는 락(lock)을 사용하며, 스레드 백엔드는 외부 스레드 호출을 금지한다. 현재 동시 커널 실행은 아키텍처별 API를 사용하는 CUDA/HIP 백엔드에서만 가능하여 이식성을 제한한다.
• SYCL: 구현이 일관되지 않았다. 컴파일러 제조사와 버전에 따라 서로 다른 스레드로부터의 동시 호출을 직렬화하는 경우가 있었다. 연구 당시 AMD 아키텍처는 동시성을 지원하는 SYCL 솔루션이 부족한 것으로 나타났다.

3. 외부 라이브러리 및 컴파일러 호환성

• Eigen: 많은 버전이 NVIDIA의 nvcc와 호환되지 않아, nvcc에 의존하는 이식성 계층(Kokkos, Alpaka)의 이식성에 영향을 미친다. 또한 현재 OpenMP 오프로드와도 호환되지 않는다.
• ROOT: NVIDIA의 nvc++로 완전히 컴파일할 수 없다. 하이브리드 컴파일(g++와 nvc++ 혼용)은 가능하지만 빌드 규칙을 복잡하게 만들며, 세심한 메모리 할당 관리가 필요하다(GPU 데이터는 nvc++에 의해 할당되어야 함).
• 빌드 복잡성: 한 프로젝트에서 여러 이식성 계층을 지원하려면 레이어 간의 헤더 충돌(예: Kokkos와 CUDA 사이)을 피하기 위해 복잡한 빌드 규칙(예: #ifdef 사용, 파일 분리)이 필요하다.

4. 데이터 구조 및 메모리 전송

• 객체 지향형 vs 평탄한 레이아웃 (Object-Oriented vs. Flat Layouts): HEP에서 흔히 볼 수 있는 복잡한 객체 지향 데이터 구조는 GPU에 잘 맞지 않으며, GPU는 평탄한 레이아웃이나 구조체 배열(SoA)을 선호한다.
• 이식성 오버헤드: Kokkos 및 SYCL과 같은 계층은 추상화(예: Kokkos::Views, sycl::buffers)를 제공하는데, 이는 특히 작은 객체에 대해 할당 및 전송 오버헤드를 증가시킨다.
• 초기화: Kokkos::Views의 기본 메모리 초기화는 성능 병목 현상이 될 수 있으며, 이는 Kokkos::ViewAllocateWithoutInitializing을 사용하여 완화할 수 있다.
• 통합 공유 메모리 (USM): std::par에서 사용되는 자동 USM 전송은 명시적 전송보다 느린 것으로 관찰되었으며, 최적의 시점에 트리거되지 않을 수 있다.
• 자개 배열 (Jagged Arrays): 어떤 API도 가변 크기의 다차원 벡터(자개 배열)를 네이티브로 지원하지 않으므로, 수동으로 제작하거나 차원 축소 또는 패딩이 필요하다.

5. RNG, FFT 및 원자 연산 (Atomics)

• 네이티브 라이브러리: 벤더 전용 라이브러리(cuRAND, rocRAND, oneMKL)는 튜닝된 성능을 제공하지만 이식성이 부족하다. Kokkos는 교차 아키텍처용 RNG/FFT 구현을 제공하지만, 네이티브 성능에는 미치지 못할 수 있다.
• HEP-CCE 기여: 그룹은 단일 API와 하드웨어 네이티브 백엔드를 갖춘 헤더 전용 라이브러리를 개발하여, SYCL이 여러 벤더에 걸쳐 네이티브 구현을 호출할 수 있도록 하였다.
• **원자 연산 (Atomics):적으로 매핑이 잘 되지 않는다. 성능은 컴파일러와 하드웨어에 따라 크게 달라진다 (예: OpenMP/clang13 환경의 NVIDIA에서 atomic<int>는 CUDA보다 35배 느렸다). 정수형 원자 연산은 일반적으로 부동 소수점/실수형 원자 연산보다 성능이 낮다.

6. 컴파일 시간

• 컴파일러 오버헤드: nvc++는 일반적으로 g++보다 23배 느리다. Kokkos는 10%20%의 오버헤드를 추가한다.
• 템플릿 중심 코드: 대규모 커널에 대해 템플릿이 많이 사용된 자동 생성 코드를 사용하는 Madgraph는 SYCL의 icpx 백엔드에서 다른 백엔드보다 1000배 이상 느린 컴파일 시간을 경험했다.

7. 런타임 프로비저닝

• Kokkos: 특정 하드웨어 식별자와 함께 빌드해야 한다. 하나의 GPU(예: V100)를 위해 빌드된 바이너리는 다른 GPU(예: A100)에서 실행되지 않을 수 있다. 공유 라이브러리와 여러 아키텍처를 지원하는 것은 복잡한 빌드 구성 매트릭스를 만들어 바이너리 배포를 어렵게 만든다.
• Alpaka: 각 백엔드 컴파일러에 대한 별도의 컴파일이 필요하다.
• SYCL/OneAPI: 다중 플랫폼 바이너리가 등장하고 있으나, 특정 아키텍처로 제한되거나 컴파일 타임 식별을 요구하는 경우가 많다.
• OpenMP: 일부 구현(LLVM Clang)은 컴파일 타임에 여러 아키텍처를 나열하고 런타임에 백엔드를 동적으로 선택하는 것을 허용한다.

의의 및 결론

본 논문은 모든 시나리오에 최적인 단일 GPU 이식성 솔루션은 존재하지 않는다고 결론짓는다. 이식성 계층의 선택은 커널 실행 시간, 병렬성 요구 사항, 데이터 구조 복잡성 및 외부 라이브러리 의존성을 포함한 애플리케이션의 특정 특성에 대한 면밀한 분석에 의해 결정되어야 한다.

저자들은 생태계가 빠르게 진화하고 있음을 강조한다. 이식성 계층(Kokkos, SYCL 등)과 네이티브 컴파일러(clang, nvcc 등) 모두 기능 지원과 최적화된 성능을 향상시키는 방향으로 발전하고 있다. 본 논문은 코드베이스와 이식성 계층 간의 최적의 매치를 보장하기 위해 이러한 기술들을 빈번하게 모니터링할 것을 권고한다. 또한, 신흥 C++ 표준(예: 비동기 실행을 위한 P2300)이 향로 이러한 선택을 단순화할 수 있으나, 광범위한 하드웨어 지원까지는 수년이 걸릴 것이라고 언급하였다.

궁극적으로, 코드의 본질, 외부 의존성 및 런타임 환경을 완전히 이해하지 못한 채 이식성 계층을 선택하는 것은 성능과 배포 유연성 측면에서 심각한 문제를 야기할 위험이 있다.