 using nvcomp::cascaded_decompress_threadblock_size;
 using nvcomp::partition_metadata_size;
 using nvcomp::compute_smem_size;
+using nvcomp::Check;
+using nvcomp::CudaUtils;
 
 namespace
 {

    nvcompStatus_t* device_statuses,
    cudaStream_t stream)
{

  // Just call kernel to perform compression. Macro for datatype happens within
  // kernel
  constexpr int threadblock_size = cascaded_decompress_threadblock_size;

  // call for all 4 possible sizes, all except the correct one will immediately
  // exit.

  // CHAR or UCHAR
  cascaded_decompression_kernel_type_check<1, size_t, threadblock_size>
      <<<batch_size, threadblock_size, 0, stream>>>(
          batch_size,
          device_compressed_ptrs,
          device_compressed_bytes,
          device_uncompressed_ptrs,
          device_uncompressed_bytes,
          device_actual_uncompressed_bytes,
          device_statuses);
  // SHORT or USHORT
  cascaded_decompression_kernel_type_check<2, size_t, threadblock_size>
      <<<batch_size, threadblock_size, 0, stream>>>(
          batch_size,
          device_compressed_ptrs,
          device_compressed_bytes,
          device_uncompressed_ptrs,
          device_uncompressed_bytes,
          device_actual_uncompressed_bytes,
          device_statuses);
  // INT or UINT
  cascaded_decompression_kernel_type_check<4, size_t, threadblock_size>
      <<<batch_size, threadblock_size, 0, stream>>>(
          batch_size,
          device_compressed_ptrs,
          device_compressed_bytes,
          device_uncompressed_ptrs,
          device_uncompressed_bytes,
          device_actual_uncompressed_bytes,
          device_statuses);
  // LONGLONG or ULONGLONG
  cascaded_decompression_kernel_type_check<8, size_t, threadblock_size>
      <<<batch_size, threadblock_size, 0, stream>>>(
          batch_size,
          device_compressed_ptrs,
          device_compressed_bytes,
          device_uncompressed_ptrs,
          device_uncompressed_bytes,
          device_actual_uncompressed_bytes,
          device_statuses);

  return nvcompSuccess;
}

    size_t batch_size,
    cudaStream_t stream)
{
  try {
    get_decompress_size_kernel<<<
        roundUpDiv(batch_size, cascaded_decompress_threadblock_size),
        cascaded_decompress_threadblock_size,
        0,
        stream>>>(
        device_compressed_ptrs,
        device_compressed_bytes,
        device_uncompressed_bytes,
        batch_size);
    CudaUtils::check_last_error();
  } catch (const std::exception& e) {
    return Check::exception_to_error(
        e, "nvcompBatchedCascadedGetDecompressSizeAsync()");
  }

  return nvcompSuccess;
}
