             // The any_sync warp primitive lets us easily check if any of the threads had an update.
             //bool loop = true;
             uint32_t loop = 0xffffffff >> 1;
+            uint16_t thread_id_to_process = 1;
             while(loop)
             {
                 //loop = false;
                 // The shfl_up lets us grab a value from the lane below.
                 //printf("thread %d line %d\n", thread_idx, __LINE__);
-                int16_t last_score = __shfl_up_sync(0xffffffff << 1, score3, 1);
+                //int16_t last_score = __shfl_up_sync(0xffffffff << 1, score3, 1);
+                //int16_t last_score = __shfl_up_sync(0xffffffff << 1, score1, 1);
+                int16_t last_score = __shfl_up_sync(0x1 << thread_id_to_process, score3, 1);
                 //printf("thread %d line %d\n", thread_idx, __LINE__);
-                score0 = max(last_score + GAP, score0);
-                score1 = max(score0 + GAP, score1);
-                score2 = max(score1 + GAP, score2);
-                score3 = max(score2 + GAP, score3);
-                //if (thread_idx == 1)
-                //{
-                //    printf("thread idx %d scores %d, %d, %d, %d\n", thread_idx, score0, score1, score2, score3);
-                //}
-                //if (thread_idx == 2)
-                //{
-                //    printf("thread idx %d scores %d, %d, %d, %d\n", thread_idx, score0, score1, score2, score3);
-                //}
-                //int16_t new_score = max(__shfl_up_sync(0xffffffff << 1, score, 1) + GAP, score);
-                //if (new_score > score)
-                //{
-                //    loop = true;
-                //    score = new_score;
-                //}
+                if (thread_idx == thread_id_to_process)
+                {
+                    score0 = max(last_score + GAP, score0);
+                    score1 = max(score0 + GAP, score1);
+                    score2 = max(score1 + GAP, score2);
+                    score3 = max(score2 + GAP, score3);
+                }
+
                 loop = loop >> 1;
+                thread_id_to_process++;
                 //printf("thread %d line %d\n", thread_idx, __LINE__);
             }
 

            uint16_t j1 = read_pos + 2;
            uint16_t j2 = read_pos + 3;
            uint16_t j3 = read_pos + 4;
            //printf("thread idx %d locations %d %d %d %d\n", thread_idx, j0, j1, j2, j3);
            int16_t score0 = max(scores_pred_i_1[j0-1] + char_profile0,
                    scores_pred_i_1[j0] + GAP);
            int16_t score1 = max(scores_pred_i_1[j1-1] + char_profile1,

        uint8_t n = nodes[node_id];

        //int16_t prev_score = init_score;
        //if (thread_idx == 0)
        //{
        //    prev_score[0] = init_score;
        //}

        //__syncthreads();

        // max_cols is the first warp boundary multiple beyond read_count. This is done
        // so all threads in the warp enter the loop.
        for(uint16_t read_pos = thread_idx * m; read_pos < max_cols; read_pos += blockDim.x * m)
        {
            //printf("updating vertical for pos %d thread %d\n", read_pos, thread_idx);
            int16_t char_profile0 = (n == read[read_pos + 0] ? MATCH : MISMATCH);
