   int i, ioff, ioffmv, ip, j, k, c, tmp, cmin, imin;
   int zero = 0;
   int sid = threadIdx.x + blockDim.x * threadIdx.y;
+
+  // Load data into shared memory
   for (i = 0; i < TLEN/1024; i++) {
     sa[sid * i*1024] = a[sid + i*1024 + TLEN*blockIdx.x];
   }

    bb[j] = sb[threadIdx.y][j + ioffmv];
    ww[j] = sw[threadIdx.y][j + ioffmv];
  }
  for (j = 0; j < TLEN*NVEC/32; j++) {
    tmp = VECLEN + ioff + j / 4;
    if (tmp - ioffmv < TLEN - VECLEN * NVEC) {
      if (j % 4 == 0) {
        aa[VECLEN] = sa[tmp];
      }
      c = 0;
#pragma unroll
      for (k = 0; k < VECLEN; k++) {    
        hammingcell(aa[k], aa[k+1], bb[k], ww[k], c, tmp, zero);
      }
      hammingend(aa[k]);
#pragma unroll
      for (k = 1; k < NVEC; k *= 2) {    
        c = c + __shfl_down(c, k);
      }
      if (c < cmin) {
        cmin = c;
        imin = ioff + j;
      }
    }
  }
  for (k = NVEC; k < 32; k *= 2) {    
    tmp = __shfl_down(cmin, k);
    if (tmp < cmin) {

  ioff = ioffmv + ip * (TLEN*NVEC/32);
  cmin = 0x7fffffff;
  imin = -1;

  // Load data for this thread into registers
#pragma unroll
  for (j = 0; j < VECLEN; j++) {
    tmp = j + ioff;
