Jak utworzyć lub zmodyfikować asembler GPU?

Question

Czy ktoś ma doświadczenie w tworzeniu/manipulowaniu kodem maszyny GPU, ewentualnie w czasie wykonywania?

Interesuje mnie modyfikowanie kodu asemblera GPU, prawdopodobnie w czasie wykonywania z minimalnym narzutem. W szczególności interesuje mnie programowanie genetyczne oparte na asemblerach.

Rozumiem, że ATI wypuściło ISA dla niektórych swoich kart, a nvidia niedawno wypuściła dezasembler dla CUDA dla starszych kart, ale nie jestem pewien, czy możliwe jest zmodyfikowanie instrukcji w pamięci w czasie pracy lub nawet przed rozdaniem.

Czy to możliwe? Wszelkie powiązane informacje są mile widziane.

Answer 1

Te linki może być dla Ciebie interesujący, chociaż łatwo jest je znaleźć, więc pewnie już widział go:

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 służy do tego celu. Podajesz program jako ciąg i ewentualnie kompilujesz go w czasie wykonywania. Zobacz linki dostarczone przez inny plakat.

Answer 3

W interfejsie API sterownika CUDA, module management functions pozwala aplikacji na ładowanie w środowisku wykonawczym "modułu", którym jest (w przybliżeniu) plik PTX lub plik cubin. PTX jest językiem pośrednim, podczas gdy cubin jest już skompilowanym zbiorem instrukcji. cuModuleLoadData() i cuModuleLoadDataEx() wydają się być zdolne do "ładowania" modułu ze wskaźnika w pamięci RAM, co oznacza, że ​​żaden rzeczywisty plik nie jest wymagany.

Twój problem wydaje się być następujący: jak programowo zbudować moduł cubin w pamięci RAM? Z tego, co wiem, firma NVIDIA nigdy nie ujawniła szczegółów instrukcji faktycznie zrozumianych przez ich sprzęt. Istnieje jednak niezależny pakiet open source o nazwie decuda, który zawiera "cudasm", asembler dla tego, co rozumie "starszy" procesor graficzny NVIDIA ("starszy" = GeForce 8xxx i 9xxx). Nie wiem, jak łatwo byłoby zintegrować się z szerszą aplikacją; jest napisany w Pythonie.

Nowsze procesory graficzne NVIDIA używają odrębnego zestawu instrukcji (jak dużo różnią się, nie wiem), więc nowy procesor GPU ("możliwości obliczania 1.x" w terminologii NVIDIA/CUDA) może nie działać na ostatnim GPU (zdolność obliczeniowa 2.x, tj. "Architektura Fermiego", taka jak GTX 480). Z tego powodu zazwyczaj preferowane jest PTX: dany plik PTX będzie przenośny w pokoleniach GPU.

Answer 4

asemblera dla NVIDIA Fermi ISA: http://code.google.com/p/asfermi

Answer 5

Znalazłem gpuocelot open-source (licencja BSD) projekt ciekawy.

To "dynamiczna struktura kompilacji dla PTX". Nazwałbym to tłumaczem cpu.

"Obecnie Ocelot pozwala na uruchamianie programów CUDA na układach GPU NVIDIA, procesorach AMD i procesorach x86".O ile mi wiadomo, ta struktura wykonuje analizę sterowania przepływem i przepływem danych w jądrze PTX, aby zastosować odpowiednie transformacje.

Answer 6

NVIDIA PTX pokolenie i modyfikacje

Nie wiem, jak niski poziom jest ona w porównaniu do sprzętu (prawdopodobnie nieudokumentowanych?), Ale może być generowane z C/C++ - jak językach GPU, modyfikowane i ponownie wykorzystane w kilka sposobów:

• OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary: minimalny uruchamialny przykład: How to use clCreateProgramWithBinary in OpenCL?

  są standaryzowane OpenC L API, które produkują i zużywają implementację zdefiniowanych formatów, które w wersji sterownika 375.39 dla Linuksa wydają się być czytelnymi dla człowieka PTX.

  Możesz więc zrzucić PTX, zmodyfikować i przeładować.

• nvcc: można skompilować CUDA kod GPU-bocznego do ptx zespół prosto z albo:

  nvcc --ptx a.cu 
  

  nvcc można również przygotować OpenCL C programów zawierających zarówno urządzenia i kod hosta: Compile and build .cl file using NVIDIA's nvcc Compiler? ale nie mogłem znaleźć sposobu pobierz ptx z nvcc. Jaki ma sens, ponieważ jest to po prostu ciągi C + C, a nie magiczny C-zestaw. Jest to również sugerowane przez: https://arrayfire.com/generating-ptx-files-from-opencl-code/

  I nie jestem pewien, jak skompilować zmodyfikowaną PTX i używać go tak jak ja z clCreateProgramWithBinary: How to compile PTX code

Korzystanie clGetProgramInfo, jądro CL wejściowe:

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

zostanie skompilowany do pewnego PTX lubię:

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

Następnie, jeśli na przykład zmodyfikować linię:

add.s32  %r3, %r2, 1; 

do:

add.s32  %r3, %r2, 2; 

i ponowne użycie PTX modyfikowane, to rzeczywiście zwiększa się o 2 zamiast 1, jak oczekiwano.