ROCmのHIPでのメモリ割り当て

こちらでhipMallocをhipHostMallocに動的に置き換えるforce-host-alloction-APUを利用してStable DifussionをRyzen APUで動作させています。

ROCmの文書よりHIPのメモリ割り当てAPIについて、hipMallocではDevice(GPU)からメモリを取得し、hipHostMallocではHost(CPU)からメモリを取得するとなっています。

CPUもGPUもメインメモリを共有するUnified Memory Architecture(UMA)のRyzen APU上のROCmのHIPで、どの領域からどれくらいメモリ取得可能か不明瞭だったので、HIP-ExcampleのHelloWorldを参考に、次の簡単なプログラムで調査してみました。

調査環境

種類	内容
CPU	Ryzen 5600G
マザーボード	ASRock B450M-HDV (BIOS 4.70)
メモリ	CFD W4U3200CS-16G (16GB×2=32GB)
OS	Fedora 40

調査用プログラム

#include <hip/hip_runtime.h>
#include <iostream>

using namespace std;

int main(int argc, char* argv[])
{
    char *buffer[256];
    int i, j;

    for (i = 0 ; i < 256; i++) {
        if (hipMalloc((void**)&buffer[i], 128*1024*1024) == hipSuccess)
            cout << "ALLOCATE: " << (i + 1) * 128 << " MB" << endl;
        else {
            cout << "DO NOT ALLOCATE: " << (i + 1) * 128 << " MB" << endl;
            break;
        }
    }
    for (j = 0; j < i - 1; j++) {
        hipFree(buffer[j]);
    }
}

Memory.cppとして保存してください。128 MB単位でメモリを確保し、成功した場合(hipSuccess)は『ALLOCATE: これまで確保した容量』を表示し、失敗した場合は『DO NOT ALLOCATE: 確保しようとした容量』と表示します。最大128×256=32768 MB(32 GB)まで確保しようとします。もっと大きく確保したいならchar *buffer[256];と for (i = 0 ; i < 256; i++) {の256を増やしてください。

調査用プログラムのビルド

hipcc等をHIPの開発環境を入れてください。Fedora 40なら

$ sudo dnf install clang hipcc lld lld-devel compiler-rt

です。

$ hipcc -o Memory Memory.cpp

でビルドしてください。

force-host-allocation-APUのビルド

force-host-alloction-APUもgitで取得してビルドしてください。

$ git clone https://github.com/segurac/force-host-alloction-APU
$ cd force-host-alloction-APU/
$ hipcc forcegttalloc.c -o libforcegttalloc.so  -shared -fPIC
$ cd ..

BIOSでの手動のVRAM割り当てなしの場合

調査環境のメインメモリは32 GBでスワップが 8 GBです。

$ free
               total        used        free      shared  buff/cache   available
Mem:        32210192     9760280     4147096      427692    19197216    22449912
Swap:        8388604        6144     8382460

dmesgコマンドでカーネルログを見ると、VRAMが512 MBでGTTが15727 MBとなっています。

$ sudo dmesg | grep amdgpu
...
[    4.585655] [drm] amdgpu: 512M of VRAM memory ready
[    4.585656] [drm] amdgpu: 15727M of GTT memory ready.
...

hipMalloc

hipMallocはVRAMからメモリを確保します。調査用プログラムをKDE Plasma環境で実行すると

$ ./Memory
ALLOCATE: 128 MB
ALLOCATE: 256 MB
ALLOCATE: 384 MB
DO NOT ALLOCATE: 512 MB

384 MBまで確保に成功し512 MBの確保は失敗します。VRAM 512 MBに対してKDEもVRAMを使用しているので、512 MB確保できないのは正しい挙動だと思います。

hipHostMalloc

force-host-alloction-APUを使ってhipMallocをhipHostMallocに置き換えて実行します。hipHostMallocはメインメモリからメモリを確保します。

$ LD_PRELOAD=./force-host-alloction-APU/libforcegttalloc.so ./Memory
ALLOCATE: 128 MB
ALLOCATE: 256 MB
ALLOCATE: 384 MB
ALLOCATE: 512 MB
ALLOCATE: 640 MB
ALLOCATE: 768 MB
ALLOCATE: 896 MB
...
ALLOCATE: 29184 MB
ALLOCATE: 29312 MB
DO NOT ALLOCATE: 29440 MB

29440 MBで確保に失敗します。30 GB弱までは確保できています。メインメモリが32 GBでスワップ 8 GBなので、他のプログラムが使用しているメモリを除くと30 GB程度で失敗したようです。

実際に確保したメモリをGPUが正しく使用できるかまではわかりませんが、ROCmのHIPでは16 GB以上確保できるようです。GTTの15727 MBは関係ないようです。

BIOSで16 GB VRAMに割り当てた場合

$ sudo dmesg | grep amdgpu
...
[    4.326371] [drm] amdgpu: 16384M of VRAM memory ready
[    4.326372] [drm] amdgpu: 7927M of GTT memory ready.
...

VRAMが16384 MBでGTTが7927 MBです。

Memoryを実行すると、

$ ./Memory
ALLOCATE: 128 MB
ALLOCATE: 256 MB
ALLOCATE: 384 MB
ALLOCATE: 512 MB
ALLOCATE: 640 MB
ALLOCATE: 768 MB
ALLOCATE: 896 MB
...
ALLOCATE: 16000 MB
ALLOCATE: 16128 MB
ALLOCATE: 16256 MB
DO NOT ALLOCATE: 16384 MB

16256 MBまで確保できて16384 MBの確保に失敗します。VRAMが16384 MBなので他のプログラムの使用分を考えると正しい挙動です。

force-host-alloction-APUを使ってhipMallocをhipHostMallocに置き換えて実行した場合、

$ LD_PRELOAD=./force-host-alloction-APU/libforcegttalloc.so./Memory
ALLOCATE: 128 MB
ALLOCATE: 256 MB
ALLOCATE: 384 MB
ALLOCATE: 512 MB
...
ALLOCATE: 13312 MB
ALLOCATE: 13440 MB
ALLOCATE: 13568 MB
ALLOCATE: 13696 MB
DO NOT ALLOCATE: 13824 MB

13824 MBで確保に失敗します。VRAMに16 GB割り当てているのでメインメモリが32 GBあっても16 GBしか使えないので他のプログラムの使用分を考えると正しい挙動だと思います。こちらもGTTの7927 MBとは関係なくメモリを確保しています。

まとめ

BIOSでVRAMを割り当てた場合最大16 GBまでの制限があるようですが、hipHostMallocでメモリを確保する場合HIPプログラムはメインメモリの空き容量分メモリを確保できます。

外付けGPUの場合GPUのVRAMとメインメモリではアクセス速度に大きな違いがあり、hipMallocをhipHostMallocに置き換えるのは有効でないと思われますが、APUの場合は最初からメインメモリしかないのでhipMallocをhipHostMallocに置き換えてもあくせす速度は変わらないはずです。メインメモリは外付けGPUのVRAMと違って容易に増設可能なので、VRAMを多く必要とする大規模言語モデル等の実行にRyzen APUが役立つ可能性があります。