Как создать или использовать ассемблер GPU?

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

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

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

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

Ответ 1

Эти ссылки могут быть вам интересны, хотя их легко найти, поэтому, возможно, вы уже видели это:

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

Ответ 2

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

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

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

Ответ 3

Я нашел gpuocelot проект с открытым исходным кодом (BSD License).

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

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

Ответ 4

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

Ответ 5

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

Ответ 6

Генерация и модификация NVIDIA PTX

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

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

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

Итак, вы можете сбросить PTX, изменить его и перезагрузить.
nvcc: можно скомпилировать код на стороне процессора CUDA для сборки ptx просто:
```
nvcc --ptx a.cu
```
nvcc также может компилировать программы OpenCL C, содержащие как код устройства, так и код хоста: Скомпилировать и построить файл .cl с использованием nvcc-компилятора NVIDIA, но я не смог найти способ чтобы получить ptx out с 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, как ожидалось.