Как создать или манипулировать графическим ассемблером? - PullRequest
Новая
       25

Как создать или манипулировать графическим ассемблером?

19 голосов
/ 11 января 2011

Есть ли у кого-нибудь опыт создания / управления машинным кодом графического процессора, возможно, во время выполнения?

Меня интересует изменение кода ассемблера графического процессора, возможно, во время выполнения с минимальными издержками.В частности, меня интересует генетическое программирование на ассемблере.

Я понимаю, что ATI выпустила ISA для некоторых своих карт, а nvidia недавно выпустила дизассемблер для CUDA для старых карт, но я не уверен, возможно лиизменять инструкции в памяти во время выполнения или даже перед этим.

Возможно ли это?Любая связанная информация приветствуется.

Ответы [ 6 ]

3 голосов
/ 13 января 2011
2 голосов
/ 15 августа 2011

Мне показался интересным gpuocelot проект с открытым исходным кодом (лицензия BSD).

Это "среда динамической компиляции для PTX".Я бы назвал это переводчиком cpu.

«В настоящее время Ocelot позволяет выполнять программы CUDA на графических процессорах NVIDIA, AMD GPU и x86-CPU».Насколько я знаю, эта структура выполняет анализ потока управления и потока данных в ядре PTX, чтобы применить правильные преобразования.

2 голосов
/ 13 января 2011

В API драйвера CUDA функции управления модулями позволяют приложению загружать во время выполнения «модуль», который (примерно) представляет собой файл PTX или кубин.PTX - это промежуточный язык, а cubin - это уже скомпилированный набор инструкций.cuModuleLoadData() и cuModuleLoadDataEx(), по-видимому, способны «загружать» модуль из указателя в ОЗУ, что означает, что фактический файл не требуется.

Итак, ваша проблема заключается в следующем: как программно построитьмодуль кубин в оперативной памяти?Насколько я знаю, NVIDIA никогда не публиковала подробности инструкций, фактически понятных их аппаратному обеспечению.Однако существует независимый пакет с открытым исходным кодом, называемый decuda , который включает в себя «cudasm», ассемблер для того, что понимают «старые» графические процессоры NVIDIA («более старые» = GeForce 8xxx и 9xxx).Я не знаю, насколько легко было бы интегрироваться в более широкое приложение;он написан на Python.

Более новые графические процессоры NVIDIA используют отдельный набор инструкций (насколько они различны, я не знаю), поэтому кубин для старого графического процессора («вычислительная способность 1.x» в NVIDIA / CUDA)терминология) может не работать на недавнем графическом процессоре (вычислительная возможность 2.x, то есть «архитектура Fermi», такая как GTX 480).Вот почему PTX обычно предпочтительнее: данный файл PTX будет переносимым для поколений графических процессоров.

1 голос
/ 13 августа 2011

Ассемблер для NVIDIA Fermi ISA: http://code.google.com/p/asfermi

1 голос
/ 13 января 2011

OpenCL сделан для этой цели. Вы предоставляете программу в виде строки и, возможно, компилируете ее во время выполнения. Смотрите ссылки предоставленные другим автором.

0 голосов

Создание и модификация NVIDIA PTX

Не уверен, насколько он низок по сравнению с аппаратным обеспечением (вероятно, недокументированным?), Но он может быть сгенерирован из C / C ++ - как языки GPU, изменен и повторно использован несколькими способами:

  • OpenCL clGetProgramInfo(program, CL_PROGRAM_BINARIES + clCreateProgramWithBinary: пример минимального запуска: Как использовать clCreateProgramWithBinary в OpenCL?

    Это стандартизированные API-интерфейсы OpenCL, которые производят и используют определенные форматы реализации, которые в версии драйвера 375.39 для Linux считаются PTX, читаемым человеком.

    Таким образом, вы можете сбросить PTX, изменить его и перезагрузить.

  • nvcc: можно скомпилировать код на стороне графического процессора CUDA в сборку ptx просто с помощью: 

    nvcc --ptx a.cu

    nvcc также может компилировать программы OpenCL C, содержащие код устройства и хоста: Компилировать и собирать файл .cl с помощью NVIDIA nvcc Compiler? , но я не смог найти, как вывести ptx с помощью nvcc. Какой вид имеет смысл, поскольку это просто строки C + C, а не магический надмножество C. Это также предлагается: https://arrayfire.com/generating-ptx-files-from-opencl-code/

    И я не уверен, как перекомпилировать измененный PTX и использовать его, как я делал с clCreateProgramWithBinary: Как скомпилировать код PTX

Используя clGetProgramInfo, ядро ​​входного CL: 

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

компилируется в некоторый 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;
}

Тогда, если, например, вы измените строку: 

add.s32     %r3, %r2, 1;

до: 

add.s32     %r3, %r2, 2;

и повторно использовать измененный PTX, он фактически увеличивается на 2 вместо 1, как и ожидалось.

...