Testing GPU Code

Java-Forum
1

Hab das hier ueber Twitter „gefunden“: https://blog.esciencecenter.nl/writing-testable-gpu-code-23bbda3a5d62

Hab den Artikel gelesen, so gut wie nix verstanden, aber irgendwie scheint mir das Geschreibsel sehr optimistisch zu sein wenn es darum geht eine „Testpyramide“ fuer Code zu erstellen der sehr nah an der Hardware ist (Mocks/Stubs/Fakes?).
Auf der anderen Seite weiss ich viel zu wenig um mir wirklich eine Meinung zu bilden.

Wollte dazu @Marco13 Meinung hoeren, er ist ja schliesslich unser CUDA Spezialist hier :slight_smile:

Offtopic: Falls es ein besseres Unterforum gibt dafuer, kann jederzeit verschoben werden.

2

Sooo viel „beef“ hat der Artikel jetzt IMHO nicht :neutral_face:

Die einführenden/abschließenden Abschnitte „Unit Testing GPU Code“ und „Designing your test case“ sind recht allgemein, könnten sich auch auf nicht-GPU-Code beziehen, und sollen demnach wohl wirklich nur die Bühne vorbereiten für „Testing GPU Code“ und den kurzen Abschnitt zu „Testing device functions“.

Device functions zu Unit-Testen kann sicher Sinn machen. Die sind zwar oft (nicht immer) sehr kurz, aber eben „building blocks“, und damit eigentlich ideal zu testen. Das Problem (das er ja auch beschreibt) ist, dass man diese Funktionen oft nicht direkt aufrufen kann: Normalerweise haben CUDA-Kernels grob so eine Struktur:

__device__ int eineDeviceFunktion(...)
{
    // So eine Device-Funktion ist oft ein kleiner "Helfer", oder Building Block.
    // Diese Funktion kann NUR vom Kernel aus aufgerufen werden. 
    // D.h. man kann sie NICHT von Java/Python/C aus aufrufen,
    // sondern nur vom Kernel selbst (siehe unten)
}


__global__ void derEigentlicheKernel(...)
{
    // Das ist der Kernel, der von Java/Python/C aus aurfgferufen werden kann. 
    // (Auch nicht "direkt", wie eine normale Funktion, aber eben mit einem 
    // speziellen "führeDiesenKernelAus"-Befehl...)
    ...
    // VON hier aus kann die Device-Funktion aufgerufen werden
    int someResult = eineDeviceFunktion(...);
    ...
}

Sein Ansatz dafür ist wohl recht einfach - nämlich eine __global__-Funktion um die Device-Funktion drumzuwickeln:

__global__ void testAufrufFürDeviceFunktion(int x, int y, int z) {
    eineDeviceFunktion(x,y,z);
}

(Ehrlich gesagt müßte ich aber nochmal genauer schauen, wie das mit den Rückgabewerten gemacht werden soll: Kernels können keine Rückgabewerte haben)

Was dann noch übrig bleibt ist der Abschnitt „Testing GPU Code“, wo er beschreibt, wie die 7 Zwischenschritte, die notwendig sind, um einen Kernel aufzurufen, durch sein Framework und diese „run_kernel()“-Funktion eingedampft werden. Das macht sicher Sinn (so viel, dass ich vor 5 Jahren auch mal einen „CLKernelTester“ erstellt habe - das ganze hat es aber nie zur Veröffentlichung gebracht).

Böswillig könnte man sagen, dass er mit diesem Beitrag vielleicht nur seine Library dort promoten will :wink: aber erstens spricht da ja nichts dagegen, und zweitens könnte es ja sein, dass die Library es wert ist, promotet zu werden - hab’ bisher noch nicht näher reingesehen.

Was genau mit der „Testpyramide“ gemeint ist, ist mir nicht ganz klar. Ein echtes „Mocking“ in diesem Sinne könnte ja nur oberhalb der echten Hardware (GPU) stattfinden, und auch nur, wenn man die richtigen Abstraktionen einzieht. Das wirft (wenn man es etwas weiter denkt) recht schnell recht allgemeine Fragen auf - wie etwa, inwieweit man um den GPU-Teil eine echte, Objektorientierte Abstraktionsschicht wickeln kann und will. Einen einzelnen Kernel testbar zu machen ist eine Sache, aber das Hin- und Herschaufeln der Daten zwischen GPU und Host auch wegzuabstrahieren, so, dass man (praktisch) „einen Integrationstest laufen lassen kann, ohne, dass man die GPU braucht“ würde wohl nochmal deutlich mehr Aufwand erfordern.

1 „Gefällt mir“