         int             offset)             ///< [in] Up-offset to pull from
     {
         long long output;
-        long long temp = input;
 
         // Use predicate set from SHFL to guard against invalid peers
         asm(
             "{"
             "  .reg .s64 r0;"
-            "  .reg .u32 lo0;"
-            "  .reg .u32 hi0;"
-            "  .reg .u32 lo1;"
-            "  .reg .u32 hi1;"
+            "  .reg .u32 lo;"
+            "  .reg .u32 hi;"
             "  .reg .pred p;"
-            "  mov.b64 {lo1, hi1}, %1;"
-            "  shfl.up.b32 lo0|p, lo1, %2, %3;"
-            "  shfl.up.b32 hi0, hi1, %2, %3;"
-            "  mov.b64 r0, {lo0, hi0};"
-            "  @p add.s64 %1, r0, %1;"
-            "  mov.s64 %0, %1;"
+            "  mov.b64 {lo, hi}, %1;"
+            "  shfl.up.b32 lo|p, lo, %2, %3;"
+            "  shfl.up.b32 hi|p, hi, %2, %3;"
+            "  mov.b64 r0, {lo, hi};"
+            "  @p add.s64 r0, r0, %4;"
+            "  mov.s64 %0, r0;"
             "}"
-            : "=l"(output) : "l"(temp), "r"(offset), "r"(first_lane));
+            : "=l"(output) : "l"(input), "r"(offset), "r"(first_lane), "l"(input));
 
         return output;
     }

        int             offset)             ///< [in] Up-offset to pull from
    {
        unsigned long long output;
        unsigned long long temp = input;

        // Use predicate set from SHFL to guard against invalid peers
        asm(
            "{"
            "  .reg .u64 r0;"
            "  .reg .u32 lo0;"
            "  .reg .u32 hi0;"
            "  .reg .u32 lo1;"
            "  .reg .u32 hi1;"
            "  .reg .pred p;"
            "  mov.b64 {lo1, hi1}, %1;"
            "  shfl.up.b32 lo0|p, lo1, %2, %3;"
            "  shfl.up.b32 hi0, hi1, %2, %3;"
            "  mov.b64 r0, {lo0, hi0};"
            "  @p add.u64 %1, r0, %1;"
            "  mov.u64 %0, %1;"
            "}"
            : "=l"(output) : "l"(temp), "r"(offset), "r"(first_lane));

        return output;
    }


        // Iterate scan steps
        InclusiveScanStep(output, scan_op, SHFL_C, Int2Type<0>());
/*
        // Iterate scan steps
        #pragma unroll
        for (int STEP = 0; STEP < STEPS; STEP++)
