   int count,
   cudaStream_t stream)
 {
-  dim3 dim_block(128, 1);  // 4 warps per stream, 1 stream per block
+  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,

 * @param[out] outputs Compression status per block
 * @param[in] count Number of blocks to compress
 **/
extern "C" __global__ void __launch_bounds__(128)
snap_kernel(
  const void* const* __restrict__ device_in_ptr,
  const uint64_t* __restrict__ device_in_bytes,

  cudaStream_t stream)
{
  uint32_t count32 = (count > 0) ? count : 0;
  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>>>(
