What to do #2

New Issue

Daniel · 2024-07-02T16:54:45+02:00

Daniel commented

2024-07-02 16:54:45 +02:00

evalGPU(expr, data)

expr = simple Julia mathematical expression
data = Vector/Matrix with values for the variables in the expression

example Vector: (2, 4, 5, 1)
example expr: x + 1
x should then be filled with each vector entry and calculated simultaniously

Briefe overview of what needs to be done:

First Step: Use CUDA.jl to write julia code that gets transpiled into PTX. This julia code is the interpreter, meaning the entire interpreter is running on the GPU -> most likely slow
Second Step: evalGPUFast(...)
Instead of running the interpreter on the GPU, run it in julia to generate the PTX code and send that via CUDA.jl to the GPU for evaluation

evalGPU(expr, data) expr = simple Julia mathematical expression data = Vector/Matrix with values for the variables in the expression - example Vector: (2, 4, 5, 1) - example expr: x + 1 - x should then be filled with each vector entry and calculated simultaniously <br> Briefe overview of what needs to be done: - First Step: Use CUDA.jl to write julia code that gets transpiled into PTX. This julia code is the interpreter, meaning the entire interpreter is running on the GPU -> most likely slow - Second Step: evalGPUFast(...) - Instead of running the interpreter on the GPU, run it in julia to generate the PTX code and send that via CUDA.jl to the GPU for evaluation

Daniel commented

2024-07-03 17:07:35 +02:00

Sehr geehrter Herr Wiplinger,

wie besprochen schicke ich ihnen noch mehr Material für die Masterarbeit.

Folgende Bücher habe ich hilfreich gefunden:

CUDA by Example - An Introduction to General-Purpose GPU Programming (Jason Sanders, Edward Kandrot)
The CUDA Handbook - A Comprehensive Guide to GPU Programming – Addison-Wesley Pearson (Wilt N., (2013))

Konkretisierung der Aufgabenstellung:
Gegeben ist eine Matrix X mit k Spalten und N Zeilen, und ein Julia Ausdruck, der Spalten der Matrix über die Symbole x1 .. xk referenziert. Beispiele für Ausrücke sind:

:(x1 * (exp(-x2) + 2.0).
:(x1^2)
:(log(abs(x1 / x2)))
:(abs(x1)^2)

Zusätzlich zu den Variablen und Konstanten können Ausdrücke auch Parameter (Vektor p = p1 .. pm) verwenden. Beispiele:

:(x1 + p1)
:(abs(x1*p1)^p2)

Es sollen zwei Funktionen zur Verfügung gestellt werden:

function interpret_gpu(exprs::Vector{Expression}, X::Matrix{Float64}, p::Vector{Vector{Float64}})::Matrix{Float64}
function evaluate_gpu(exprs::Vector{Expression}, X::Matrix{Float64}, p::Vector{Vector{Float64}})::Matrix{Float64}

Beide Funktionen werten mehrere Expressions auf der GPU aus damit die GPU möglichst gut genutzt werden kann. Die X Matrix ist für alle Expressions gleich und ist eine k x N Matrix. P enthält den Parameter Vektor für jede Expression. Es ist ein Vektor von Vektoren, da Expressions unterschiedlich viele Parameter haben können. p[i] ist der Vektor der Parameterwerte für expr[i]. Das Ergebnis ist eine Matrix mit length(exprs) Spalten und N Zeilen, jede Spalte ist das Ergebnis einer Expression.

Folgende Funktionen sollen unterstützt werden: *, /, +, -, ^, abs, log, exp, sqrt.

Zwei Varianten sollen getestet werden:

Implementierung eines einfachen Interpretierers der vollständig auf der GPU läuft. Hier muss man vorsichtig sein, dass alle Threads innerhalb eines Warps (meistens 32 threads) alle denselben Branch nehmen. Der Interpreter verarbeitet die Expressions sinnvollerweise in einer Postfix-Repräsentation der Expressions die man noch auf der CPU vorbereiten kann. Z.B.: „x1 + x2 * p1“ -> „x1, x2, p1, *, +“. Der Interpreter läuft über die Postfix-Repräsentation und führt dann in einer Mehrfachverzeigung anhand den richtigen Code für die Operation aus (wie in der LVA FCW). Die ganze Interpreterschleife für alle Expressions soll vollständig auf der GPU laufen. Hier wird die Schwierigkeit sein die GPU Cores maximal zu belasten.
Implementierung eines Compilers / Transpilers von den Julia Expressions in PTX code. https://docs.nvidia.com/cuda/parallel-thread-execution/index.html) https://stackoverflow.com/questions/67752857/how-can-i-create-an-executable-to-run-a-kernel-in-a-given-ptx-file. Die Übersetzung ist vermutlich etwas schwieriger zu implementieren weil PTX eine Register-Maschine annimmt und Sie deshalb bei der Übersetzung von Julia Expressions auf PTX eine einfache Strategie für Registerallokation implementieren müssen. Dazu ist das Compiler-Buch (Dragon Book) hilfreich. Hier könnte der Aufwand für die Übersetzung von PTX auf nativen Code (wird vom NVIDIA CUDA Treiber gemacht) ein Falschenhals sein.

Die beiden Ansätze sollen mit verschiedenen Ausdrücken und unterschiedlich langen Matrizen (N=1000, 10000, 10^5, 10^6) getestet werden und die Ressourcennutzung der GPU analysiert werden.

Folgende Papers scheinen interessant zu sein:
https://umtl.cs.uni-saarland.de/paper_preprints/paper_koester_mpie_2020.pdf
https://ieeexplore.ieee.org/document/9370323

ich habe noch ein paar andere Papers zu dem Thema gefunden, aber nicht weiterverfolgt.

Hier noch ein einfaches Beispiel wie man in Julia zur Laufzeit einen PTX String erzeugen und aufrufen kann. Die Verwendung der Funktion ist im Kommentar der ersten Zeile gezeigt. In dem Einfachen Beispiel wird an der Stelle op im String einfach eine der 4 Operationen aufgerufen. Der Rest des PTX-Codes ist für das Durchlaufen der N Elemente auf verschiedenen Threads notwendig. Der Code verwenden das CUDA.jl Paket

# culoadtest(N, rand(["add.f32", "sub.f32", "mul.f32", "div.approx.f32"]))
function culoadtest(N::Int32, op = "add.f32")
vadd_code = ".version 7.1
.target sm_52
.address_size 64

        // .globl       VecAdd_kernel

.visible .entry VecAdd_kernel(
        .param .u64 VecAdd_kernel_param_0,
        .param .u64 VecAdd_kernel_param_1,
        .param .u64 VecAdd_kernel_param_2,
        .param .u32 VecAdd_kernel_param_3
)
{
        .reg .pred      %p<2>;
        .reg .f32       %f<4>;
        .reg .b32       %r<6>;
        .reg .b64       %rd<11>;

        ld.param.u64    %rd1, [VecAdd_kernel_param_0];
        ld.param.u64    %rd2, [VecAdd_kernel_param_1];
        ld.param.u64    %rd3, [VecAdd_kernel_param_2];
        ld.param.u32    %r2, [VecAdd_kernel_param_3];
        mov.u32         %r3, %ntid.x;
        mov.u32         %r4, %ctaid.x;
        mov.u32         %r5, %tid.x;
        mad.lo.s32      %r1, %r3, %r4, %r5;
        setp.ge.s32     %p1, %r1, %r2;
        @%p1 bra        \$L__BB0_2;

        cvta.to.global.u64      %rd4, %rd1;
        mul.wide.s32    %rd5, %r1, 4;
        add.s64         %rd6, %rd4, %rd5;
        cvta.to.global.u64      %rd7, %rd2;
        add.s64         %rd8, %rd7, %rd5;
        ld.global.f32   %f1, [%rd8];
        ld.global.f32   %f2, [%rd6];" *
        op *
        "               %f3, %f2, %f1;
        cvta.to.global.u64      %rd9, %rd3;
        add.s64         %rd10, %rd9, %rd5;
        st.global.f32   [%rd10], %f3;
 
\$L__BB0_2:
        ret;
}"

    linker = CuLink()
    add_data!(linker, "VecAdd_kernel", vadd_code)

    image = complete(linker)

    mod = CuModule(image)
    func = CuFunction(mod, "VecAdd_kernel")

    d_a = CUDA.fill(1.0f0, N)
    d_b = CUDA.fill(2.0f0, N)
    d_c = CUDA.fill(0.0f0, N)

    # Grid/Block configuration
    threadsPerBlock = 256;
    blocksPerGrid  = (N + threadsPerBlock - 1) ÷ threadsPerBlock;

    @time CUDA.@sync cudacall(func, Tuple{CuPtr{Cfloat},CuPtr{Cfloat},CuPtr{Cfloat},Cint}, d_a, d_b, d_c, N; threads=threadsPerBlock, blocks=blocksPerGrid)

end

Für die Arbeit mit Julia empfehle ich entweder auf einem Linux System oder mit WSL zu arbeiten. VS Code eignet sich gut als Editor. Das Revise Paket https://timholy.github.io/Revise.jl/stable/ ist sinnvoll, um Änderungen der Code-Files sofort in der REPL verfügbar zu machen (für Sie völlig transparent). Es ist sinnvoll gleich von Beginn an den Code in einem Paket zu schreiben. https://julialang.org/contribute/developing_package/ Das Paket können Sie vorerst lokal entwickeln, es muss nicht veröffentlicht werden. Wenn Sie mir einen Zugang zu dem github Repository geben, kann ich die Entwicklung ein wenig mitverfolgen. Mein Account ist @gkronber.

Falls Sie Fragen haben, können Sie sich jederzeit gerne melden.

Freundliche Grüße,
Gabriel Kronberger

Sehr geehrter Herr Wiplinger, wie besprochen schicke ich ihnen noch mehr Material für die Masterarbeit. Folgende Bücher habe ich hilfreich gefunden: - CUDA by Example - An Introduction to General-Purpose GPU Programming (Jason Sanders, Edward Kandrot) - The CUDA Handbook - A Comprehensive Guide to GPU Programming – Addison-Wesley Pearson (Wilt N., (2013)) Konkretisierung der Aufgabenstellung: Gegeben ist eine Matrix X mit k Spalten und N Zeilen, und ein Julia Ausdruck, der Spalten der Matrix über die Symbole x1 .. xk referenziert. Beispiele für Ausrücke sind: - :(x1 * (exp(-x2) + 2.0). - :(x1^2) - :(log(abs(x1 / x2))) - :(abs(x1)^2) Zusätzlich zu den Variablen und Konstanten können Ausdrücke auch Parameter (Vektor p = p1 .. pm) verwenden. Beispiele: - :(x1 + p1) - :(abs(x1*p1)^p2) Es sollen zwei Funktionen zur Verfügung gestellt werden: - function ```interpret_gpu(exprs::Vector{Expression}, X::Matrix{Float64}, p::Vector{Vector{Float64}})::Matrix{Float64}``` - function ```evaluate_gpu(exprs::Vector{Expression}, X::Matrix{Float64}, p::Vector{Vector{Float64}})::Matrix{Float64}``` Beide Funktionen werten mehrere Expressions auf der GPU aus damit die GPU möglichst gut genutzt werden kann. Die X Matrix ist für alle Expressions gleich und ist eine k x N Matrix. P enthält den Parameter Vektor für jede Expression. Es ist ein Vektor von Vektoren, da Expressions unterschiedlich viele Parameter haben können. p[i] ist der Vektor der Parameterwerte für expr[i]. Das Ergebnis ist eine Matrix mit length(exprs) Spalten und N Zeilen, jede Spalte ist das Ergebnis einer Expression. Folgende Funktionen sollen unterstützt werden: *, /, +, -, ^, abs, log, exp, sqrt. Zwei Varianten sollen getestet werden: - Implementierung eines einfachen Interpretierers der vollständig auf der GPU läuft. Hier muss man vorsichtig sein, dass alle Threads innerhalb eines Warps (meistens 32 threads) alle denselben Branch nehmen. Der Interpreter verarbeitet die Expressions sinnvollerweise in einer Postfix-Repräsentation der Expressions die man noch auf der CPU vorbereiten kann. Z.B.: „x1 + x2 * p1“ -> „x1, x2, p1, *, +“. Der Interpreter läuft über die Postfix-Repräsentation und führt dann in einer Mehrfachverzeigung anhand den richtigen Code für die Operation aus (wie in der LVA FCW). Die ganze Interpreterschleife für alle Expressions soll vollständig auf der GPU laufen. Hier wird die Schwierigkeit sein die GPU Cores maximal zu belasten. - Implementierung eines Compilers / Transpilers von den Julia Expressions in PTX code. https://docs.nvidia.com/cuda/parallel-thread-execution/index.html) https://stackoverflow.com/questions/67752857/how-can-i-create-an-executable-to-run-a-kernel-in-a-given-ptx-file. Die Übersetzung ist vermutlich etwas schwieriger zu implementieren weil PTX eine Register-Maschine annimmt und Sie deshalb bei der Übersetzung von Julia Expressions auf PTX eine einfache Strategie für Registerallokation implementieren müssen. Dazu ist das Compiler-Buch (Dragon Book) hilfreich. Hier könnte der Aufwand für die Übersetzung von PTX auf nativen Code (wird vom NVIDIA CUDA Treiber gemacht) ein Falschenhals sein. Die beiden Ansätze sollen mit verschiedenen Ausdrücken und unterschiedlich langen Matrizen (N=1000, 10000, 10^5, 10^6) getestet werden und die Ressourcennutzung der GPU analysiert werden. Folgende Papers scheinen interessant zu sein: https://umtl.cs.uni-saarland.de/paper_preprints/paper_koester_mpie_2020.pdf https://ieeexplore.ieee.org/document/9370323 ich habe noch ein paar andere Papers zu dem Thema gefunden, aber nicht weiterverfolgt. Hier noch ein einfaches Beispiel wie man in Julia zur Laufzeit einen PTX String erzeugen und aufrufen kann. Die Verwendung der Funktion ist im Kommentar der ersten Zeile gezeigt. In dem Einfachen Beispiel wird an der Stelle op im String einfach eine der 4 Operationen aufgerufen. Der Rest des PTX-Codes ist für das Durchlaufen der N Elemente auf verschiedenen Threads notwendig. Der Code verwenden das CUDA.jl Paket ```cuda # culoadtest(N, rand(["add.f32", "sub.f32", "mul.f32", "div.approx.f32"])) function culoadtest(N::Int32, op = "add.f32") vadd_code = ".version 7.1 .target sm_52 .address_size 64 // .globl VecAdd_kernel .visible .entry VecAdd_kernel( .param .u64 VecAdd_kernel_param_0, .param .u64 VecAdd_kernel_param_1, .param .u64 VecAdd_kernel_param_2, .param .u32 VecAdd_kernel_param_3 ) { .reg .pred %p<2>; .reg .f32 %f<4>; .reg .b32 %r<6>; .reg .b64 %rd<11>; ld.param.u64 %rd1, [VecAdd_kernel_param_0]; ld.param.u64 %rd2, [VecAdd_kernel_param_1]; ld.param.u64 %rd3, [VecAdd_kernel_param_2]; ld.param.u32 %r2, [VecAdd_kernel_param_3]; mov.u32 %r3, %ntid.x; mov.u32 %r4, %ctaid.x; mov.u32 %r5, %tid.x; mad.lo.s32 %r1, %r3, %r4, %r5; setp.ge.s32 %p1, %r1, %r2; @%p1 bra \$L__BB0_2; cvta.to.global.u64 %rd4, %rd1; mul.wide.s32 %rd5, %r1, 4; add.s64 %rd6, %rd4, %rd5; cvta.to.global.u64 %rd7, %rd2; add.s64 %rd8, %rd7, %rd5; ld.global.f32 %f1, [%rd8]; ld.global.f32 %f2, [%rd6];" * op * " %f3, %f2, %f1; cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd5; st.global.f32 [%rd10], %f3; \$L__BB0_2: ret; }" linker = CuLink() add_data!(linker, "VecAdd_kernel", vadd_code) image = complete(linker) mod = CuModule(image) func = CuFunction(mod, "VecAdd_kernel") d_a = CUDA.fill(1.0f0, N) d_b = CUDA.fill(2.0f0, N) d_c = CUDA.fill(0.0f0, N) # Grid/Block configuration threadsPerBlock = 256; blocksPerGrid = (N + threadsPerBlock - 1) ÷ threadsPerBlock; @time CUDA.@sync cudacall(func, Tuple{CuPtr{Cfloat},CuPtr{Cfloat},CuPtr{Cfloat},Cint}, d_a, d_b, d_c, N; threads=threadsPerBlock, blocks=blocksPerGrid) end ``` Für die Arbeit mit Julia empfehle ich entweder auf einem Linux System oder mit WSL zu arbeiten. VS Code eignet sich gut als Editor. Das Revise Paket https://timholy.github.io/Revise.jl/stable/ ist sinnvoll, um Änderungen der Code-Files sofort in der REPL verfügbar zu machen (für Sie völlig transparent). Es ist sinnvoll gleich von Beginn an den Code in einem Paket zu schreiben. https://julialang.org/contribute/developing_package/ Das Paket können Sie vorerst lokal entwickeln, es muss nicht veröffentlicht werden. Wenn Sie mir einen Zugang zu dem github Repository geben, kann ich die Entwicklung ein wenig mitverfolgen. Mein Account ist @gkronber. Falls Sie Fragen haben, können Sie sich jederzeit gerne melden. Freundliche Grüße, Gabriel Kronberger

Daniel closed this issue

2024-09-01 11:17:02 +02:00

Sign in to join this conversation.

No Label

No Milestone

No project

No Assignees

1 Participants

Notifications

Due Date

The due date is invalid or out of range. Please use the format 'yyyy-mm-dd'.

No due date set.

Dependencies

No dependencies set.

Reference: Daniel/master-thesis#2