¿Cómo crear o manipular el ensamblador de GPU?

Question

¿Alguien tiene experiencia en crear/manipular códigos de máquina GPU, posiblemente en tiempo de ejecución?

Estoy interesado en modificar el código de ensamblador de la GPU, posiblemente en tiempo de ejecución con una sobrecarga mínima. Específicamente, estoy interesado en la programación genética basada en ensambladores.

Entiendo que ATI ha publicado ISA para algunas de sus tarjetas, y nvidia recientemente lanzó un desensamblador para CUDA para tarjetas antiguas, pero no estoy seguro de si es posible modificar instrucciones en memoria en tiempo de ejecución o incluso antes.

¿Esto es posible? Cualquier información relacionada es bienvenida.

Answer 1

Estos enlaces podrían ser interesantes para usted, a pesar de que es fácil de encontrar, así que probablemente ya 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 está hecho para ese propósito. Usted proporciona un programa como una cadena y posiblemente compila en tiempo de ejecución. Vea los enlaces provistos por otro afiche.

Answer 3

En la API del controlador CUDA, el module management functions permite que una aplicación cargue en tiempo de ejecución un "módulo", que es (más o menos) un archivo PTX o cubin. PTX es el idioma intermedio, mientras que cubin es un conjunto de instrucciones ya compilado. cuModuleLoadData() y cuModuleLoadDataEx() parecen ser capaces de "cargar" el módulo desde un puntero en la RAM, lo que significa que no se requiere ningún archivo.

Así que su problema parece ser: ¿cómo construir mediante programación un módulo de cubos en la RAM? Hasta donde yo sé, NVIDIA nunca dio a conocer detalles sobre las instrucciones realmente entendidas por su hardware. Sin embargo, hay un paquete de código abierto independiente llamado decuda que incluye "cudasm", un ensamblador de lo que la GPU NVIDIA "anterior" entiende ("más antiguo" = GeForce 8xxx y 9xxx). No sé cuán fácil sería integrar en una aplicación más amplia; está escrito en Python.

GPU NVIDIA más nuevos usan un conjunto de instrucciones distintas (qué tan distintas, no sé), por lo que un cubin para una antigua GPU ("capacidad informática 1.x" en terminología NVIDIA/CUDA) puede no funcionar en un GPU (capacidad informática 2.x, es decir, "arquitectura Fermi" como GTX 480). Por eso es preferible usar PTX: un archivo PTX determinado será portátil a través de las generaciones de GPU.

Answer 4

Un ensamblador para el ISA de Fermi de NVIDIA: http://code.google.com/p/asfermi

Answer 5

que he encontrado gpuocelot de código abierto (licencia BSD) proyecto interesante.

Es "un marco de compilación dinámico para PTX". Yo lo llamaría traductor de cpu.

"Ocelot actualmente permite que los programas CUDA se ejecuten en GPU NVIDIA, GPU AMD y CPU x86".Hasta donde yo sé, este marco de trabajo hace un flujo de control y análisis de flujo de datos en Kernel PTX para aplicar transformaciones apropiadas.

Answer 6

generación de NVIDIA PTX y modificación

No

seguro de cómo bajo nivel que se compara con el hardware (probablemente indocumentado?), Pero puede ser generada a partir de C/C++ - como los lenguajes de la GPU, modificados y reutilizados en algunas de las formas:

• OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary: mínima ejemplo ejecutable: How to use clCreateProgramWithBinary in OpenCL?

  Estos son estandarizados OpenC L API's, que producen y consumen formatos de implementación definidos, que en el controlador versión 375.39 para Linux pasa a ser PTX legible para el ser humano.

  Para poder descargar el PTX, modificarlo y volver a cargarlo.

• nvcc: puede compilar CUDA código del lado del GPU para PTX montaje simple, ya sea con:

  nvcc --ptx a.cu 
  

  nvcc también puede compilar programas OpenCL C que contengan tanto del dispositivo como código de host: Compile and build .cl file using NVIDIA's nvcc Compiler? pero no pude encontrar la manera de saca el ptx con nvcc. ¿Qué tipo de sentido tiene? Ya que solo se trata de cadenas C + C, y no de un superconjunto C mágico. Esto también es sugerido por: https://arrayfire.com/generating-ptx-files-from-opencl-code/

  y no estoy seguro de cómo volver a compilar el PTX modificado y usarlo como lo hice con clCreateProgramWithBinary: How to compile PTX code

Usando clGetProgramInfo, un núcleo de entrada CL:

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

se compila a algunos les gusta PTX:

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

Entonces, si por ejemplo se modifica la línea:

add.s32  %r3, %r2, 1; 

a:

add.s32  %r3, %r2, 2; 

y reutilizar el PTX modificado, en realidad se incrementa en 2 en vez de 1 como se esperaba.