ROCm 패키지를 빌드하고 유지보수하면서 마주친 문제들.
ROCm 패키지를 빌드하고 유지보수하면서 마주친 문제들 2026-02-22 작성자 Luna Nova
2026-02-22 작성자 Luna Nova
ROC m
:7.1.1
빌드할 수 있다(없다).
ROCmacronymnt ROCm… 이게 Radeon Open Co m pute의 약자일까? Compute modules? Compute ecosysteM? 아니다! 알고 보니 Radeon Open Compute Platform이다. 하지만 P는 m보다 어감이 별로였다고 한다. AMD의 공식 입장은 "ROCm no longer functions as an acronym."이다. "더 이상". 그러니까. 누군가 m이 어떻게 Platform을 뜻하냐고 묻기 전까지는 약자로 기능했지만, 이제 그 괴리가 눈에 띄었으니 앞으로는 기능할 수 없다는 뜻이다. 마치 Wile E. Coyote식 약자학이다.
은 GPU 가속 컴퓨팅을 위한 AMD의 오픈 소스 생태계다. ROCm은 blender 레이 트레이싱과 머신 러닝을 포함한 여러 작업에서 AMD에 매우 중요하다. 운이 좋다면 대신 vulkan을 쓸 수도 있다.
나는 1년 정도 nixpkgs에서 ROCm 패키지 집합을 유지보수해 왔다.
지원되는 모든 그래픽스 타깃용 ROCm 패키지를 전부 빌드하거나, 아니면 어쩌면 하나만 빌드할 당신의 앞으로의 경험을 위해 도움이 될 만한 요점을 정리하면 다음과 같다:
CPU 코어가 충분하지 않을 것이다
Zen 5의 32 스레드로는 쾌적한 경험이 불가능하다
EPYC Milan의 128t로도 쾌적한 경험이 불가능하다
5.6GHz로 오버클럭된 EPYC 엔지니어링 샘플의 256t로도 쾌적한 경험이 불가능하다ck
composable_kernel은 겉으로는 최적화된 커널을 제공하는 머신 러닝 라이브러리다. 실제로는 컴파일러 고문 테스트에 가깝다. 템플릿 인스턴스화가 너무 복잡해서 커널 하나 device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance(수천 개 중 하나)를 컴파일하는 데 --offload-arch마다 clang++ 시간이 15분씩 걸린다. 업스트림은 우발적으로 만들어진 엄청나게 뒤엉킨 템플릿 텐서 타일 고문 테스트를 좀 더 현대적이고 빠른 C++ 사용으로 정리하면서 이 문제를 해결하기 시작하고 있다.
전화번호처럼 보이는 load average를 각오하라
RAM도 충분하지 않을 것이다
32GiB는 불가능하다hipblaslt-ram hipBLASlt의 Tensile 커널 생성기는 OOM이 날 때까지 RAM을 먹어치운다. 정확한 한계는 불분명하며 시간이 지나며 달라져 왔지만, 최악의 경우 supermassive RAM usage with TensileCreateLibrary #316에서 96GiB도 부족하다고 기록되어 있다. 단일 ISA용으로 빌드하고 그 ISA가 gfx942가 아니라면 32GB 미만으로도 아마 어떻게든 될 수 있다. 아마도. 상황이 더 나빠지지 않았다면.
96GiB로도 쾌적한 경험에는 부족하다
512GiB여도 전체 병렬화를 허용하기엔 여전히 부족하다
계속하려면 $1,000ram-shortage를 넣어 달라
주요 다른 패키지들의 필수 링크 의존성인 패키지들이 빌드 시점에 당신의 GPU 모델을 지원하지 않는다pytorch-hipblaslt pytorch는 hipblaslt와의 링크에 강하게 의존하지만, hipblaslt는 gfx1030이나 gfx90c 같은 아키텍처용으로는 빌드할 수 없다. 이것 때문에 많은 문제가 발생했다. ROCm loses some supported GPUs by requiring hipblaslt #119081을 보라. 사용자와 배포판은 디바이스 커널이 없는 스텁 빌드를 생성할 수 있도록 hipblaslt를 패치하거나, 런타임에는 사용되지 않을 여분의 버리는 아키텍처 하나를 추가로 빌드하는 방법에 의존해 왔다.
GPU_TARGETS를 아무 지원 모델로나 설정해도 무시된다.AMDGPU_TARGETS 또는 GPU_ARCHS 또는 SUPPORTED_TARGETS 또는 CMAKE_HIP_ARCHITECTURES 또는 PYTORCH_ROCM_ARCH 또는 AOTRITON_TARGET_ARCH 또는 HIP_ARCHITECTURES를 원했다패키지 출력물이 배포판 인프라에 비해 너무 크다hipblaslt-metadata nixpkgs가 들고 가는 패치가 없으면 hipBLASLt의 Tensile 커널 메타데이터 파일(.dat)은 장황한 형식으로 저장되어, 일부 ISA의 lazy library 하나당 80MiB+에 달하고 전체적으로는 10GiB를 넘었다. 그 패치는 msgpack 형식 메타데이터에 zstd 압축을 가능하게 해서 파일 크기를 극적으로 줄인다. rocm-libraries issue #1327와 ROCm 6.4부터 적용된 nixpkgs의 messagepack-compression-support.patch를 보라. 소스 tarball 자체도 7.2 이후에는 번들된 YAML 파일을 fetch 후 압축해야 하는데, CDNA4/gfx950용으로 추가된 2GiB 때문에 한계를 넘어섰기 때문이다. nixpkgs PR #497589을 보라.
중요한 패키지 하나가 사실은 컴파일러 템플릿 고문 테스트이면서 동시에 머신 러닝 라이브러리라는 사실을 알게 된다ck
composable_kernel은 겉으로는 최적화된 커널을 제공하는 머신 러닝 라이브러리다. 실제로는 컴파일러 고문 테스트에 가깝다. 템플릿 인스턴스화가 너무 복잡해서 커널 하나 device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance(수천 개 중 하나)를 컴파일하는 데 --offload-arch마다 clang++ 시간이 15분씩 걸린다. 업스트림은 우발적으로 만들어진 엄청나게 뒤엉킨 템플릿 텐서 타일 고문 테스트를 좀 더 현대적이고 빠른 C++ 사용으로 정리하면서 이 문제를 해결하기 시작하고 있다.
핵심 행렬 곱셈 패키지 하나가 빌드 중간에 200GiB의 어셈블리 파일을 내뿜어 공간이 바닥난다hipblaslt-disk-usagehipBLASLt의 Tensile 커널 생성기는 임시 디스크 공간 사용량이 파국적이다. 빌드 과정에서 어셈블리 파일에 엄청나게 장황한 주석을 뿜어낸다. 예를 들면 // 1 wait state required when next inst writes vgprs held by previous dwordx4 store inst 같은 친절한 설명이 반복된다. ROCm 6.4.3 기준으로 빌드는 모든 어셈블리 파일을 코드 오브젝트로 변환하기 전에 한꺼번에 실체화하며, 이 과정에서 임시 공간 240GB를 소비한다. 그중 ≈90GB는 주석 스팸만으로 차지한다.
nixpkgs PR #449985에서의 조사로 장황한 주석을 억제해 피크 사용량을 150GB로 줄였다. 후속 작업(PR #451188)은 Tensile에 대한 큰 리팩터링을 적용했다. 모든 어셈블리 파일을 한 번에 실체화하는 대신, 즉시 코드 오브젝트로 변환하고 unlink하는 방식이다. 이로써 빌드 피크 공간은 25GB까지 내려가 원래의 240GB에서 90% 감소했다. 이런 개선 이후에도 hipblaslt는 빌드에 여전히 25GB+의 임시 공간을 요구한다.
이 패치들은 아직 병합되지 않았고, hipblaslt의 전반적 상태 때문에 끝없이 draft 상태에 머물러 있다. 업스트림은 최근 어셈블리는 unlink하지만 오브젝트 파일은 그대로 두는 더 작은 변경을 병합했다. 약간이나마 도움이 된다. 이것은 ROCm 7.2.0에 들어왔다. * 어셈블리 주석 속에서 『전쟁과 평화』를 25,000번쯤 발견하게 된다war-and-peace-math War and Peace 평문은 대략 3.7 MiB다. hipBLASLt의 빌드 출력물(모든 ISA 전체)에 포함된 ≈90 GiB의 어셈블리 주석은 대략 25,000권 분량에 해당한다.
ROCm의 LLVM 포크가 x86 호스트 AVX512 코드를 컴파일하다 무한 루프에 빠진다
rocmSupport를 켠다. whisper-cpp 빌드의 CPU 부분이 10시간 넘게 멈춘다llvm-infinite-loop
ROCm의 pre-release LLVM rev 기반 LLVM 22 포크에는 v64i8 ↔ v32i16 lowering이 무한 루프에 빠지는 x86 vector shuffle legalization 버그가 있었다. 문제의 원인은 ggml-cpu/amx/mmq.cpp의 AVX-512 경로가 lowerShuffleAsLanePermuteAndPermute 안의 무한 루프를 밟았기 때문이다. 업스트림 수정 사항은 nixpkgs PR #497818에서 백포트되었다.누군가 Gentoo용 ROCm 패키지를 고치느라 Google Summer of Code 프로젝트 하나를 통째로 쓴다gsoc-gentoo 12주짜리 Google Summer of Code 프로젝트 — "Refining ROCm Packages in Gentoo" — 는 ROCm 관련 문제를 많이 기록해 두었다. 2주차 보고서는 순환 의존성과 관련한 많은 문제를 되짚었다. 개발자는 "battling against hip"라는, rebase되지 않은 git 히스토리를 경고담처럼 브랜치에 보존해 두었다. 프로젝트 요약에 따르면 원래 제안서는 ROCm을 vanilla clang으로 포팅하는 데 1주를 배정했다. 5주차 보고서는 새 문제가 계속 나타나 TensorFlow와 JAX 지원을 포기해야 했다고 기록한다.
80개가 넘는 패키지 집합 안에서 많은 순환 의존성을 발견한다circular aotriton ↔ torch. aqlprofile ↔ rocm-runtime. rocprofiler-register ↔ clr.
주요 라이브러리의 커널이 git에 보이지 않는다. 한참 혼란스러워하다가 업스트림의 lfsconfig가 모든 파일을 제외하고 있다는 걸 깨닫는다miopen-fetch 'FIXME: 여기의 끔찍함 수준을 누가 좀 낮출 수 있다면 정말 좋겠다.' remote를 다시 추가하고, 태그를 fetch하고, clean하고, 브랜치를 전환하고, 망가진 config를 제거하고, *.kdb.bz2를 다시 추적하고, LFS 오브젝트를 fetch 및 pull한 다음, .git을 삭제하는 데 20줄이 필요하다.
AArch64용으로 빌드한다. 일부 패키지는 x64 전용 플래그를 설정한다rdc
rdc는 CMakeLists에서 -m64 -msse -msse2를 설정한다. 이유는 불분명하다. 이를 고치기 위한 PR이 project() 이전 순서로 인한 플래그 덮어쓰기 문제와 함께 제출되었다.
이유를 도무지 알 수 없는데도 말이다.
GPU 두 개를 사용하려고 한다. 핵심 집단 통신 라이브러리에 최적화 수준을 낮춰야 하는 정의되지 않은 동작이 들어 있다rccl-ub
RCCL(AMD의 NCCL 대응물)은 업스트림 코드의 UB 때문에 nixpkgs에서는 -O2 -fno-strict-aliasing로 빌드된다. 이 UB는 런타임에서 무너지는 원인이 되었다.
업스트림 빌드 과정을 개선하려고 시도한다
draft PR 수백 줄쯤 들어간 시점에서 재작성 없이는 안전하게 진행할 방법이 없다는 결론에 이른다hipblaslt-draft nixpkgs는 Tensile의 Python 프로세스 메모리 및 디스크 사용을 최적화하는 draft PR의 일부를 vendor한다. 왜 이것이 안전하지 않은지에 대해서는 논의 내용을 보라. 기존 코드는 mutation과 방어적 cloning에 의존하고 있어서, 메모리 사용량을 줄이기 위해 deep clone을 제거하는 일은 위험하다. 이 타입이 없는 Python 코드베이스에서 변경이 올바른지 검사할 수 있는 정적 분석 도구는 없다.
전체 빌드 시스템을 rust로 다시 쓸까 고민한다
전체 빌드 시스템을 rust로 다시 쓴다
rocmPackages를 건드리는 PR에 대해 nixpkgs-review 실행을 트리거한다.
업스트림도 문제를 해결하려고 애쓰고 있다. 언젠가는 거기에 도달하길 바란다.
생태계와 여러 배포판 전반에서 도움을 주고 있는 모든 분들께, 순서는 특별히 없지만 큰 감사를 전한다:
@bstefanuk — AMD에서 ROCm 전반의 CMake 설정을 개선하기 위한 대규모 리팩터링을 진행.
@stellaraccident — AMD에서 새로운 kpack 형식으로 호스트 코드와 커널을 분리하면서, 패키지를 ISA별로 더 원칙적으로 쪼개려 노력 중.
@GZGavinZhao — Solus/NixOS에서 1031이 1030을 로드하는 식으로 충분히 비슷한 ISA끼리 로드되도록 해, 지원되지 않는 ISA에도 최선의 지원을 제공하게 하는 패치 집합을 유지보수하는 등 Solus와 NixOS의 ROCm 지원에 큰 기여를 하고 있다.
@AngryLoki — Gentoo에서 꾸준히 패치를 업스트림으로 올리고 있으며, 그중 많은 것들이 모든 배포판에 도움이 된다.
@Madouura — 처음에 ROCm의 상당 부분을 nixpkgs에 들여오는 헤라클레스 같은 작업을 해냈다. 한동안 보이지 않는다. 이 규모와 혼란에 지쳤다고 해도 이해할 만하다.
@Flakebi, @mschwaig — NixOS에서 리뷰(현재까지)와 기여(주로 5.x 및 6.x 시기)에 감사.
@06kellyjac 및 @Wulfsta — NixOS에서 Strix Halo와 Radeon VII 테스트를 도와주었다.
@acowley — Nixpkgs 본격 지원 이전에 존재했던 nixos-rocm overlay 작업에 기여.
ROCm… 이게 Radeon Open Co m pute의 약자일까? Compute modules? Compute ecosysteM? 아니다! 알고 보니 Radeon Open Compute Platform이다. 하지만 P는 m보다 어감이 별로였다고 한다. AMD의 공식 입장은 "ROCm no longer functions as an acronym."이다. "더 이상". 그러니까. 누군가 m이 어떻게 Platform을 뜻하냐고 묻기 전까지는 약자로 기능했지만, 이제 그 괴리가 눈에 띄었으니 앞으로는 기능할 수 없다는 뜻이다. 마치 Wile E. Coyote식 약자학이다.
composable_kernel은 겉으로는 최적화된 커널을 제공하는 머신 러닝 라이브러리다. 실제로는 컴파일러 고문 테스트에 가깝다. 템플릿 인스턴스화가 너무 복잡해서 커널 하나 device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance(수천 개 중 하나)를 컴파일하는 데 --offload-arch마다 clang++ 시간이 15분씩 걸린다. 업스트림은 우발적으로 만들어진 엄청나게 뒤엉킨 템플릿 텐서 타일 고문 테스트를 좀 더 현대적이고 빠른 C++ 사용으로 정리하면서 이 문제를 해결하기 시작하고 있다.
hipBLASlt의 Tensile 커널 생성기는 OOM이 날 때까지 RAM을 먹어치운다. 정확한 한계는 불분명하며 시간이 지나며 달라져 왔지만, 최악의 경우 supermassive RAM usage with TensileCreateLibrary #316에서 96GiB도 부족하다고 기록되어 있다. 단일 ISA용으로 빌드하고 그 ISA가 gfx942가 아니라면 32GB 미만으로도 아마 어떻게든 될 수 있다. 아마도. 상황이 더 나빠지지 않았다면.
pytorch는 hipblaslt와의 링크에 강하게 의존하지만, hipblaslt는 gfx1030이나 gfx90c 같은 아키텍처용으로는 빌드할 수 없다. 이것 때문에 많은 문제가 발생했다. ROCm loses some supported GPUs by requiring hipblaslt #119081을 보라. 사용자와 배포판은 디바이스 커널이 없는 스텁 빌드를 생성할 수 있도록 hipblaslt를 패치하거나, 런타임에는 사용되지 않을 여분의 버리는 아키텍처 하나를 추가로 빌드하는 방법에 의존해 왔다.
hydra의 10시간 빌드 타임아웃에 걸리는 단일 빌드를 피하고, 여러 박스 빌드 팜 전반에서 병렬화할 수 있게 하려고 .o 파일을 nix derivation 사이로 밀반입한다. nixpkgs의 composable_kernel 디렉터리의 default.nix와 base.nix를 보라.
hipBLASLt의 Tensile 커널 생성기는 임시 디스크 공간 사용량이 파국적이다. 빌드 과정에서 어셈블리 파일에 엄청나게 장황한 주석을 뿜어낸다. 반복되는 친절한 주석의 예로 // 1 wait state required when next inst writes vgprs held by previous dwordx4 store inst 같은 것이 있다. ROCm 6.4.3 기준으로 빌드는 모든 어셈블리 파일을 코드 오브젝트로 변환하기 전에 한꺼번에 실체화하며, 이 과정에서 임시 공간 240GB를 소비한다. 그중 ≈90GB는 주석 스팸만으로 차지한다.
nixpkgs PR #449985에서의 조사로 장황한 주석을 억제해 피크 사용량을 150GB로 줄였다. 후속 작업(PR #451188)은 Tensile에 대한 큰 리팩터링을 적용했다. 모든 어셈블리 파일을 한 번에 실체화하는 대신, 즉시 코드 오브젝트로 변환하고 unlink하는 방식이다. 이로써 빌드 피크 공간은 25GB까지 내려가 원래의 240GB에서 90% 감소했다. 이런 개선 이후에도 hipblaslt는 빌드에 여전히 25GB+의 임시 공간을 요구한다.
이 패치들은 아직 병합되지 않았고, hipblaslt의 전반적 상태 때문에 끝없이 draft 상태에 머물러 있다. 업스트림은 최근 어셈블리는 unlink하지만 오브젝트 파일은 그대로 두는 더 작은 변경을 병합했다. 약간이나마 도움이 된다. 이것은 ROCm 7.2.0에 들어왔다.
평문 상태의 War and Peace는 대략 3.7 MiB다. hipBLASLt의 빌드 출력물(모든 ISA 전체)에 포함된 ≈90 GiB의 어셈블리 주석은 대략 25,000권 분량에 해당한다.
12주짜리 Google Summer of Code 프로젝트 — "Refining ROCm Packages in Gentoo" — 는 ROCm 관련 문제를 많이 기록해 두었다. 2주차 보고서는 순환 의존성과 관련한 많은 문제를 되짚었다. 개발자는 "battling against hip"라는, rebase되지 않은 git 히스토리를 경고담처럼 브랜치에 보존해 두었다. 프로젝트 요약에 따르면 원래 제안서는 ROCm을 vanilla clang으로 포팅하는 데 1주를 배정했다. 5주차 보고서는 새 문제가 계속 나타나 TensorFlow와 JAX 지원을 포기해야 했다고 기록한다.
aotriton ↔ torch. aqlprofile ↔ rocm-runtime. rocprofiler-register ↔ clr.
'FIXME: 여기의 끔찍함 수준을 누가 좀 낮출 수 있다면 정말 좋겠다.' remote를 다시 추가하고, 태그를 fetch하고, clean하고, 브랜치를 전환하고, 망가진 config를 제거하고, *.kdb.bz2를 다시 추적하고, LFS 오브젝트를 fetch 및 pull한 다음, .git을 삭제하는 데 20줄이 필요하다.
rdc는 CMakeLists에서 -m64 -msse -msse2를 설정한다. 이유는 불분명하다. 이를 고치기 위한 PR이 project() 이전 순서로 인한 플래그 덮어쓰기 문제와 함께 제출되었다.
nixpkgs는 Tensile의 Python 프로세스 메모리 및 디스크 사용을 최적화하는 draft PR의 일부를 vendor한다. 왜 이것이 안전하지 않은지에 대해서는 논의 내용을 보라. 기존 코드는 mutation과 방어적 cloning에 의존하고 있어서, 메모리 사용량을 줄이기 위해 deep clone을 제거하는 일은 위험하다. 이 타입이 없는 Python 코드베이스에서 변경이 올바른지 검사할 수 있는 정적 분석 도구는 없다.
ROCm의 pre-release LLVM rev 기반 LLVM 22 포크에는 v64i8 ↔ v32i16 lowering이 무한 루프에 빠지는 x86 vector shuffle legalization 버그가 있었다. 문제의 원인은 ggml-cpu/amx/mmq.cpp의 AVX-512 경로가 lowerShuffleAsLanePermuteAndPermute 안의 무한 루프를 밟았기 때문이다. 업스트림 수정 사항은 nixpkgs PR #497818에서 백포트되었다.
BibTeX로 인용
@online{rocm-711-you-can-not-build,
author = {Luna Nova},
title = {ROCm 7.1.1: you can (not) build},
date = {2026-02-22},
year = {2026},
url = {https://lunnova.dev/articles/rocm-711-you-can-not-build/},
}