   int mask = 0;
 
   // full search
+  assert(blockDim.x == 32);
   for (int d = 1; d < 32; ++d) {
     const int nbr_id = (threadIdx.x + d) & 31;
     mask |= (val == __shfl_sync(participants, val, nbr_id)) << nbr_id;

      const int match_mask_warp = warpBallot(
          match_mask_self && __clz(__brev(match_mask_self)) != threadIdx.x);

      if (match_mask_warp) {
        // find the byte offset (thread id) within the warp where the first
        // match is located

        assert(candidate_first_match_thread != threadIdx.x || match_found);
        assert(!match_found || candidate_first_match_thread <= threadIdx.x);

        if (candidate_first_match_thread < first_match_thread) {
          // if we found a valid match, and it occurs before a previously found
          // match, use that
          first_match_thread = candidate_first_match_thread;
          hashPos = __shfl_sync(0xffffffff, hashPos, first_match_thread);
          match_location
              = convertIdx(hashTable[hashPos], decomp_idx + first_match_thread);
        }
      }

