Простое добавление векторов в Inline PTX CUDA - PullRequest
Новая
       46

Простое добавление векторов в Inline PTX CUDA

0 голосов
/ 23 июня 2019

Пытаюсь сделать простой код, который добавит V1 (вектор) с V2 и сохранит значение в V3. Это работа в CUDA, но я не могу написать это в PTX, кто-то может помочь? 

__global__ void addKernelPTXv4(float4  *cc, const float4 *aa, const float4 *bb)
{

    int i = threadIdx.x;

    cc[i].x = aa[i].x + bb[i].x;
    cc[i].y = aa[i].y + bb[i].y;
    cc[i].z = aa[i].z + bb[i].z;
    cc[i].w = aa[i].w + bb[i].w;
}

__global__ void addKernelPTXv4(float4  *cc, const float4 *aa, const float4 *bb)
{
    asm( 
        ".reg.f32  a1, a2, a3, a4;"
        ".reg.f32  b1, b2, b3, b4;"
        ".reg.f32  c1, c2, c3, c4;"
        ".reg.r32  r1;"
        ".reg.s64 rd1, rd2, rd3;"

        "mov.r32 r1, %tid.x;"
        "mul.s32 r1, r1, 16;"
        "add.s64 rd1, %0, r1;"
        "add.s64 rd2, %1, r1;"
        "add.s64 rd3, %2, r1;"

        "ld.global.v4.f32 { a1, a2, a3, a4 }, [rd2];"
        "ld.global.v4.f32 { b1, b2, b3, b4 }, [rd3];"
        "add.f32 c1, a1, b2;"
        "add.f32 c2, a2, b2;"
        "add.f32 c3, a3, b3;"
        "add.f32 c4, a4, b4;"
        "st.global.v4.f32  [rd1], { c1, c2, c3, c4 };"

        : "=l"(cc)
        : "l"(aa), "l"(bb) 
    );
}

Когда я печатаю Vector CC, после запуска кода PTX он везде равен нулю. Не могли бы вы сказать мне, почему и где проблема в коде?

////////////// ADD

После компиляции моего кода ptx / asm в Visual Studio 2015 я проверяю выходной файл PTX: 

    // .globl   _Z14addKernelPTXv4P6float4PKS_S2_
.visible .entry _Z14addKernelPTXv4P6float4PKS_S2_(
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_0,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_1,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_2
)
{
    .reg .b64   %rd<5>;


    ret;
}

Это .... ПУСТОЙ? Так в чем же проблема?

Когда я компилирую версию CUDA, выходной PTX будет: 

    // .globl   _Z14addKernelPTXv4P6float4PKS_S2_
.visible .entry _Z14addKernelPTXv4P6float4PKS_S2_(
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_0,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_1,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_2
)
{
    .reg .f32   %f<21>;
    .reg .b32   %r<2>;
    .reg .b64   %rd<11>;


    ld.param.u64    %rd1, [_Z14addKernelPTXv4P6float4PKS_S2__param_0];
    ld.param.u64    %rd2, [_Z14addKernelPTXv4P6float4PKS_S2__param_1];
    ld.param.u64    %rd3, [_Z14addKernelPTXv4P6float4PKS_S2__param_2];
    cvta.to.global.u64  %rd4, %rd1;
    cvta.to.global.u64  %rd5, %rd3;
    cvta.to.global.u64  %rd6, %rd2;
    mov.u32     %r1, %tid.x;
    mul.wide.s32    %rd7, %r1, 16;
    add.s64     %rd8, %rd6, %rd7;
    add.s64     %rd9, %rd5, %rd7;
    add.s64     %rd10, %rd4, %rd7;
    ld.global.v4.f32    {%f1, %f2, %f3, %f4}, [%rd8];
    ld.global.v4.f32    {%f5, %f6, %f7, %f8}, [%rd9];
    add.f32     %f11, %f4, %f8;
    add.f32     %f14, %f3, %f7;
    add.f32     %f17, %f2, %f6;
    add.f32     %f20, %f1, %f5;
    st.global.v4.f32    [%rd10], {%f20, %f17, %f14, %f11};
    ret;
}

Почему VS2015 очищает мой код? Что не так?

1 Ответ

1 голос
/ 23 июня 2019

Ее код разрешения: 

__global__ void addKernelPTXv4(float4  *ccc, const float4 *aaa, const float4 *bbb)
{
    asm volatile ("{    \n\t"
        ".reg.f32  aa<4>, bb<4>, cc<4>;    \n\t"
        ".reg.s32  rr0;    \n\t"
        ".reg.s64 rrd<4>;    \n\t"

        "mov.s32 rr0, %tid.x;    \n\t"
        "mul.wide.s32 rrd0, rr0, 16;    \n\t"

        "add.s64 rrd1, %0, rrd0;    \n\t"
        "add.s64 rrd2, %1, rrd0;    \n\t"
        "add.s64 rrd3, %2, rrd0;    \n\t"

        "ld.global.v4.f32 { aa0, aa1, aa2, aa3 }, [rrd2];    \n\t"
        "ld.global.v4.f32 { bb0, bb1, bb2, bb3 }, [rrd3];    \n\t"

        "add.f32 cc0, aa0, bb0;    \n\t"
        "add.f32 cc1, aa1, bb1;    \n\t"
        "add.f32 cc2, aa2, bb2;    \n\t"
        "add.f32 cc3, aa3, bb3;    \n\t"

        "st.global.v4.f32  [rrd1], { cc0, cc1, cc2, cc3 };    \n\t"
        "}    \n\t"
        :
        : "l"(ccc), "l"(aaa), "l"(bbb)
        : "memory"
        );
}

После компиляции: 

    // .globl   _Z14addKernelPTXv4P6float4PKS_S2_
.visible .entry _Z14addKernelPTXv4P6float4PKS_S2_(
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_0,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_1,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_2
)
{
    .reg .b64   %rd<4>;


    ld.param.u64    %rd1, [_Z14addKernelPTXv4P6float4PKS_S2__param_0];
    ld.param.u64    %rd2, [_Z14addKernelPTXv4P6float4PKS_S2__param_1];
    ld.param.u64    %rd3, [_Z14addKernelPTXv4P6float4PKS_S2__param_2];
    // inline asm
    {    
    .reg.f32  aa<4>, bb<4>, cc<4>;    
    .reg.s32  rr0;    
    .reg.s64 rrd<4>;    
    mov.s32 rr0, %tid.x;    
    mul.wide.s32 rrd0, rr0, 16;    
    add.s64 rrd1, %rd1, rrd0;    
    add.s64 rrd2, %rd2, rrd0;    
    add.s64 rrd3, %rd3, rrd0;    
    ld.global.v4.f32 { aa0, aa1, aa2, aa3 }, [rrd2];    
    ld.global.v4.f32 { bb0, bb1, bb2, bb3 }, [rrd3];    
    add.f32 cc0, aa0, bb0;    
    add.f32 cc1, aa1, bb1;    
    add.f32 cc2, aa2, bb2;    
    add.f32 cc3, aa3, bb3;    
    st.global.v4.f32  [rrd1], { cc0, cc1, cc2, cc3 };    
    }    

    // inline asm
    ret;
}

И оригинальный код CUDA: 

__global__ void addKernelPTXv4(float4  *c, const float4 *a, const float4 *b)
{
    int i = threadIdx.x;
    c[i].x = a[i].x + b[i].x;
    c[i].y = a[i].y + b[i].y;
    c[i].z = a[i].z + b[i].z;
    c[i].w = a[i].w + b[i].w;
}

После компиляции: 

    // .globl   _Z14addKernelPTXv4P6float4PKS_S2_
.visible .entry _Z14addKernelPTXv4P6float4PKS_S2_(
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_0,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_1,
    .param .u64 _Z14addKernelPTXv4P6float4PKS_S2__param_2
)
{
    .reg .f32   %f<21>;
    .reg .b32   %r<2>;
    .reg .b64   %rd<11>;


    ld.param.u64    %rd1, [_Z14addKernelPTXv4P6float4PKS_S2__param_0];
    ld.param.u64    %rd2, [_Z14addKernelPTXv4P6float4PKS_S2__param_1];
    ld.param.u64    %rd3, [_Z14addKernelPTXv4P6float4PKS_S2__param_2];
    cvta.to.global.u64  %rd4, %rd1;
    cvta.to.global.u64  %rd5, %rd3;
    cvta.to.global.u64  %rd6, %rd2;
    mov.u32     %r1, %tid.x;
    mul.wide.s32    %rd7, %r1, 16;
    add.s64     %rd8, %rd6, %rd7;
    add.s64     %rd9, %rd5, %rd7;
    add.s64     %rd10, %rd4, %rd7;
    ld.global.v4.f32    {%f1, %f2, %f3, %f4}, [%rd8];
    ld.global.v4.f32    {%f5, %f6, %f7, %f8}, [%rd9];
    add.f32     %f11, %f4, %f8;
    add.f32     %f14, %f3, %f7;
    add.f32     %f17, %f2, %f6;
    add.f32     %f20, %f1, %f5;
    st.global.v4.f32    [%rd10], {%f20, %f17, %f14, %f11};
    ret;
}

Итак, можно сделать чуть-чуть код PTX, чем VS.2015, скомпилированный из CUDA; -)

...