some question about global_load

[hlq1025]

@kerrmudgeon the comment says that the redundant mov PTX instruction is used to enforce the compiler to initialize data to zero before ld.global. but the code don't init the data to zero. here is the code below.
template
struct global_load<AccessType,
32
> {
CUTLASS_DEVICE
global_load(AccessType &D, void const *ptr, bool pred_guard) {
uint4 *data = reinterpret_cast<uint4 *>(&D);

asm volatile(
    "{\n"
    "  .reg .pred p;\n"
    "  setp.ne.b32 p, %9, 0;\n"
    "  mov.b32 %0, %10;\n"
    "  mov.b32 %1, %11;\n"
    "  mov.b32 %2, %12;\n"
    "  mov.b32 %3, %13;\n"
    "  mov.b32 %4, %14;\n"
    "  mov.b32 %5, %15;\n"
    "  mov.b32 %6, %16;\n"
    "  mov.b32 %7, %17;\n"

#if CUTLASS_ENABLE_L2_PREFETCH
" @p ld.global.L2::128B.v4.u32 {%0, %1, %2, %3}, [%8];\n"
" @p ld.global.L2::128B.v4.u32 {%4, %5, %6, %7}, [%18];\n"
#else
" @p ld.global.v4.u32 {%0, %1, %2, %3}, [%8];\n"
" @p ld.global.v4.u32 {%4, %5, %6, %7}, [%18];\n"
#endif
"}\n"
: "=r"(data[0].x), "=r"(data[0].y), "=r"(data[0].z), "=r"(data[0].w),
"=r"(data[1].x), "=r"(data[1].y), "=r"(data[1].z), "=r"(data[1].w)
: "l"(ptr), "r"((int)pred_guard), "r"(data[0].x), "r"(data[0].y),
"r"(data[0].z), "r"(data[0].w), "r"(data[1].x), "r"(data[1].y),
"r"(data[1].z), "r"(data[1].w), "l"(((uint8_t *)ptr) + 16));
}
};
could you explain it in detail? Thanks a lot.

[hwu36]

the comment is actually not accurate. it helps C++ to initialize the data properly. Not always forcing it to be 0. It is just a trick to fool the compiler.

[hlq1025]

@hwu36 the mov ptx just assign the self value to itself. why not delete the mov ptx instructions. could you explain the trick in detail.

[hwu36]

it basically tells the compiler that the reg is not only written, but also read. so the compiler would not delete any earlier assignment.

[hlq1025]

it basically tells the compiler that the reg is not only written, but also read. so the compiler would not delete any earlier assignment.

@hwu36 what is the benefit of tell the compiler that the reg is also read?

[hwu36]

x = 0;
x = load(addr, cond);

the compiler can delete x=0 if it thinks load is write only.