   cudaStream_t stream)
 {
   uint32_t count32 = (count > 0) ? count : 0;
-  dim3 dim_block(96, 1);     // 4 warps per stream, 1 stream per block
+  dim3 dim_block(96, 1);     // 3 warps per stream, 1 stream per block
   dim3 dim_grid(count32, 1);  // TODO: Check max grid dimensions vs max expected count
 
   unsnap_kernel<<<dim_grid, dim_block, 0, stream>>>(

#define BATCH_COUNT (1 << LOG2_BATCH_COUNT)
#define LOG2_PREFETCH_SIZE 10
#define PREFETCH_SIZE (1 << LOG2_PREFETCH_SIZE)  // 512B, in 32B chunks
#define PREFETCH_SECTORS 8
#define LITERAL_SECTORS 8

#define LOG_CYCLECOUNT 0


  int count,
  cudaStream_t stream)
{
  dim3 dim_block(64, 1);  // 2 warps per stream, 1 stream per block
  dim3 dim_grid(count, 1);
  if (count > 0) { snap_kernel<<<dim_grid, dim_block, 0, stream>>>(
    device_in_ptr, device_in_bytes, device_out_ptr, device_out_available_bytes,
