Wie erstellt oder manipuliert man GPU-Assembler?

Question

Hat jemand Erfahrung in der Erstellung/Manipulation von GPU-Maschinencode, möglicherweise zur Laufzeit?

Ich bin daran interessiert, GPU-Assembler-Code möglicherweise zur Laufzeit mit minimalen Aufwand zu ändern. Insbesondere interessiere ich mich für Assembler-basierte genetische Programmierung.

Ich verstehe ATI hat ISAs für einige ihrer Karten veröffentlicht, und nvidia hat vor kurzem einen Disassembler für CUDA für ältere Karten veröffentlicht, aber ich bin nicht sicher, ob es möglich ist, Anweisungen im Speicher zur Laufzeit oder sogar vorher zu ändern.

Ist das möglich? Alle damit verbundenen Informationen sind willkommen.

Answer 1

Diese Links könnten für Sie interessant sein, obwohl seine leicht, sie zu finden, so haben Sie wahrscheinlich schon gesehen:

http://www.worldlingo.com/ma/enwiki/en/ARB_(GPU_assembly_language)

http://developer.nvidia.com/object/gpu_programming_guide.html

http://developer.amd.com/gpu/Pages/default.aspx

http://msdn.microsoft.com/en-us/library/bb219840.aspx

http://www.khronos.org/opencl/

http://www.comp.nus.edu.sg/~ashwinna/docs/CS6282_Modeling_the_GPU.pdf

Answer 2

OpenCL wird zu diesem Zweck getan. Sie geben ein Programm als String an und übersetzen es möglicherweise zur Laufzeit. Siehe Links von anderen Postern.

Answer 3

In der CUDA-Treiber-API erlaubt es die module management functions einer Anwendung, zur Laufzeit ein "Modul" zu laden, das (ungefähr) eine PTX- oder Cubin-Datei ist. PTX ist die Zwischensprache, während Cubin ein bereits kompilierter Satz von Anweisungen ist. cuModuleLoadData() und cuModuleLoadDataEx() scheinen in der Lage zu sein, das Modul von einem Zeiger im RAM zu "laden", was bedeutet, dass keine tatsächliche Datei benötigt wird.

Ihr Problem scheint also zu sein: Wie programmiere ich ein Cubin-Modul im RAM? Soweit ich weiß, hat NVIDIA niemals Details zu den Anweisungen veröffentlicht, die von ihrer Hardware tatsächlich verstanden werden. Es gibt jedoch ein unabhängiges Open Source-Paket namens decuda, das "cudasm" enthält, einen Assembler für das, was die "ältere" NVIDIA-GPU versteht ("älter" = GeForce 8xxx und 9xxx). Ich weiß nicht, wie einfach es wäre, in eine breitere Anwendung zu integrieren; Es ist in Python geschrieben.

Neuere NVIDIA GPU verwenden einen eindeutigen Befehlssatz (wie viel distinct, ich weiß nicht), so ein Cubin für eine alte GPU ("Rechenfähigkeit 1.x" in NVIDIA/CUDA Terminologie) funktioniert möglicherweise nicht auf einem aktuellen GPU (Rechenfähigkeit 2.x, dh "Fermi-Architektur" wie eine GTX 480). Aus diesem Grund wird PTX in der Regel bevorzugt: Eine bestimmte PTX-Datei ist über GPU-Generationen hinweg portierbar.

Answer 4

Ein Assembler für die NVIDIA Fermi ISA: http://code.google.com/p/asfermi

Answer 5

ich interessant gpuocelot Open-Source (BSD-Lizenz) Projekt gefunden haben.

Es ist "ein dynamisches Kompilierungs-Framework für PTX". Ich würde es cpu-Übersetzer nennen.

"Ocelot ermöglicht derzeit die Ausführung von CUDA-Programmen auf NVIDIA GPUs, AMD GPUs und x86-CPUs".Soweit ich weiß, führt dieses Framework eine Kontrollfluss- und Datenflussanalyse auf dem PTX-Kernel durch, um geeignete Transformationen anzuwenden.

Answer 6

NVIDIA PTX Generation und Modifikation

Nicht sicher, wie niedriges Niveau an die Hardware verglichen wird (wahrscheinlich ohne Papiere?), Aber es kann von C erzeugt werden/C++ - wie GPU Sprachen, modifiziert und wiederverwendet in ein paar Möglichkeiten:

• OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary: minimal runnable Beispiel: How to use clCreateProgramWithBinary in OpenCL?

  Dies sind standardisierte OpenC L-APIs, die implementation-definierte Formate erzeugen und konsumieren, die in Treiber-Version 375.39 für Linux zufällig lesbare PTX-Dateien sind.

  So können Sie die PTX dump, ändern und neu laden.

• nvcc: kann CUDA GPU-Seite Code kompilieren Baugruppe ptx einfach mit entweder:

  nvcc --ptx a.cu 
  

  nvcc kann OpenCL C-Programme auch beide Gerät und dem Host-Code enthält kompilieren: Compile and build .cl file using NVIDIA's nvcc Compiler? aber ich konnte nicht finden, wie man Holen Sie den PTX mit nvcc heraus. Welche Art von Sinn macht es, da es nur einfache C + C-Strings und kein magischer C-Supersatz ist. Dies wird auch durch vorgeschlagen: https://arrayfire.com/generating-ptx-files-from-opencl-code/

  Und ich bin nicht sicher, wie die modifizierte PTX neu kompilieren und es verwenden, wie ich mit clCreateProgramWithBinary tat: How to compile PTX code

clGetProgramInfo verwenden, einen Eingang CL-Kernel:

__kernel void kmain(__global int *out) { 
    out[get_global_id(0)]++; 
} 

wird bis zu einem gewissen PTX zusammengestellt mag:

// 
// Generated by NVIDIA NVVM Compiler 
// 
// Compiler Build ID: CL-21124049 
// Cuda compilation tools, release 8.0, V8.0.44 
// Based on LLVM 3.4svn 
// 

.version 5.0 
.target sm_20 
.address_size 64 

    // .globl _Z3incPi 

.visible .entry _Z3incPi(
    .param .u64 _Z3incPi_param_0 
) 
{ 
    .reg .pred %p<2>; 
    .reg .b32 %r<4>; 
    .reg .b64 %rd<5>; 


    ld.param.u64 %rd1, [_Z3incPi_param_0]; 
    mov.u32  %r1, %ctaid.x; 
    setp.gt.s32 %p1, %r1, 2; 
    @%p1 bra BB0_2; 

    cvta.to.global.u64 %rd2, %rd1; 
    mul.wide.s32 %rd3, %r1, 4; 
    add.s64  %rd4, %rd2, %rd3; 
    ldu.global.u32 %r2, [%rd4]; 
    add.s32  %r3, %r2, 1; 
    st.global.u32 [%rd4], %r3; 

BB0_2: 
    ret; 
} 

Dann, wenn zum Beispiel die Zeile ändern:

add.s32  %r3, %r2, 1; 

zu:

add.s32  %r3, %r2, 2; 

und Wiederverwendung der PTX geändert, erhöht er tatsächlich um 2 statt 1 wie erwartet.