Я пытаюсь сделать простой код, который будет добавлять 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, он везде равен нулю. Не могли бы вы сказать мне, почему и где проблема в коде?
////////////// ДОБАВЛЯТЬ
После компиляции моего кода 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 очищает мой код? Что не так ?
=l
). Поскольку вы читаете значение, возможно,+l
? - person David Wohlferd   schedule 23.06.2019asm volatile
. В нынешнем виде ваш ассемблерный код объявляется неэффективным, поэтому компилятор может его оптимизировать. - person tera   schedule 23.06.2019