Escrevendo uma especificação para Nvidia Kepler (binários CUDA, versão em idioma sm_30) para Ghidra

Muitas especificações para o Ghidra já foram escritas para linguagens comuns de processador, mas nada para linguagens gráficas. É compreensível, porque tem suas próprias especificidades: predicados, constantes através das quais os parâmetros são transmitidos e outras coisas herdadas dos shaders. Além disso, o formato usado para armazenar o código geralmente é proprietário e precisa ser revertido por conta própria.



Neste artigo, descobriremos o que é o quê para dois exemplos.



O primeiro programa é o axpy mais simples (analógico do hello world para GPGPU). O segundo ajuda a entender a implementação de condições e saltos na GPU, já que tudo é diferente lá.



Como todos os idiomas da Nvidia usam pouca codificação endian, copie imediatamente os bytes do editor hexadecimal para algum notebook (por exemplo, o Notepad ++) na ordem inversa de 8 partes (o comprimento das instruções é constante aqui). Então, através de uma calculadora de programador (a da Microsoft é adequada), nós a traduzimos em código binário. Em seguida, procuramos correspondências, compomos a máscara de instruções e depois os operandos. Para decodificar e procurar uma máscara, foram utilizados o editor hexadecimal e o desmontador cuobjdump, às vezes é necessário um montador, como no AMDGPU (porque o desmontador não está disponível lá, mas este é um tópico para um artigo separado). Funciona assim: tentamos inverter sequencialmente todos os bits suspeitos na calculadora, obtemos um novo valor hexadecimal para os bytes, substituímos por um binário compilado via nvcc ou um assembler, se existir, o que nem sempre é o caso.Em seguida, verificamos o cuobjdump.



Publico o código-fonte no formato (principalmente em C, sem vantagens e OOP para uma conexão mais próxima com o código da GPU da máquina), depois desastre + bytes de uma só vez, porque é mais conveniente, eles simplesmente não precisam ser trocados.



Copie-o para axpy.cu e compile-o via cmd: nvcc axpy.cu --cubin --gpu-architecture sm_30 Desmonte o

arquivo ELF resultante chamado axpy.cubin no mesmo local: cuobjdump axpy.cubin -sass



Exemplo 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];
}

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

É fácil adivinhar que os testes foram originalmente adaptados ao código da máquina para que o compilador não tivesse nada para otimizar. Para todo o resto, você teria que cancelar manualmente as otimizações. Em exemplos complexos, isso pode não ser possível, portanto, nesses casos, você precisará confiar no descompilador e no frontend.



Em geral, a regra é esta: para testar o frontend, tomamos qualquer exemplo simples (com o mínimo de otimizações possíveis) primeiro adequado (reprodução de erros). Quanto ao resto, o código descompilado já estará com otimizações (ou apenas de alguma forma correto através da refatoração). Mas, por enquanto, a tarefa principal é pelo menos apenas o código correto que faz a mesma coisa que o código de máquina. Isso é modelagem de software. A "modelagem de software" em si não implica refatoração, tradução de C para C ++, restauração de classes e, mais ainda, identificação de modelos.



Agora estamos procurando padrões para mnemônicos, operandos e modificadores.



Para fazer isso, compare os bits (em binário) entre as instruções suspeitas (ou seqüências de caracteres, se for mais conveniente chamar). Você também pode tirar proveito do que outros usuários publicam em suas perguntas sobre o stackoverflow, como “ajudar a entender o código binário / sass / de máquina”, usar tutoriais (inclusive em chinês) e outros recursos. Portanto, o número da operação principal é armazenado nos bits 58-63, mas existem bits adicionais 0-4 (eles distinguem entre as instruções “I2F”, “ISETP”, “MOV32I”), em algum lugar em seu lugar 0-2 (para negligenciar 3- 4 bits nas instruções vazias, na especificação eles são marcados como “UNK”).



Para registros e números constantes, você pode experimentar o desmontador para encontrar todos os bits que afetam a saída de despejo, como o disposto sob o spoiler. Todos os campos que eu consegui encontrar estão na especificação no Github, arquivo CUDA.slaspec, seção token.



Então você precisa criar endereços para os registros, novamente eles estão no Github. Isso é necessário porque no nível micro, Sleigh considera os registradores como variáveis ​​globais em um espaço com o tipo "register_space", mas desde seu espaço não está marcado como "inferível" (e provavelmente não pode ser), no descompilador elas se tornam variáveis ​​locais (geralmente com a interface "Var", mas às vezes havia também um prefixo "local") ou parâmetros (o " param_ "). O SP nunca foi útil, é necessário principalmente formalmente para garantir que o descompilador esteja funcionando. Um PC (algo como IP do x86) é necessário para emulação.



Depois, existem registros de predicado, algo como sinalizadores, mas mais "objetivo geral" do que para um propósito predeterminado, como excesso, igualdade (in) a zero etc.

Em seguida, um registro de bloqueio para modelar várias instruções ISCADD .CC e IMAD.HI, porque o primeiro deles em minha implementação realiza a contagem para si e para o segundo, a fim de evitar transferir parte da soma para os 4 bytes superiores, pois isso vai atrapalhar a descompilação. Mas você precisa bloquear o próximo registro até que a operação IMAD.HI seja concluída. Algo semelhante, ou seja, a discrepância entre a documentação oficial e a saída esperada do descompilador já estava na SPU para o mesmo Ghidra.



Depois, há registros especiais, que até agora são implementados através do cpool. No futuro, pretendo substituí-los pelos símbolos padrão para algum espaço "inferível". Estes são o mesmo threadIdx, blockIdx.



Então ligamos as variáveis ​​aos campos dest, par0, par1, par2, res. Depois, há sub-tabelas e depois delas - o que era tudo - as tabelas principais (raiz) com as instruções principais.



Aqui você precisa seguir rigorosamente o formato de "operandos mnemônicos", no entanto, é dada uma isenção para modificadores, que, no entanto, devem ser anexados a um mnemônico ou a uma seção com operandos. Nenhum outro formato é permitido, mesmo o mesmo Hexagon DSP asm precisará ser adaptado a essa sintaxe, que, no entanto, não é muito difícil.



O passo final será escrever uma implementação para obter instruções no idioma do firmware do Pcode. A única coisa que gostaria de observar no primeiro exemplo são as instruções ISCADD .CC e IMAD.HI, onde o primeiro deles pega um ponteiro para registrá-los e os desreferencia como ponteiros para 8 bytes em vez de 4. Isso foi feito intencionalmente para melhor se adaptar ao descompilador e seu comportamento, apesar do que está escrito na documentação da Nvidia sobre a transferência de parte do valor.



Para o segundo exemplo, é melhor marcar a caixa nas configurações do descompilador, ao lado da inscrição "Simplificar predicação". O ponto é que os predicados são a mesma condição para instruções diferentes; na verdade, nada mais que o conhecido "SIMD" ou seu próximo análogo. Essa. se o bit predicado estiver definido, as instruções serão executadas, além disso, em uma linha.



Você também precisa criar o hábito de escrever imediatamente uma implementação para cada instrução, e não apenas um protótipo (operando mnemônicos), porque também há um decompilador, emulador e outros analisadores.

Mas, em geral, escrever uma implementação no Pcode é uma tarefa ainda mais simples do que escrever uma gramática para um decodificador de bytes. Foi rápido corrigir a implementação de algumas instruções complexas do x86 (e não apenas), graças a uma linguagem intermediária muito conveniente, uma única região central (otimizador), 2 back-end (principalmente C; como alternativa - Java / C #, mais parecido com o último, etc.) (já que um goto aparece de tempos em tempos, mas não uma quebra rotulada).

Nos artigos a seguir, também pode haver front-ends para linguagens gerenciadas como DXBC, SPIR-V, eles usarão o back-end Java / C #. Mas até agora apenas os códigos de máquina estão nos planos. bytecodes requerem uma abordagem especial.



Projeto de

Ajuda Ghidra



:



Pcode

Sleigh