 
     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));
+    const uint32_t cells_per_thread = 4;
+    uint16_t max_cols = max(read_count, (((read_count - 1) / (blockDim.x * cells_per_thread)) + 1) * (blockDim.x * cells_per_thread));
     //if (thread_idx == 0)
     //{
     //    printf("max cols %d\n", max_cols);


        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);

    //__shared__ int16_t score_i[1024];
    //__shared__ int16_t score_prev_i[1024];

    uint32_t thread_idx = threadIdx.x;

    long long int start = clock64();
