Come creare o manipolare l'assemblatore GPU?

Question

Qualcuno ha esperienza nella creazione/manipolazione del codice macchina GPU, possibilmente in fase di esecuzione?

Sono interessato a modificare il codice dell'assemblatore GPU, possibilmente in fase di esecuzione con un sovraccarico minimo. Specificamente sono interessato alla programmazione genetica basata sugli assemblatori.

Comprendo che ATI ha rilasciato ISA per alcune delle sue schede e nvidia ha recentemente rilasciato un disassemblatore per CUDA per le vecchie schede, ma non sono sicuro che sia possibile modificare le istruzioni in memoria in fase di esecuzione o anche prima.

È possibile? Qualsiasi informazione correlata è benvenuta.

Answer 1

Questi collegamenti potrebbero essere interessanti per voi, anche se la sua facile trovarli, quindi probabilmente già visto:

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 è fatto a tal fine. Fornisci un programma come stringa e, eventualmente, lo compili in fase di runtime. Vedi i collegamenti forniti da altri poster.

Answer 3

Nell'API del driver CUDA, lo module management functions consente a un'applicazione di caricare in fase di esecuzione un "modulo", che è (approssimativamente) un file PTX o cubino. PTX è la lingua intermedia, mentre cubin è un insieme di istruzioni già compilato. cuModuleLoadData() e cuModuleLoadDataEx() sembrano essere in grado di "caricare" il modulo da un puntatore nella RAM, il che significa che non è richiesto alcun file effettivo.

Quindi il problema sembra essere: come creare un modulo cubin in modo programmatico nella RAM? Per quanto ne so, NVIDIA non ha mai rilasciato dettagli sulle istruzioni effettivamente comprese dal loro hardware. Esiste, tuttavia, un pacchetto opensource indipendente chiamato decuda che include "cudasm", un assemblatore per ciò che la GPU NVIDIA "vecchia" comprende ("precedente" = GeForce 8xxx e 9xxx). Non so quanto sarebbe facile integrarlo in un'applicazione più ampia; è scritto in Python.

Le GPU NVIDIA più recenti utilizzano un set di istruzioni distinto (quanto distinto, non lo so), quindi un cubin per una GPU precedente ("capacità di calcolo 1.x" nella terminologia NVIDIA/CUDA) potrebbe non funzionare su un recente GPU (capacità di calcolo 2.x, ovvero "architettura Fermi" come una GTX 480). Per questo motivo PTX è generalmente preferito: un determinato file PTX sarà trasferibile attraverso le generazioni di GPU.

Answer 4

Un assemblatore per la NVIDIA Fermi ISA: http://code.google.com/p/asfermi

Answer 5

ho trovato gpuocelot open-source (licenza BSD) progetto interessante.

È "una struttura di compilazione dinamica per PTX". Lo chiamerei traduttore cpu.

"Attualmente Ocelot consente l'esecuzione dei programmi CUDA su GPU NVIDIA, GPU AMD e CPU x86".Per quanto ne so, questo framework fa analisi del flusso di controllo e del flusso di dati sul kernel PTX per applicare le opportune trasformazioni.

Answer 6

generazione NVIDIA PTX e la modifica

Non

sicuro di come basso livello che viene confrontato con l'hardware (probabilmente non documentato?), Ma può essere generato da C/C++ - come le lingue GPU, modificati e riutilizzati in alcuni modi:

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

  Questi sono standardizzati OPENÇ L API, che producono e consumano formati definiti dall'implementazione, che nella versione driver 375.39 per Linux sembra essere PTX leggibile dall'uomo.

  Quindi è possibile scaricare il PTX, modificarlo e ricaricarlo.

• nvcc: può compilare CUDA codice GPU-side per PTX montaggio semplice con:

  nvcc --ptx a.cu 
  

  nvcc può anche compilare i programmi OpenCL C contenenti sia dispositivo e il codice host: Compile and build .cl file using NVIDIA's nvcc Compiler? ma non riuscivo a trovare il modo di ottenere il ptx con nvcc. Che tipo di ha senso dal momento che si tratta di semplici stringhe C + C, e non di un superset C magico. Questo è suggerito anche da: https://arrayfire.com/generating-ptx-files-from-opencl-code/

  E io non sono sicuro di come ricompilare il PTX modificato e utilizzarlo come ho fatto con clCreateProgramWithBinary: How to compile PTX code

Utilizzando clGetProgramInfo, un kernel di ingresso CL:

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

viene compilato in una certa PTX piace:

// 
// 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; 
} 

Quindi se per esempio si modifica la linea:

add.s32  %r3, %r2, 1; 

a:

add.s32  %r3, %r2, 2; 

e riutilizzare il PTX modificato, in realtà aumenta di 2 invece di 1 come previsto.