Estamos escribiendo una especificación para Nvidia Kepler (binarios de CUDA, versión de idioma sm_30) para Ghidra

Para los lenguajes de procesador comunes, ya se han escrito muchas especificaciones para Ghidra, pero nada para las gráficas. Es comprensible, porque tiene sus propios detalles: predicados, constantes a través de los cuales se transmiten los parámetros, incluidas otras cosas heredadas de los sombreadores. Además, el formato utilizado para almacenar el código a menudo es exclusivo y debe revertirlo usted mismo.



En este artículo, usando dos ejemplos, descubriremos qué es qué.



El primer programa es el axpy más simple (análogo de hello world para GPGPU). El segundo ayuda a comprender la implementación de condiciones y saltos en la GPU, ya que Todo es diferente allí.



Todos los idiomas de Nvidia usan una pequeña codificación endian, por lo que copie inmediatamente los bytes del editor hexadecimal en algún cuaderno (por ejemplo, Notepad ++) en orden inverso de 8 piezas (la longitud de las instrucciones es constante aquí). Luego, a través de la calculadora del programador (la de Microsoft es adecuada), la traducimos a código binario. A continuación, buscamos coincidencias, conformamos la máscara de la instrucción, luego los operandos. Para decodificar y buscar una máscara, se utilizaron el editor hexadecimal y el desensamblador cuobjdump, a veces se requiere un ensamblador, como en AMDGPU (porque el desensamblador no está disponible allí, pero este es un tema para un artículo separado). Funciona así: intente invertir secuencialmente todos los bits sospechosos en la calculadora, luego obtenemos un nuevo valor hexadecimal para los bytes, los sustituimos en un binario compilado a través de nvcc o ensamblador si existe, lo cual no siempre es el caso.Luego a través de cuobjdump lo comprobamos.



Difundí la fuente en el formato (principalmente en C, sin más y OOP para una conexión más cercana con el código de la GPU de la máquina), luego desmonte + bytes a la vez, porque es más conveniente, simplemente no necesitan intercambiarse.



Lo copiamos a axpy.cu y lo compilamos a través de cmd: nvcc axpy.cu --cubin --gpu-architecture sm_30

El archivo ELF resultante llamado axpy.cubin se desmonta en el mismo lugar: cuobjdump axpy.cubin -sass



Ejemplo 1:

__global__ void axpy(float param_1, float* param_2, float* param_3) {
unsigned int uVar1 = threadIdx.x;
param_2[uVar1] = param_1 * param_3[uVar1];
}

Ejemplo 2

__global__ void predicates(float* param_1, float* param_2) {
    unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
    if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
    else param_2[uVar1] = uVar1;
}

Es fácil adivinar que las pruebas se adaptaron originalmente al código de máquina para que el compilador no tuviera nada que optimizar. Para todo lo demás, tendría que cancelar manualmente las optimizaciones. En ejemplos complejos, esto puede no ser posible en absoluto, por lo que para tales casos tendrá que confiar en el descompilador y la interfaz.



En general, la regla es esta: para probar el front-end, tomamos cualquier ejemplo simple (con un mínimo de optimizaciones posibles) adecuado (errores de reproducción). Por lo demás, el código descompilado ya tendrá optimizaciones (o solo lo corregirá de alguna manera mediante la refactorización). Pero por ahora, la tarea principal es al menos el código correcto que hace lo mismo que el código de máquina. Esto es "modelado de software". El "modelado de software" en sí mismo no implica refactorización, traducción de C a C ++, restauración de clases, y aún más cosas como la identificación de plantillas.



Ahora estamos buscando patrones para mnemónicos, operandos y modificadores.



Para hacer esto, compare los bits (en representación binaria) entre instrucciones sospechosas (o cadenas, si es más conveniente llamarlas de esa manera). También puede usar lo que otros usuarios publican en sus preguntas en stackoverflow como "ayudarme a comprender el código binario / sass / máquina", usar tutoriales (incluso en chino) y otros recursos. Entonces, el número de operación principal se almacena en los bits 58-63, pero también hay bits adicionales 0-4 (distinguen las instrucciones "I2F", "ISETP", "MOV32I"), en algún lugar en lugar de ellos 0-2 (por negligencia, 3- 4 bits en instrucciones en blanco, están marcados como "UNK" en la especificación).



Para registros y números constantes, puede experimentar con un desensamblador para encontrar todos los bits que afectan la salida de volcado, como el que se encuentra debajo del spoiler. Todos los campos que logré encontrar están en la especificación de Github, archivo CUDA.slaspec, token de sección.



Luego debe encontrar direcciones para los registros, nuevamente están en Github. Esto es necesario porque a nivel micro, Sleigh considera los registros como variables globales en un espacio con el tipo "register_space", pero desde su espacio no está marcado como "inflable" (y lo más probable es que no pueda serlo), luego en el descompilador se convierten en variables locales (la mayoría de las veces con el interfijo "Var", pero a veces también existía el prefijo "local") o parámetros (el " param_ "). SP nunca fue útil, se necesita principalmente formalmente para garantizar que el descompilador funcione. Se necesita una PC (algo así como IP de x86) para emular.



Luego hay registros de predicados, algo así como banderas, pero más "propósito general" que para un propósito predeterminado, como desbordamiento, (in) igualdad a cero, etc.

Luego, un registro de bloqueo para modelar un montón de instrucciones ISCADD .CC e IMAD.HI, porque el primero de ellos en mi implementación realiza el conteo por sí mismo y el segundo, para evitar transferir parte de la suma a los 4 bytes superiores, ya que Esto estropeará la descompilación. Pero luego debe bloquear el siguiente registro hasta que se complete la operación IMAD.HI. Algo similar, es decir La discrepancia entre la documentación oficial y la salida esperada del descompilador ya estaba en el módulo SPU para el mismo Ghidra.



Luego hay registros especiales, que hasta ahora se implementan a través de cpool. En el futuro, planeo reemplazarlos con los símbolos predeterminados para algún espacio "inferable". Estos son el mismo threadIdx, blockIdx.



Luego vinculamos las variables a los campos dest, par0, par1, par2, res. Luego están las sub-tablas, y después de ellas, de qué se trataba, las tablas principales (raíz) con las instrucciones principales.



Aquí es necesario seguir estrictamente el formato de "mnemotecnia de operandos", sin embargo, se da alivio a los modificadores, que, sin embargo, deben adjuntarse a la mnemotecnia oa la sección con operandos. No se permiten otros formatos, incluso el mismo Hexagon DSP asm tendrá que adaptarse a esta sintaxis, lo que, sin embargo, no es muy difícil.



La etapa final será escribir la implementación de las instrucciones en el lenguaje de microprogramación Pcode. Lo único que me gustaría señalar del primer ejemplo son las instrucciones ISCADD .CC e IMAD.HI, donde el primero toma un puntero para registrar y desreferenciarlos como punteros para 8 bytes en lugar de 4. Esto se hace intencionalmente para adaptarse mejor al descompilador y su comportamiento, a pesar de lo que está escrito en la documentación de Nvidia sobre la transferencia de parte de la cantidad.



Para el segundo ejemplo, es mejor verificar la configuración del descompilador opuesta a la inscripción "Simplificar predicación". El punto es que los predicados son una y la misma condición para diferentes instrucciones, de hecho, nada más que el conocido "SIMD", o su próximo análogo. Aquellos. si se establece el bit de predicado, las instrucciones se ejecutan, además, en una fila.



También debe acostumbrarse a escribir inmediatamente una implementación para cada instrucción, y no solo un prototipo (operandos mnemotécnicos), porque también hay un descompilador, un emulador y otros analizadores.

Pero, en general, escribir una implementación en Pcode es una tarea aún más fácil que escribir una gramática para un decodificador de bytes. Resultó rápidamente solucionar la implementación de algunas instrucciones complejas de x86 (y no solo), gracias a un lenguaje intermedio muy conveniente, un solo middleland (optimizador), 2 backends (principalmente C; como alternativa: Java / C #, más como este último, es decir, K. goto aparece de vez en cuando, pero no rotura rotulada).

En los siguientes artículos, también puede haber interfaces para lenguajes administrados como DXBC, SPIR-V, usarán el backend Java / C #. Pero hasta ahora solo los códigos de máquina están en los planes. Los bytecodes requieren un enfoque especial. Referencias del



Proyecto

Ghidra



:



Pcode

Sleigh