llama.cpp: CDNA 희소 어텐션 지원, 커널 선택 조정

1. CUDA 공유 코드에서 AMD 실행 장벽을 풀었다

이번 변경은 10월 29일 llama.cpp에 병합됐다. CUDA와 AMD GPU용 실행·개발 환경인 HIP가 공유하는 백엔드에서 HIP 빌드의 희소 어텐션 경로를 막던 조건을 제거한 것이 핵심이다. 기존 CUDA 경로의 선택 조건을 유지하면서 AMD CDNA까지 적용 범위를 넓혔다.

이전에는 HIP에서 희소 어텐션 코드가 컴파일 대상에서 빠졌고, 필요한 위치를 모으는 함수도 실행을 중단시켰다. 변경 뒤에는 HIP에서도 이 준비 작업과 희소 계산을 이어서 수행한다. 커널 선택 규칙도 바꿔 CDNA에서 희소 조건을 만족하면 기존의 작업 크기별 선택 규칙보다 희소 행렬 연산 커널을 우선한다.

2. 필요한 위치를 먼저 모으는 비용으로 계산량을 줄인다

희소 어텐션은 계산에 포함할 위치를 나타내는 마스크에서 참조할 수 있는 키 위치를 목록으로 모은 뒤, 선택된 키·값을 처리하는 방식이다. 가려진 위치까지 계산하는 대신 필요한 위치만 골라 처리하므로, 어텐션 연산을 줄이는 이익과 위치 목록을 먼저 만드는 준비 비용을 맞바꾼다.

이번 구현은 GPU에서 함께 움직이는 실행 단위의 크기에 맞춰 위치 선택 결과를 저장한다. CDNA의 GPU가 병렬 계산을 나눠 맡는 64개 실행 레인을 담을 수 있도록 64비트 마스크를 쓰며, 선택된 위치의 수를 세는 처리도 이에 맞췄다.

3. 작업 형태에 맞춰 희소 커널을 고른다

코드 주석에 따르면 CDNA에서 희소 지원을 컴파일하는 대상은 질의·키 차원을 나타내는 DKQ가 576이고, 커널의 열 묶음 설정인 ncols2가 16인 변형이다. 이 변형을 선택할 수 있는지를 먼저 확인한 뒤 실제 작업 조건을 검사한다. RDNA 등 다른 아키텍처의 지원은 작성자가 후속 작업으로 남겼다.

실행 시에는 마스크가 존재하고 계산 결과에 더하는 바이어스와 값 상한 설정이 0인 작업을 희소 경로 후보로 삼는다. CDNA의 키 길이는 최소 4096이면서, 희소 목록에 담을 최대 키 개수의 두 배 이상이어야 한다. 이처럼 소프트웨어가 작업의 크기와 형태를 검사해 실행 경로를 결정한다.

4. 합성 테스트의 처리량과 변동 폭

pp512·tg64는 시험 구분이며, d16384·d32768은 시험 깊이 표기다. 아래 두 열은 원문에 실린 측정 순서를 따른다. 처리량 옆의 ± 값도 원문 그대로다.

표: qwen4exp 합성 테스트의 두 측정 결과

조건: PR 작성자의 합성 테스트. qwen4exp A3B Q4_K - Medium, 103.68 GiB, 176.94 B, ROCm, 50회 반복. 2x mi100s, 1x r9700, 1x rx7900xtx에 분산 로드. 희소 경로 활성화용 별도 패치 포함.

단위: t/s(초당 처리 토큰 수)

첫 번째 측정에서 깊이가 d16384에서 d32768로 늘면 pp512 처리량은 760.52에서 633.18 t/s로, tg64는 35.81에서 30.78 t/s로 낮아진다. pp512의 ± 값도 55.29에서 99.63으로 커져, 이 측정에서는 깊은 시험의 처리량이 낮고 변동 폭은 더 크다.

작성자는 이 테스트에서 성능 개선이 분명하다고 평가했다. 또 모델의 절반가량이 이번 변경의 영향을 받지 않는 아키텍처에 올라갔다는 점을 들어, 측정 결과가 개선 효과를 작게 보여 준다고 주장했다. 성능표에는 적용 전후를 식별하는 표시가 없어 두 열만으로 개선율을 계산하기 어렵다.

5. 이전 ROCm 빌드도 고려했다

ROCm 7.0보다 낮은 버전에서는 레인별 선택 결과를 모으는 __ballot_sync 대신 기존 __ballot을 사용하도록 호환 정의를 추가했다. 작성자는 ROCm 7.0 이전에는 __ballot_sync 내장 기능이 HIP_ENABLE_WARP_SYNC_BUILTINS 뒤에 숨겨져 있었다고 짚었다. 원문에 추가된 코드는 다음과 같다.

#if HIP_VERSION < 70000000
#define __ballot_sync(mask, var) __ballot(var)
#endif //HIP_VERSION < 70000000

작성자는 이 경우 이 커널에서는 __ballot만 쓰면 된다고 설명했다. 다만 이 호환 매크로 자체는 커널별 제한 없이 공유 HIP 헤더 ggml/src/ggml-cuda/vendors/hip.h에 HIP_VERSION < 70000000 조건으로 정의되어 있다. 동시에 다른 파일에서도 __ballot_sync를 쓰고 있는데 그곳에서는 wave64 장치에서 마스크가 실제로 사용되므로, ROCm 6.3-6.4의 wave64에서는 그 사용이 올바르지 않다고 경고했다. 그 사용은 후속으로 처리하겠다고 했지만 해당 후속 커밋은 이 PR에 포함되지 않았으므로, 위의 호환 처리는 남아 있는 이 정확성 문제를 해결하지 못한다.

이번 변경의 핵심은 CDNA에서 필요한 위치만 계산하는 경로를 열고, 작업 조건에 맞춰 그 커널을 선택하도록 연결한 데 있다.