Вопрос по cuda, assembly, inline-assembly – Как я могу вызвать функцию PTX из CUDA C?

3

Я пытаюсь найти способ вызова функции ptx (.func) из CUDA C. Скажем, у меня была функция PTX, как это:

.func (.reg .s32 %res) inc_ptr ( .reg .s32 %ptr, .reg .s32 %inc )
{
    add.s32 %res, %ptr, %inc;
    ret;
}

Я знаю, что могу назвать это из ptx так:

call (%d), inc_ptr, (%s, %d);

Но я понятия не имею, как назвать это из CUDA C. Я знаю, что могу встроить ptx-сборку с помощью asm (), но я не нашел способа встроить функцию. Надеюсь, кто-то может помочь!

Спасибо!

@RogerDahl Да, это то, чего я тоже боялся. Однако похоже, что в CUDA 5 может появиться компоновщик кода устройства (developer.download.nvidia.com/assets/cuda/files/CUDADownloads/…). Я не узнал, как это сделать в CUDA 5 Preview. fursund
К сожалению, я думаю, что это может быть невозможно. Проблема состоит в том, что CUDA C не поставляется с компоновщиком для кода устройства. Поэтому все, что запускается из одного ядра, должно находиться в одном и том же файле .cu. Надеюсь, я ошибаюсь, потому что, IMO, NVIDIA плохо справилась с синтаксисом встроенного PTX. Roger Dahl

Ваш Ответ

2   ответа
2

представленных в CUDA 5.0. Я не верю, что есть способ сделать это в «целом» режим компиляции программы или в версиях инструментария до CUDA 5.0 или в ревизиях PTX до 3.1.

Вероятно, проще всего проиллюстрировать, как это сделать, на работающем примере. Давайте начнем с простой функции PTX для увеличения указателей, как в вашем примере:

.version 3.1
.target sm_30
.address_size 32
.visible .func inc_ptr(.param .b32 ptr, .param .b32 inc)
{
    .reg .s32   %r<6>;
    ld.param.u32 %r1, [ptr];
    ld.param.u32 %r2, [inc];
    ld.u32 %r3, [%r1];
    ld.u32 %r4, [%r3];
    add.s32 %r5, %r4, %r2;
    st.u32  [%r3], %r5;
    ret;
}

Это может быть скомпилировано в объект перемещаемого устройства с помощьюptxas и затем упакован в файл-контейнер. Последний шаг кажется критическим. По умолчаниюptxas выход только перемещаемыйelf объект, жировой контейнер не производится. Кажется, что фаза связывания кода устройства, которую запускает nvcc (по крайней мере, в CUDA 5), ожидает, что весь код устройства присутствует в толстых двоичных контейнерах. В противном случае связь потерпит неудачу. Результат выглядит так:

$ ptxas -arch=sm_30 -c -o inc_ptr.gpu.o inc_ptr.ptx
$ fatbinary -arch=sm_30 -create inc_ptr.fatbin -elf inc_ptr.gpu.o 
$ cuobjdump -sass inc_ptr.fatbin 

Fatbin elf code:
================
arch = sm_30
code version = [1,6]
producer = <unknown>
host = mac
compile_size = 32bit

    code for sm_30
        Function : inc_ptr
    /*0008*/     /*0x0040dc8580000000*/     LD R3, [R4];
    /*0010*/     /*0x00301c8580000000*/     LD R0, [R3];
    /*0018*/     /*0x14001c0348000000*/     IADD R0, R0, R5;
    /*0020*/     /*0x00301c8590000000*/     ST [R3], R0;
    /*0028*/     /*0x00001de790000000*/     RET;
    /*0030*/     /*0x00001de440000000*/     NOP CC.T;
    /*0038*/     /*0x00001de440000000*/     NOP CC.T;
    /*0040*/     /*0xe0001de74003ffff*/     BRA 0x40;
    /*0048*/     /*0x00001de440000000*/     NOP CC.T;
    /*0050*/     /*0x00001de440000000*/     NOP CC.T;
    /*0058*/     /*0x00001de440000000*/     NOP CC.T;
    /*0060*/     /*0x00001de440000000*/     NOP CC.T;
    /*0068*/     /*0x00001de440000000*/     NOP CC.T;
    /*0070*/     /*0x00001de440000000*/     NOP CC.T;
    /*0078*/     /*0x00001de440000000*/     NOP CC.T;
        ........................

Вы можете видеть, что fatbinary содержит микрокод из собранного PTX. С подготовленной функцией устройства fatbin вы можете сделать что-то вроде этого в коде CUDA C:

extern "C" __device__ void inc_ptr(int* &ptr, const int inc);

__global__
void memsetkernel(int *inout, const int val, const int N)
{
    int stride = blockDim.x * gridDim.x;
    int *p = inout;
    inc_ptr(p, threadIdx.x + blockDim.x*blockIdx.x);

    for(; p < inout+N; inc_ptr(p, stride)) *p = val;
}  


int main(void)
{
    const int n=10;
    int *p;
    cudaMalloc((void**)&p, sizeof(int)*size_t(n));
    memsetkernel<<<1,32>>>(p, 5, n);

    return 0;
}

В отдельном режиме компиляции инструментальная цепочка кода устройства будет учитыватьextern объявление и (если вы получаете контроль над символами), функция устройства fatbinary может быть связана с другим устройством и кодом хоста для создания конечного объекта:

$ nvcc -arch=sm_30 -Xptxas="-v" -dlink -o memset.out inc_ptr.fatbin memset_kernel.cu 

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z12memsetkernelPiii' for 'sm_30'
ptxas info    : Function properties for _Z12memsetkernelPiii
    8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 20 registers, 332 bytes cmem[0]

$ cuobjdump -sass memset.out 

Fatbin elf code:
================
arch = sm_30
code version = [1,6]
producer = <unknown>
host = mac
compile_size = 32bit
identifier = inc_ptr.fatbin memset_kernel.cu 

    code for sm_30
        Function : _Z12memsetkernelPiii
    /*0008*/     /*0x10005de428004001*/     MOV R1, c [0x0] [0x44];
    /*0010*/     /*0x20105d034800c000*/     IADD R1, R1, -0x8;
    /*0018*/     /*0x00019de428004005*/     MOV R6, c [0x0] [0x140];
    /*0020*/     /*0x10101c034800c000*/     IADD R0, R1, 0x4;
    /*0028*/     /*0x8400dc042c000000*/     S2R R3, SR_Tid_X;
    /*0030*/     /*0x90041c0348004000*/     IADD R16, R0, c [0x0] [0x24];
    /*0038*/     /*0x94001c042c000000*/     S2R R0, SR_CTAid_X;
    /*0048*/     /*0xd0009de428004000*/     MOV R2, c [0x0] [0x34];
    /*0050*/     /*0x91045d0348004000*/     IADD R17, R16, -c [0x0] [0x24];
    /*0058*/     /*0x40011de428000000*/     MOV R4, R16;
    /*0060*/     /*0xa0015ca320064000*/     IMAD R5, R0, c [0x0] [0x28], R3;
    /*0068*/     /*0x01119c85c8000000*/     STL [R17], R6;
    /*0070*/     /*0xa0209ca350004000*/     IMUL R2, R2, c [0x0] [0x28];
    /*0078*/     /*0x0001000710000000*/     JCAL 0x0;
    /*0088*/     /*0x0110dc85c0000000*/     LDL R3, [R17];
    /*0090*/     /*0x20001de428004005*/     MOV R0, c [0x0] [0x148];
    /*0098*/     /*0x00049c4340004005*/     ISCADD R18, R0, c [0x0] [0x140], 0x2;
    /*00a0*/     /*0x4831dc031b0e0000*/     ISETP.GE.U32.AND P0, pt, R3, R18, pt;
    /*00a8*/     /*0x000001e780000000*/     @P0 EXIT;
    /*00b0*/     /*0x1004dde428004005*/     MOV R19, c [0x0] [0x144];
    /*00b8*/     /*0x0034dc8590000000*/     ST [R3], R19;
    /*00c8*/     /*0x40011de428000000*/     MOV R4, R16;
    /*00d0*/     /*0x08015de428000000*/     MOV R5, R2;
    /*00d8*/     /*0x0001000710000000*/     JCAL 0x0;
    /*00e0*/     /*0x0110dc85c0000000*/     LDL R3, [R17];
    /*00e8*/     /*0x4831dc03188e0000*/     ISETP.LT.U32.AND P0, pt, R3, R18, pt;
    /*00f0*/     /*0x000001e74003ffff*/     @P0 BRA 0xb8;
    /*00f8*/     /*0x00001de780000000*/     EXIT;
    /*0100*/     /*0xe0001de74003ffff*/     BRA 0x100;
    /*0108*/     /*0x00001de440000000*/     NOP CC.T;
    /*0110*/     /*0x00001de440000000*/     NOP CC.T;
    /*0118*/     /*0x00001de440000000*/     NOP CC.T;
    /*0120*/     /*0x00001de440000000*/     NOP CC.T;
    /*0128*/     /*0x00001de440000000*/     NOP CC.T;
    /*0130*/     /*0x00001de440000000*/     NOP CC.T;
    /*0138*/     /*0x00001de440000000*/     NOP CC.T;
        .....................................


        Function : inc_ptr
    /*0008*/     /*0x0040dc8580000000*/     LD R3, [R4];
    /*0010*/     /*0x00301c8580000000*/     LD R0, [R3];
    /*0018*/     /*0x14001c0348000000*/     IADD R0, R0, R5;
    /*0020*/     /*0x00301c8590000000*/     ST [R3], R0;
    /*0028*/     /*0x00001de790000000*/     RET;
    /*0030*/     /*0x00001de440000000*/     NOP CC.T;
    /*0038*/     /*0x00001de440000000*/     NOP CC.T;
    /*0040*/     /*0xe0001de74003ffff*/     BRA 0x40;
    /*0048*/     /*0x00001de440000000*/     NOP CC.T;
    /*0050*/     /*0x00001de440000000*/     NOP CC.T;
    /*0058*/     /*0x00001de440000000*/     NOP CC.T;
    /*0060*/     /*0x00001de440000000*/     NOP CC.T;
    /*0068*/     /*0x00001de440000000*/     NOP CC.T;
    /*0070*/     /*0x00001de440000000*/     NOP CC.T;
    /*0078*/     /*0x00001de440000000*/     NOP CC.T;
        ........................

Для достижения этой цели могут использоваться другие приемы, которые можно использовать с помощью цепочки инструментов, но такой подход, безусловно, работает.

Добавлен ответ, чтобы убрать этот вопрос из списка вопросов без ответа, если кто-то будет так любезен, чтобы выразить его и / или принять.
Красиво сделано, гений!
ОК, один вопрос. Как мне создать исполняемый файл, который включает в себя объект, связанный с устройством (memset.out)?
-1

CUDA C поддерживает asm, есть документ, который находится в каталоге doc после установки набора инструментов cuda.

Да, CUDA C поддерживает asm, и, как я уже писал, я знаю, что могу встроить ассемблерный код в CUDA C, но я не знаю, как встроить функцию сборки ptx (.func). fursund

Похожие вопросы