 
     long long int serial = 0;
 
+    const uint32_t m = 4;
+    uint16_t max_cols = max(read_count, (((read_count - 1) / (blockDim.x * m)) + 1) * (blockDim.x * m));
+    //if (thread_idx == 0)
+    //{
+    //    printf("max cols %d\n", max_cols);
+    //    //printf("row %d\n", i);
+    //}
     // Run DP loop for calculating scores. Process each row at a time, and
     // compute vertical and diagonal values in parallel.
     for(uint16_t graph_pos = 0; graph_pos < graph_count; graph_pos++)

            // 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;
            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);
                //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;
                //}
                loop = loop >> 1;
                //printf("thread %d line %d\n", thread_idx, __LINE__);
            }


            //    prev_score[0] = score3;
            //}

            init_score = __shfl_sync(0xffffffff, score3, 31);

            //if (thread_idx == 0)
            //{
