MCP로 살펴본 MI300X — 매트릭스 코어가 실행하는 것과 실행하지 않는 것

AI 요약

이 글은 AMD Developer Cloud에서 시간당 1.99달러로 임대한 AMD Instinct MI300X 가상 함수(VF) 한 장을 대상으로, 클라우드 제어와 GPU 상태 확인을 MCP(Model Context Protocol) 서버에 묶고 실제 하드웨어가 어떤 수치 형식을 실행하는지 측정합니다. 작성 환경에는 AMD GPU가 없으며, 모든 작업은 태그로 범위를 제한한 12개 MCP 도구를 통해 원격 GPU 드롭릿에서 수행합니다. 측정과 하드웨어 확인은 2026년 9월 16일에 진행합니다.

■ MCP 서버가 관리하는 범위와 안전장치

서버는 DigitalOcean API를 사용하는 AMD Developer Cloud 드롭릿을 대상으로 합니다. `list_droplets`, `droplet_status`, `start_droplet`, `stop_droplet`, `reboot_droplet`, `action_status`, `ssh_command`, `run_on_droplet`, `gpu_status`, `hardware_scan`, `list_gpu_sizes`, `get_help`를 제공합니다. 모든 조회는 `tag_name=gemma`로 제한하므로 태그가 없는 드롭릿은 도구에서 보이지 않으며, 다른 대상의 태그가 아닌 한 이름이나 ID만으로 임의의 드롭릿을 조작할 수 없습니다. 계정 토큰 자체에는 전체 권한이 있지만, 태그가 재부팅이나 전원 조작의 경계를 만듭니다.

생성(create)과 삭제(destroy) 도구는 의도적으로 제공하지 않습니다. 생성은 시간당 과금을 시작하고 삭제는 상태를 없애므로, 모델의 단일 도구 호출에 달러 단위의 결정을 맡기지 않고 DigitalOcean 콘솔에서 직접 수행하도록 합니다. 다만 드롭릿을 중지해도 과금은 멈추지 않으며, 리소스가 예약된 상태로 유지되어 전체 요금이 계속 부과됩니다. 실제 과금을 멈추려면 삭제해야 하며, `stop_droplet`은 이 사실을 반환 메시지에 매번 포함합니다.

서버는 하드코딩한 IP를 저장하지 않고 각 호출마다 현재 주소를 다시 확인합니다. `hardware_scan`은 여러 SSH 연결을 반복하는 대신 한 번의 연결에서 호스트, GPU, 펌웨어, ROCm 패키지, 설치 도구를 수집하고, 마커로 구분한 셸 출력물을 파싱합니다. 결과는 20 vCPU Intel Xeon Platinum 8568Y+, 사용 가능한 RAM 236GB, 720GB 디스크, Debian 13과 커널 6.12.94를 보여줍니다. GPU는 `gfx942` 아키텍처, 304개 Compute Unit(CU), 4개 SIMD/CU, 64-wide wavefront, 최대 2100MHz 클록, PCIe Gen5 x16 구성입니다.

MI300X의 VRAM은 205,822,885,888바이트, 즉 191.69GiB이며 그중 168.31GiB가 사용 중입니다. `VIS_VRAM`과 `VRAM`이 동일하므로 전체 프레임버퍼가 CPU에서 매핑 가능한 large-BAR 구성입니다. `rocminfo`가 보고하는 coarse-grained, fine-grained, extended fine-grained 세 GLOBAL 메모리 풀도 서로 다른 세 할당이 아니라 동일한 191.69GiB를 세 방식으로 표현한 값입니다. 이 장치는 가상 함수이기 때문에 ASD, PFP, MES, SOS 같은 일부 펌웨어 조회는 지원되지 않으며, MEC·RLC·SDMA와 일부 RAS/XGMI 구성요소만 보고됩니다.

■ 실패한 측정에서 고친 제어-plane 설계

초기 구현은 명령의 종료 코드만 보고 GPU 상태를 판정했습니다. 그러나 드라이버가 초기화되지 않은 상황에서 `rocm-smi`는 stderr에 “Driver not initialized”를 출력하면서도 종료 코드 0을 반환했고, `amd-smi list` 역시 오류 세 줄을 출력하고 종료 코드 0을 반환했습니다. 이 때문에 빈 결과를 성공으로 보고하는 문제가 발생했습니다. 현재 `gpu_status`는 종료 코드를 신뢰하지 않고 `rocm-smi --json`을 파싱한 뒤, 카드가 하나도 추출되지 않으면 `amd-smi list`로 넘어갑니다.

메모리 사용량도 처음에는 잘못된 키를 읽었습니다. 실제 출력은 `GPU Memory Allocated (VRAM%)`였지만 요약기가 존재하지 않는 `GPU Memory Use (%)`를 찾았고, 그 결과 168.31GiB를 점유한 vLLM 프로세스를 유휴 상태처럼 보고했습니다. 현재는 실제 장비에서 얻은 JSON을 테스트 픽스처로 사용하고, 다른 ROCm 버전을 위한 대체 키도 함께 처리합니다. `rocminfo`에서는 에이전트와 ISA 모두 `Name:` 키를 사용해 `gfx942`가 두 번 세어지는 문제도 발생했으며, 현재 스캔은 `gfx<digits>` 형식의 타깃만 세어 GPU 수를 계산합니다.

새로 생성한 드롭릿에서는 `lspci`에 카드가 보이고 `/dev/dri/renderD128`도 있지만 `/dev/kfd`가 없어 GPU가 작동하지 않는 상황이 나타났습니다. 원인은 설치 누락이 아니라 VF에 `amdgpu`가 바인딩되지 않은 상태였으며, Debian의 커널 내장 amdgpu와 이미지의 ROCm 사용자 공간만으로 충분했습니다. `reboot_droplet`을 실행하고 `action_status`로 완료를 확인한 뒤 `gpu_status`를 호출하는 세 단계 절차로 복구했으며, 측정 당시 SSH는 액션 완료 약 20초 후 돌아왔습니다. VF에서 `amdgpu/psp_13_0_6_cap.bin` 로드 실패나 지원되지 않는 TA 유형 같은 `dmesg` 메시지는 문제의 원인이 아니었습니다.

도구는 셸 실행 시 `shell=True`를 사용하지 않고 `asyncio.create_subprocess_exec`를 통해 인자 목록으로 실행합니다. SSH는 `BatchMode=yes`로 설정해 잘못된 키가 비밀번호 입력을 기다리는 대신 즉시 실패하도록 합니다. 도구는 예외를 서버 밖으로 던지지 않고 오류 메시지로 반환하며, 읽기 전용·쓰기·파괴적 작업을 MCP annotation으로 구분합니다. 테스트는 토큰이나 네트워크 없이 패키지를 모킹한 상태에서 실행합니다.

■ 매트릭스 코어의 실제 수치 형식 측정

8192³ 행렬 곱을 30회 반복해 측정한 결과, 당시 카드가 동시에 서비스를 제공하고 있어 절대 성능은 낮아졌지만 형식 간 비율은 비교할 수 있었습니다. `torch 2.12.0+rocm10.0.0` 기준으로 BF16은 1.655ms, 664.3TFLOP/s, FP16은 1.660ms, 662.5TFLOP/s를 기록했습니다. 두 형식은 모두 이 카드에서 1307.4TFLOP/s의 동일한 피크를 사용하므로 실측 성능도 사실상 같습니다. FP16이 BF16보다 빠른 형식이라는 선택지는 이 하드웨어에서는 성립하지 않으며, 지수 범위가 더 좁은 FP16을 선택해도 속도 이득이 없습니다.

FP8 E4M3FNUZ는 0.938ms, 1172.5TFLOP/s로 BF16 대비 1.77배 빨랐습니다. 이론상 2배인 2614.9TFLOP/s에는 미치지 못하지만, 행렬 곱 전체에서 곱셈 이외의 부분을 고려하면 네이티브 경로에 해당하는 결과로 설명합니다. 반면 INT8은 2.413ms, 455.6TFLOP/s에 그쳐 BF16의 0.69배였습니다. AMD 사양표는 INT8을 FP8과 같은 2614.9TOPS로 제시하지만, 이 스택의 `torch._int_mm`은 튜닝된 경로에 도달하지 못했습니다. 따라서 사양표에서 같은 피크를 제시한다고 실제 머신에서 같은 성능이 나오는 것은 아닙니다.

MXFP4와 MXFP8은 기능 목록과 PyTorch 타입 정의에는 나타나지만 MI300X에서 네이티브로 실행되지 않습니다. `torch._scaled_mm`은 `Float8_e8m0fnu` 블록 스케일링이 `gfx950`, `gfx1250`에서만 지원된다고 `NotImplementedError`를 반환합니다. vLLM도 `gfx95`와 `gfx1250`에서만 `supports_mx()`를 참으로 판단합니다. 따라서 `gfx942`인 CDNA 3에서는 MX 형식을 거부하는 대신 에뮬레이션 커널로 넘어가며, 이 커널은 매 순방향 패스에서 4비트 가중치를 BF16으로 복원하고 활성화를 양자화·역양자화한 뒤 일반적인 BF16 `F.linear`를 실행합니다. 결과적으로 계산 중 메모리 절약이나 속도 이득 없이 복원 비용과 양자화 오차를 함께 부담합니다.

■ AMD용 FP8과 NVIDIA용 FP8은 같은 형식이 아닙니다

MI300X의 네이티브 FP8은 NVIDIA와 OCP에서 사용하는 `e4m3fn`이 아니라 CDNA 3의 `e4m3fnuz`입니다. `e4m3fn`의 최대 유한값은 448.0이고 최소 정규값은 0.015625인 반면, `e4m3fnuz`는 각각 240.0과 0.0078125입니다. 지수 바이어스가 1만큼 다르기 때문에 동일한 1바이트 비트 패턴 `0b01000000`도 전자는 2.0, 후자는 1.0으로 해석합니다. NVIDIA용 FP8 체크포인트의 바이트를 MI300X 형식으로 재해석하면 값이 달라지며, 448.0은 MI300X에서 NaN이 될 수 있습니다.

실제 호출에서도 `torch.float8_e4m3fn`은 `HIPBLAS_STATUS_NOT_SUPPORTED`로 실패하고, `torch.float8_e4m3fnuz`만 `_scaled_mm`을 통과했습니다. vLLM은 `gfx94`를 확인해 `torch.float8_e4m3fnuz`를 선택합니다. 따라서 이 카드에서는 다른 GPU용 FP8 체크포인트를 그대로 가져오기보다 BF16 가중치에서 실행 중 온라인 양자화하고, 실행될 장비에서 스케일을 계산하는 경로를 사용합니다.

■ vLLM·GGUF·실제 배포 상태

검사 시점의 환경에는 ROCm을 인식하는 `torch 2.9.1+rocm6.4`가 설치되어 있었지만 vLLM은 설치되어 있지 않았습니다. ROCm용 vLLM 사전 빌드 휠이 없고, 소스 빌드에는 `hipcc`, `rocblas`, `hipblaslt`, `miopen`, `rccl` 등이 필요하지만 해당 이미지에는 `hipcc`도 없었습니다. 실제 서비스는 `vllm/vllm-openai-rocm:nightly-rocm100` 컨테이너에서 실행 중이었고, `google/gemma-4-E2B-it` 모델이 3시간 동안 VRAM 168.31GiB를 점유하고 있었습니다.

이 빌드에는 GGUF 지원도 없습니다. `supported_quantization`과 전역 `QUANTIZATION_METHODS`에 `gguf`가 없고, 관련 모듈 import도 실패하며 `_custom_ops`에 ggml/gguf 심볼도 없습니다. 과거 빌드에서 `--quantization gguf`가 받아들여졌더라도, 현재는 엔진을 바꾸지 않는 한 사용할 수 없습니다. 또한 GGUF의 k-quants는 가중치 전용이며 커널 내부에서 복원된 뒤 산술은 BF16으로 수행되므로, MI300X 매트릭스 코어가 네이티브로 제공하지 않는 수치 형식을 새로 만들어 주지도 않습니다.

글의 결론은 MI300X에서 사용할 실용적인 선택을 `FP8 E4M3FNUZ`, 온라인 양자화, 그 외 선택지는 배제하는 한 줄로 정리합니다. AMD의 피크 표, 플랫폼 기능 목록, 모델 허브의 FP8 체크포인트가 각각 다른 방향을 가리킬 수 있으므로, 사양표나 라이브러리의 지원 목록만 보지 않고 실제 카드에서 커널 호출과 포맷 변환을 측정해야 한다고 설명합니다. DigitalOcean 공식 MCP 서버가 드롭릿 API와 플릿 관리를 담당하는 반면, 이 서버는 API 바깥의 원격 실행, GPU 상태, 하드웨어 인벤토리를 보완하는 구조로 제시합니다.

원문: dev.to / 번역·요약: Trawling