CUDA Unified Memory が便利すぎて生きるのが辛い。

わけがなかった。メモリ無駄遣いしすぎるのでそんなに便利ではないという印象だった。

CUDA 6.0がリリースされた。CUDA 6.0の目玉機能のひとつが、Unified Memoryだろう。Unified Memory については、

http://topsecret.hpc.co.jp/wiki/index.php/CUDA_6%E9%80%9F%E5%A0%B1%281%29:_Unified_Memory

で熟知すべし。

つまり、これまではGPU呼び出しの前にメモリ転送する必要があったが、それが必要無くなる。

https://devblogs.nvidia.com/parallelforall/unified-memory-in-cuda-6/

には、operator new を overload するだけで何もしなくていいぜ、ぐへへ。ぐらいに書いてある。

実装方法はページングを活用してコピーをユーザーから隠蔽するみたいな感じで、

http://news.mynavi.jp/column/sc13/004/

に解説がある。

と、いうところまでは、まあいいのだが、本当にoperator newをoverloadして大丈夫なのかは心配が残る。

調べた。まあ大丈夫ではなかった。メモリを無駄遣いしすぎ。

/* -*- c++ -*- */
#include <windows.h>
#include <psapi.h>
#include <stdio.h>
#include <stdio.h>

int
main(int argc, char **argv)
{
    PROCESS_MEMORY_COUNTERS pmc = {0};
    int alloc_size = 4;
    int alloc_count = 1024;

    if (argc > 1) {
        alloc_size = atoi(argv[1]);
    }

    if (argc > 2) {
        alloc_count = atoi(argv[2]) * 1024;
    }

    void **ptr_list = new void *[alloc_count];

    GetProcessMemoryInfo(GetCurrentProcess(), &pmc, sizeof(pmc));

    SIZE_T init_peak = pmc.PeakWorkingSetSize;
    SIZE_T init_ws = pmc.WorkingSetSize;
    SIZE_T diff;

    for (int li=0; li<5; li++) {
        for (int i=0; i<alloc_count; i++) {
            ptr_list[i] = malloc(alloc_size);
            memset(ptr_list[i], 0, alloc_size);
        }

        GetProcessMemoryInfo(GetCurrentProcess(), &pmc, sizeof(pmc));
        diff = (pmc.WorkingSetSize-init_ws) - (alloc_size * alloc_count);

        printf("%d : libc malloc(size(%d) * count=(%d) = %f[MiB]) ws=%f[MiB], overhead=%f[%%]\n",
               li,
               alloc_size,
               alloc_count,
               (alloc_size * alloc_count)/(1024.0*1024.0),
               (pmc.WorkingSetSize-init_ws)/(1024.0*1024.0),
               (diff*100.0) / (double)(alloc_size * alloc_count));

        for (int i=0; i<alloc_count; i++) {
            free(ptr_list[i]);
        }

        GetProcessMemoryInfo(GetCurrentProcess(), &pmc, sizeof(pmc));
        printf("%d : libc free : ws=%f[MiB], peak=%f[MiB]\n",
               li,
               (pmc.WorkingSetSize-init_ws)/(1024.0*1024.0),
               pmc.PeakWorkingSetSize/(1024.0*1024.0));
    }


    for (int li=0; li<5; li++) {
        for (int i=0; i<alloc_count; i++) {
            cudaError_t er = cudaMallocManaged(&ptr_list[i], alloc_size);
            memset(ptr_list[i], 0, alloc_size);
        }

        GetProcessMemoryInfo(GetCurrentProcess(), &pmc, sizeof(pmc));
        diff = (pmc.WorkingSetSize-init_ws) - (alloc_size * alloc_count);

        printf("%d : cuda malloc(size(%d) * count=(%d) = %f[MiB]) ws=%f[MiB], overhead=%f[%%]\n",
               li,
               alloc_size,
               alloc_count,
               (alloc_size * alloc_count)/(1024.0*1024.0),
               (pmc.WorkingSetSize-init_ws)/(1024.0*1024.0),
               (diff*100.0) / (double)(alloc_size * alloc_count));

        for (int i=0; i<alloc_count; i++) {
            cudaFree(ptr_list[i]);
        }

        GetProcessMemoryInfo(GetCurrentProcess(), &pmc, sizeof(pmc));
        printf("%d : cuda free : ws=%f[MiB], peak=%f[MiB]\n",
               li,
               (pmc.WorkingSetSize-init_ws)/(1024.0*1024.0),
               pmc.PeakWorkingSetSize/(1024.0*1024.0));
    }

}

4byte のメモリを1024回割り当てたときのWorkingSetSizeを見ると、796200%程のオーバーヘッドがある。cudaMallocManagedのポインタ値を見ると、どうやら32KB単位で割り当てているらしい。

まあ、アプリケーションの大部分がループで、巨大な配列を一個割り当てるだけみたいなアプリケーションであれば、全然問題無いけど、Unified Memory の目的を考えると、そうではないアプリケーション、つまり、それなりに巨大なプログラムで全ポインタの素性を調べるのがめんどいようなアプリケーションだと、小さいオブジェクト割り当てがメモリを圧迫するような状況も十分考えられるし、そういうような場合気軽にoperator newをoverloadするとか難しいと思う。