Portable Acceleration of Learning With Errors KEMs for Post-Quantum Cryptography

기술 요약: 양자 내성 암호(PQC)를 위한 휴대 가능한 학습 오차(LWE) 기반 KEM 가속화

문제 정의
양자 내성 암호(PQC)로의 전환은 실제 응용 프로그램의 계산 요구 사항을 충족할 수 있는 구현 능력을 필요로 합니다. PQC 후보군 중, 학습 오차(Learning With Errors, LWE) 기반 키 캡슐화 메커니즘(KEM)은 강력한 보안 기초를 제공하지만 상당한 계산 오버헤드를 발생시킵니다. 이 오버헤드는 주로 대규모 행렬-벡터 연산과 암호학적으로 안전한 난수 생성에 의해 발생합니다. 그래픽 처리 장치(GPU)는 이러한 병렬 워크로드를 가속화하기 위한 자연스러운 후보이지만, 기존의 GPU 구현은 주로 CUDA 전용 최적화에 의존하여 벤더 종속성(vendor lock-in)을 유발하고, 이기종 가속기 생태계(예: NVIDIA 및 AMD) 간의 휴대성을 제한합니다. 또한, OpenACC를 사용한 이전의 휴대 가능 시도들은 제한적인 교차 벤더 지원을 보여주었습니다. 해결하고자 하는 핵심 과제는, 난수 생성과 같은 핵심 구성 요소를 장치(device)에 상주시켜 유지하면서도, 서로 다른 GPU 아키텍처 전반에 걸쳐 소스 레벨의 휴대성을 유지하며 밀집된 LWE 기반 KEM에 대한 고성능 가속을 달로하는 것입니다.

방법론
저자들은 OpenMP Target 오프로딩을 사용하는 순수 LWE 기반 KEM의 휴대 가능한 GPU 구현을 제시합니다. 방법론은 세 가지 주요 기술 전략을 포함합니다:

1. 휴대 가능한 오프로딩 전략: 구현에는 데이터 이동 및 커널 실행을 관리하기 위해 OpenMP Target 지침을 활용합니다. GPU 상주성을 보장하고 호스트-장치 간 통신을 최소화하기 위해, 저자들은 지속적인 장치 메모리 관리(target enter data map(alloc:...))를 사용합니다. 이를 통해 주요 행렬, 암호문 및 임시 버퍼를 KeyGen, Encapsulation, Decapsulation 단계 동안 GPU에 상주하도록 유지합니다.
2. 이기종 난수 생성(RNG): 워크플로의 핵심 구성 요소는 암호학적으로 안전한 난수 생성입니다. 저자들은 원래 CUDA를 위해 설계된 기존 RNGonGPU 라이브러리를 HIP(Heterogeneous-computing Interface for Portability)을 지원하도록 확장했습니다. HIP 호환 어댑터 레이어를 OpenMP Target 워크플로와 통합함으로써, NVIDIA 및 AMD 가속기 모두에서 완전한 GPU 상주형 난수 생성을 가능하게 했습니다. 이는 #pragma omp target data use_device_ptr를 사용하여 OpenMP가 관리하는 장치 버퍼를 외부 HIP 기반 RNG 백엔드에 노출함으로써 달성되며, 무작위성을 호스트에서 생성하여 장치로 복사할 필요를 방지합니다.
3. 커널 병렬화: 계산 커널은 OpenMP teams distribute parallel for 구문을 사용하여 병렬화됩니다. 키 생성은 비밀 벡터와 오차 벡터의 병렬 생성과 그에 따른 밀집 행렬-벡터 곱을 포함합니다. 암호화 및 복호화는 독립적인 메시지 비트 또는 암호문 슬롯을 동시에 처리하도록 배치(batching)되어 스레드 수준의 병렬성과 버퍼 재사용을 극대화합니다. Fujisaki-Okamoto(FO) 변환의 동일성 검사 또한 전체 과정이 GPU에서 수행되며, 호스트에는 단일 상태 플래그만을 반환합니다.

실험 설정
본 연구는 네 가지 뚜렷한 가속기 아키텍처를 대상으로 구현을 평가했습니다:

• NVIDIA: A100-SXM-64GB (PCIe 4.0, HBM2e) 및 GH200 Grace Hopper (NVLink-C2C, HBM3).
• AMD: MI300X (PCIe 5.0, HBM3) 및 MI300A (Unified APU, 공유 HBM3).

성능은 멀티코어 OpenMP CPU 베이스라인과 비교하여 최대 16,384의 격자 차원($n$)까지 다양한 플랫폼에서 측정되었습니다. 측정 지표로는 실행 시간, 확장성, 에너지 효율(energy-to-solution), 그리고 벤더 도구(Nsight Systems 및 ROCm Systems Profiler)를 이용한 상세 런타임 프로파일링이 포함되었습니다.

주요 결과

• CPU 대비 가속: OpenMP Target 오프로딩은 멀티코어 CPU 베이스라인에 비해 상당한 가속을 입증했습니다. 큰 격자 차원(예: $n=4096$)의 경우, 워크로드가 메모리 집약적(memory-bound)이기 때문에 GPU가 스트리밍 행렬 연산을 위한 우수한 메모리 대역폭과 병렬성을 제공하여 2개 차수 이상의 속도 향상(120배)을 달성했습니다.
• 교차 플랫폼 성능:
  • NVIDIA GH200은 높은 HBM3 대역폭(4 TB/s)과 CPU-GPU 통신 오버헤드를 최소화하는 NVLink-C2C 덕분에 전반적으로 가장 우수한 성능을 달성했습니다.
  • AMD MI300X는 작은 격자 차원에서 최적에 가까운 성능과 강력한 확장성을 보이며 근소한 차이로 뒤를 이었으나, 더 큰 $n$에 대해서는 GH200에 비해 약간 더 긴 실행 시간을 보였습니다.
  • NVIDIA A100은 낮은 HBM2e 대역폭과 PCIe 통신 오버헤드로 인해 GH200 대비 중간 정도의 성능 저하를 보였습니다.
  • AMD MI300A는 가장 약한 확장성과 높은 실행 시간을 나타냈습니다. 프로파일링 결과, 높은 대역폭의 통합 메모리를 보유하고 있음에도 불구하고, 공유 HBM3 서브시스템이 CPU 측 작업(SHA3 해싱, FO 키 유도)과 GPU 커널 사이의 경합을 유발하여, MI300X 대비 유효 대역폭이 약 38% 감소하는 것으로 밝혀졌습니다.
• 에너지 효율성: NVIDIA GH200은 짧은 실행 시간과 낮은 평균 전력 소비에 힘입어, 큰 격자 크기에 대해 AMD MI300X보다 약 2.5배 더 에너지 효율적인 것으로 나타났습니다.
• 휴대성: 단일 소스 코드는 아키텍처별 코드 수정 없이 테스트된 모든 플랫폼에서 성공적으로 실행되었으며, 이는 다양한 하드웨어 환경에서 PQC 가속을 위한 OpenMP Target의 타당성을 입증합니다.

의의 및 주장
본 논문은 휴대 가능한 GPU 가속이 벤더 종속성 없이 PQC의 계산 오버헤드를 줄이기 위한 실행 가능한 전략임을 주장합니다. 연구는 다음을 입증합니다:

1. 소스 레벨의 휴대성 달성 가능: OpenMP Target과 HIP 통합 RNG 백엔드의 결리된 조합을 통해, 단일 코드베이스가 이기종 GPU 생태계를 효과적으로 활용할 수 있습니다.
2. 메모리 아키텍처가 성능을 결정함: 메모리 집약적인 LWE 워크로드의 경우, 피크 계산 능력보다 메모리 대역폭, 메모리 시스템 구조 및 호스트-장치 간 인터커넥트 효율성이 더 중요합니다. 결과는 통합 메모리 아키텍처(MI300A와 같은)가 CPU와 GPU 트래픽이 동일한 메모리 자원을 두고 경합할 경우 자동으로 우수한 성능을 보장하지 않는다는 점을 강조합니다.
3. 실제 배포: CUDA 전용 의존성을 피함으로써, 이 접근 방식은 다양한 하드웨어 환경 전반에 걸쳐 양자 내성 암호 인프라를 배포하는 것을 용이하게 합니다.

저자들은 OpenMP Target이 효과적인 소스 레벨 휴대성을 제공하지만, 런타임 성능은 컴파일러 툴체인의 성숙도, 상호 운용되는 장치 라이브러리의 동작(특히 HIP RNG 어댑터의 동적 메모리 할당 관련), 그리고 대상 하드웨어의 특정 메모리 계층 구조에 따라 달라질 수 있다고 결론지었습니다.