   }
 
   // Check for errors. Any error should be reported in the global status value
-  if (threadIdx.x == 0) {
+  if (cg_group.thread_rank() == 0) {
     if (decompressor.get_output_status() != nvcompSuccess) {
       *kernel_output_status = decompressor.get_output_status();
     }

    DecompressT& decompressor,
    GroupT&& cg_group)
{
  assert(chunks_per_block == blockDim.y);
  
  __shared__ uint32_t ix_chunks[chunks_per_block];

  volatile uint32_t& this_ix_chunk = *(ix_chunks + threadIdx.y);
  if (threadIdx.x == 0) {
    this_ix_chunk = blockIdx.x * chunks_per_block + threadIdx.y;
  }

  cg_group.sync();

        comp_chunk_sizes[this_ix_chunk],
        raw_chunk_size); 

    if (cg_group.thread_rank() == 0) {
      this_ix_chunk = initial_chunks + atomicAdd(ix_chunk, uint32_t{1});
    }

