ptxas жалуется на (печатает) мою печальную функцию устройства - PullRequest
Новая
       50

ptxas жалуется на (печатает) мою печальную функцию устройства

0 голосов
/ 16 февраля 2020

Рассмотрим следующий код PTX: 

//
// Generated by NVIDIA NVVM Compiler... sort of
//
// Compiler Build ID: CL-25769353
// Cuda compilation tools, release 10.1, V10.1.105
// Based on LLVM 3.4svn
//

.version 6.4
.target sm_30
.address_size 64

.func  (.param .b32 func_retval0) foo(
        .param .b32 foo_param_0,
        .param .b32 foo_param_1,
        .param .b32 foo_param_2
)
{
        .reg .b16       %rs<3>;
        .reg .b32       %r<3>;


        ld.param.u16    %rs1, [foo_param_0];
        ld.param.u16    %rs2, [foo_param_1];
        ld.param.u32    %r2, [foo_param_2];
        // inline asm
        sad.s16 %r1, %rs1, %rs2, %r2;
        // inline asm
        st.param.b32    [func_retval0+0], %r1;
        ret;
}

При попытке скомпилировать его с помощью ptxas (CUDA 10.1) я получаю: 

ptxas /tmp/a.ptx, line 27; error   : Arguments mismatch for instruction 'sad'
ptxas fatal   : Ptx assembly aborted due to errors

Почему это так? Что не так с этой комбинацией типов?

Ссылка PTX говорит: 

sad.type  d, a, b, c;

.type = { .u16, .u32, .u64,
          .s16, .s32, .s64 };

и кажется d и c u32 всегда и type относится к a и b. Вот что в любом случае есть у __sad() функций в device_functions.h.

1 Ответ

1 голос
/ 16 февраля 2020

На самом деле тип d и a должен быть "беззнаковой версией типа b и c". Итак, это должно работать: 

.func  (.param .b32 func_retval0) foo(
        .param .b32 foo_param_0,
        .param .b32 foo_param_1,
        .param .b32 foo_param_2
)
{
        .reg .b32       %r<5>;


        ld.param.u32    %r4, [foo_param_2];
        ld.param.s16    %r2, [foo_param_0];
        ld.param.s16    %r3, [foo_param_1];
        // inline asm
        sad.s32 %r1, %r2, %r3, %r4;
        // inline asm
        st.param.b32    [func_retval0+0], %r1;
        ret;
}
...