 #ifdef ENABLE_BITCOMP
 #include <bitcomp.h>
 
+#define BTCHK(call)                                                            \
+  {                                                                            \
+    bitcompResult_t err = call;                                                \
+    if (BITCOMP_SUCCESS != err) {                                              \
+      return nvcompErrorInternal;                                              \
+    }                                                                          \
+  }
+
 nvcompStatus_t nvcompBatchedBitcompCompressGetMaxOutputChunkSize(
     size_t max_chunk_size,
     nvcompBitcompFormatOpts format_opts,

    if (format_opts)
        algo = static_cast<bitcompAlgorithm_t>(format_opts->algorithm_type);
    bitcompHandle_t plan;
    if (bitcompCreateBatchPlan(&plan, batch_size, dataType, BITCOMP_LOSSLESS, algo) != BITCOMP_SUCCESS ||
        bitcompSetStream(plan, stream) != BITCOMP_SUCCESS)
    return nvcompErrorInternal;

    // Launch the Bitcomp async batch compression
    if (bitcompBatchCompressLossless(
            plan,
            device_uncompressed_ptrs,
            device_compressed_ptrs,
            device_uncompressed_bytes,
            device_compressed_bytes) != BITCOMP_SUCCESS)
        return nvcompErrorInternal;

    // Once launched, the handle can be destroyed
    if (bitcompDestroyPlan (plan) != BITCOMP_SUCCESS)
        return nvcompErrorInternal;
    
    return nvcompSuccess;
}

    if (format_opts)
        algo = static_cast<bitcompAlgorithm_t>(format_opts->algorithm_type);
    bitcompHandle_t plan;
    BTCHK(bitcompCreateBatchPlan(&plan, batch_size, dataType, BITCOMP_LOSSLESS, algo));
    BTCHK(bitcompSetStream(plan, stream));

    // Launch the Bitcomp async batch decompression with extra checks
    BTCHK(bitcompBatchUncompressCheck(
        plan,
        device_compressed_ptrs,
        device_uncompressed_ptrs,
        device_uncompressed_bytes,
        (bitcompResult_t*)device_statuses));

    // Need a separate kernel to query the actual uncompressed size,
    // as bitcomp doesn't write the uncompressed size during decompression
    BTCHK(bitcompBatchGetUncompressedSizesAsync(
        device_compressed_ptrs,
        device_actual_uncompressed_bytes,
        batch_size,
        stream));

    // Also launch a kernel to convert the output statuses
    const int threads = 512;
    int blocks = (batch_size - 1) / threads + 1;
    convertOutputStatuses <<<blocks,threads,0,stream>>> (device_statuses, batch_size);

    // Once launched, the handle can be destroyed
    BTCHK(bitcompDestroyPlan(plan));
    return nvcompSuccess;
}

