             // vertex for out-of-bound items, but zero-valued
             LoadWarpStriped<LOAD_DEFAULT>(d_columns + block_offset, guarded_items, VertexId(0), columns);
             LoadWarpStriped<LOAD_DEFAULT>(d_values + block_offset, guarded_items, Value(0), values);
-            LoadWarpStriped<LOAD_DEFAULT>(d_rows + block_offset, guarded_items, last_block_row, rows);
+            LoadWarpStriped<LOAD_DEFAULT>(d_rows + block_offset, guarded_items, temp_storage.last_block_row, rows);
         }
         else
         {

        if (threadIdx.x == 0)
        {
            VertexId first_block_row            = d_rows[block_offset];
            last_block_row                      = d_rows[block_oob - 1];

            // Initialize carry to identity
            carry.running_prefix.row            = first_block_row;
            carry.running_prefix.partial        = Value(0);
            temp_storage.identity               = carry.running_prefix;

            temp_storage.first_scatter.row      = last_block_row;
            temp_storage.first_scatter.partial  = Value(0);
            temp_storage.last_block_row         = last_block_row;
        }

        __syncthreads();

        last_block_row = temp_storage.last_block_row;
    }



                typename BlockScan::TempStorage             scan;               // Smem needed for reduce-value-by-row scan
                typename BlockDiscontinuity::TempStorage    discontinuity;      // Smem needed for head-flagging
            };
        };

        VertexId        first_block_row;
        VertexId        last_block_row;

        PartialSum      identity;
        Value           first_scatter_partial;
    };

    //---------------------------------------------------------------------
