__device__ __forceinline__ bhalf8 to_bhalf()

in src/ew_op_gpu.h [291:329]

• 39 lines of code
• 1 McCabe index (conditional complexity)

__device__ __forceinline__ bhalf8 to_bhalf(float8 v)
{
    bhalf8 r;
    asm("{\n\t"
        ".reg .f32 exp<8>, f<8>;\n\t"
        ".reg .u32 u<8>;\n\t"
        "and.b32 exp0, %4 , 0xff800000;\n\t"
        "and.b32 exp1, %5 , 0xff800000;\n\t"
        "and.b32 exp2, %6 , 0xff800000;\n\t"
        "and.b32 exp3, %7 , 0xff800000;\n\t"
        "and.b32 exp4, %8 , 0xff800000;\n\t"
        "and.b32 exp5, %9 , 0xff800000;\n\t"
        "and.b32 exp6, %10, 0xff800000;\n\t"
        "and.b32 exp7, %11, 0xff800000;\n\t"
        "fma.rz.ftz.f32 f0, exp0, 0F3b800000, %4 ;\n\t"  // "0F%08x" % ((127 - 7 - 1) << 23)
        "fma.rz.ftz.f32 f1, exp1, 0F3b800000, %5 ;\n\t"
        "fma.rz.ftz.f32 f2, exp2, 0F3b800000, %6 ;\n\t"
        "fma.rz.ftz.f32 f3, exp3, 0F3b800000, %7 ;\n\t"
        "fma.rz.ftz.f32 f4, exp4, 0F3b800000, %8 ;\n\t"
        "fma.rz.ftz.f32 f5, exp5, 0F3b800000, %9 ;\n\t"
        "fma.rz.ftz.f32 f6, exp6, 0F3b800000, %10;\n\t"
        "fma.rz.ftz.f32 f7, exp7, 0F3b800000, %11;\n\t"
        "mov.b32 u0, f0;\n\t"
        "mov.b32 u1, f1;\n\t"
        "mov.b32 u2, f2;\n\t"
        "mov.b32 u3, f3;\n\t"
        "mov.b32 u4, f4;\n\t"
        "mov.b32 u5, f5;\n\t"
        "mov.b32 u6, f6;\n\t"
        "mov.b32 u7, f7;\n\t"
        "vadd.u32.u32.u32 %0.h0, u0.h1, 0, u1;\n\t" // use 16 bit merge functionality of vadd
        "vadd.u32.u32.u32 %1.h0, u2.h1, 0, u3;\n\t"
        "vadd.u32.u32.u32 %2.h0, u4.h1, 0, u5;\n\t"
        "vadd.u32.u32.u32 %3.h0, u6.h1, 0, u7;\n\t"
        "}" : "=r"(r.x), "=r"(r.y), "=r"(r.z), "=r"(r.w)
            :  "f"(v.a.x),"f"(v.a.y),"f"(v.a.z),"f"(v.a.w),
               "f"(v.b.x),"f"(v.b.y),"f"(v.b.z),"f"(v.b.w));
    return r;
}