                          d_values.Current() + portion * PORTION_SIZE,
                          portion_num_items, current_bit, num_bits);
                     if (CubDebug(error = cudaPeekAtLastError())) break;
+                    if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
                 }
                 
                 // use the temporary buffers if no overwrite is allowed

                MaxPolicyT, IS_DESCENDING, KeyT, OffsetT>;
            if (CubDebug(error = cudaOccupancyMaxActiveBlocksPerMultiprocessor(
                &histo_blocks_per_sm, histogram_kernel, HISTO_BLOCK_THREADS, 0))) break;
            histogram_kernel<<<histo_blocks_per_sm * num_sms, HISTO_BLOCK_THREADS, 0, stream>>>
                (d_bins, d_keys.Current(), num_items, begin_bit, end_bit);
            if (CubDebug(error = cudaPeekAtLastError())) break;

            // exclusive sums to determine starts
            const int SCAN_BLOCK_THREADS = ActivePolicyT::ExclusiveSumPolicy::BLOCK_THREADS;
            DeviceRadixSortExclusiveSumKernel<MaxPolicyT, OffsetT>
                <<<num_passes, SCAN_BLOCK_THREADS, 0, stream>>>(d_bins);
            if (CubDebug(error = cudaPeekAtLastError())) break;

            // use the other buffer if no overwrite is allowed
            KeyT* d_keys_tmp = d_keys.Alternate();

                    if (CubDebug(error = cudaMemsetAsync(
                           d_lookback, 0, num_blocks * RADIX_DIGITS * sizeof(AtomicOffsetT),
                           stream))) break;

                    // log onesweep_kernel configuration
                    if (debug_synchronous)
                    {
                        _CubLog("Invoking onesweep_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread, "
                                "current bit %d, bit_grain %d, portion %d/%d\n",
                                num_blocks, ONESWEEP_BLOCK_THREADS, reinterpret_cast<long long>(stream),
                                ActivePolicyT::OnesweepPolicy::ITEMS_PER_THREAD, current_bit,
                                num_bits, static_cast<int>(portion), static_cast<int>(num_portions));
                    }

                    auto onesweep_kernel = DeviceRadixSortOnesweepKernel<
                        MaxPolicyT, IS_DESCENDING, KeyT, ValueT, OffsetT, PortionOffsetT>;
                    onesweep_kernel<<<num_blocks, ONESWEEP_BLOCK_THREADS, 0, stream>>>
