Можно ли поместить инструкции по сборке в код CUDA? - PullRequest
Новая
       43

Можно ли поместить инструкции по сборке в код CUDA?

7 голосов
/ 09 сентября 2010

Я хочу использовать ассемблерный код в коде CUDA C, чтобы сократить дорогостоящие исполнения, как мы используем asm в программировании на c.

Возможно ли это?

Ответы [ 2 ]

18 голосов
/ 16 августа 2011

Начиная с CUDA 4.0, встроенный PTX поддерживается набором инструментов CUDA. В наборе инструментов есть документ, который описывает его: Using_Inline_PTX_Assembly_In_CUDA.pdf

Ниже приведен код, демонстрирующий использование встроенного PTX в CUDA 4.0. Обратите внимание, что этот код не должен использоваться в качестве замены встроенной функции CUDA __clz (), я просто написал его для изучения аспектов новой встроенной функции PTX. 

__device__ __forceinline__ int my_clz (unsigned int x)
{
    int res;

    asm ("{\n"
         "        .reg .pred iszero, gezero;\n"
         "        .reg .u32 t1, t2;\n"
         "        mov.b32         t1, %1;\n"
         "        shr.u32         %0, t1, 16;\n"
         "        setp.eq.b32     iszero, %0, 0;\n"
         "        mov.b32         %0, 0;\n"
         "@iszero shl.b32         t1, t1, 16;\n"
         "@iszero or.b32          %0, %0, 16;\n"
         "        and.b32         t2, t1, 0xff000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 8;\n"
         "@iszero or.b32          %0, %0, 8;\n"
         "        and.b32         t2, t1, 0xf0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 4;\n"
         "@iszero or.b32          %0, %0, 4;\n"
         "        and.b32         t2, t1, 0xc0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 2;\n"
         "@iszero or.b32          %0, %0, 2;\n"
         "        setp.ge.s32     gezero, t1, 0;\n"
         "        setp.eq.b32     iszero, t1, 0;\n"
         "@gezero or.b32          %0, %0, 1;\n"
         "@iszero add.u32         %0, %0, 1;\n\t"
         "}"
         : "=r"(res)
         : "r"(x));
    return res;
}
4 голосов
/ 09 сентября 2010

Нет, вы не можете, нет ничего похожего на конструкции asm из C / C ++.Что вы можете сделать, это настроить сгенерированную сборку PTX, а затем использовать ее с CUDA.

См. this для примера.

Но для графических процессоров оптимизация сборки не требуется, вы должны сначала выполнить другие оптимизации, такие как объединение памяти и заполнение.См. Руководство по рекомендациям CUDA для получения дополнительной информации.

...