assembly - programming - first cuda program

¿Cómo crear o manipular el ensamblador GPU?

Question

¿Alguien tiene experiencia en la creación / manipulación de código de máquina GPU, posiblemente en tiempo de ejecución?

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

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

es posible? Cualquier información relacionada es bienvenida.

Answer 1

Estos enlaces pueden ser interesantes para usted, aunque es fácil encontrarlos, por lo que probablemente ya lo haya 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

Me ha gpuocelot proyecto de código abierto gpuocelot (Licencia BSD).

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

"Ocelot actualmente permite que los programas CUDA se ejecuten en GPU NVIDIA, GPU AMD y CPU x86". Por lo que sé, este marco de trabajo realiza análisis de flujo de control y flujo de datos en el núcleo de PTX para aplicar las transformaciones adecuadas.

Answer 3

OpenCL se hace para ese propósito. Proporciona un programa como una cadena y posiblemente lo compila en tiempo de ejecución. Vea los enlaces proporcionados por otro cartel.

Answer 4

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

Answer 5

Generación y modificación de NVIDIA PTX.

No estoy seguro de su nivel bajo en comparación con el hardware (¿es probable que no esté documentado?), Pero se puede generar desde lenguajes de GPU similares a C / C ++, modificados y reutilizados de varias maneras:

• OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary : ejemplo ejecutable mínimo: ¿Cómo usar clCreateProgramWithBinary en OpenCL?

  Estas son las API de OpenCL estandarizadas, que producen y consumen formatos definidos de implementación, que en la versión de controlador 375.39 para Linux resultan ser PTX legibles para los humanos.

  Así que puedes volcar el PTX, modificarlo y recargar.

• nvcc : puede compilar el código del lado de la GPU de CUDA en el ensamblaje de ptx simplemente con:

  nvcc --ptx a.cu

  nvcc también puede compilar programas OpenCL C que contienen tanto el dispositivo como el código del host: ¿ Compilar y compilar el archivo .cl usando el compilador nvcc de NVIDIA? pero no pude encontrar cómo sacar el ptx con nvcc. Qué tipo de sentido tiene, ya que son simples cadenas C + C, y no 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 : cómo compilar el código PTX

Usando clGetProgramInfo , un kernel de entrada CL:

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

se compila a algunos PTX como:

//
// 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 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 lugar de 1 como se esperaba.

Answer 6

En la API del controlador CUDA, las funciones de administración del módulo permiten que una aplicación cargue en tiempo de ejecución un "módulo", que es (aproximadamente) un archivo PTX o cubin. PTX es el lenguaje 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 real.

Entonces, su problema parece ser: ¿cómo construir un módulo cubin mediante programación en la RAM? Que yo sepa, 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 para lo que entiende la GPU NVIDIA "más antigua" ("más antigua" = GeForce 8xxx y 9xxx). No sé qué tan fácil sería integrarme en una aplicación más amplia; Está escrito en Python.

Las GPU NVIDIA más nuevas utilizan un conjunto de instrucciones distinto (no sé cuánto), por lo que una cubina para una GPU antigua ("capacidad de computación 1.x" en la terminología de NVIDIA / CUDA) puede no funcionar en una GPU reciente (capacidad de computación 2.x, es decir, "arquitectura Fermi", como una GTX 480). Es por eso que generalmente se prefiere PTX: un archivo PTX dado será portátil en todas las generaciones de GPU.