DeviceSegmentedSort synchronizes default stream and produces wrong results when launched from a kernel

[fkallen]

1. DeviceSegmentedSort performs cudaMemcpy on the default stream instead of the supplied stream.
2. In device code, memcpy may execute before group_sizes and num_selected_groups are calculated. This leads to wrong results.

    if (CUB_IS_HOST_CODE)
    {
      #if CUB_INCLUDE_HOST_CODE
      if (CubDebug(error = cudaMemcpy(h_group_sizes,
                                      group_sizes.get(),
                                      num_selected_groups * sizeof(unsigned int),
                                      cudaMemcpyDeviceToHost)))
      {
        return error;
      }
      #endif
    }
    else
    {
      #if CUB_INCLUDE_DEVICE_CODE
      memcpy(h_group_sizes,
             group_sizes.get(),
             num_selected_groups * sizeof(unsigned int));
      #endif
    }

[gevtushenko]

Thank you for reporting this! I'll create PR soon.

[gevtushenko]

@fkallen could you check if the fix works for you?

[fkallen]

Your fix does work for me.