Move sync from for_each to parallel_for to fix CDP usages.

The device synchronization was decoupled from `THRUST_CDP_DISPATCH`
and was trying to sync regardless of CDP state. This led to
device syncs being invoked from device code when CDP is disabled
and the thread-serial implementation was used.

Some other algorithms that are implemented with `parallel_for`
have also been updated.

Old behavior:

1. `for_each`: calls `parallel_for`
2. `parallel_for`: calls appropriate impl using `THRUST_CDP_DISPATCH`
3. `parallel_for`: returns
4. `for_each`: calls `cub::detail::device_synchronize`

New behavior:

1. `for_each`: calls `parallel_for`
2. `parallel_for`: calls appropriate impl using `THRUST_CDP_DISPATCH`
4. `parallel_for`: calls `cub::detail::device_synchronize`
3. `parallel_for`: returns

Author: alliepiper

thrust/system/cuda/detail/fill.h

@@ -71,11 +71,6 @@ fill_n(execution_policy<Derived>& policy,
                          value),
                          count);
 
-  cuda_cub::throw_on_error(
-    cuda_cub::synchronize_optional(policy)
-  , "fill_n: failed to synchronize"
-  );
-
   return first + count;
 }    // func fill_n
 

thrust/system/cuda/detail/for_each.h

@@ -81,11 +81,6 @@ namespace cuda_cub {
                            for_each_f<Input, wrapped_t>(first, wrapped_op),
                            count);
 
-    cuda_cub::throw_on_error(
-      cuda_cub::synchronize_optional(policy)
-    , "for_each: failed to synchronize"
-    );
-
     return first + count;
   }
 

thrust/system/cuda/detail/parallel_for.h

@@ -161,7 +161,9 @@ parallel_for(execution_policy<Derived> &policy,
   THRUST_CDP_DISPATCH(
     (cudaStream_t stream = cuda_cub::stream(policy);
      cudaError_t  status = __parallel_for::parallel_for(count, f, stream);
-     cuda_cub::throw_on_error(status, "parallel_for failed");),
+     cuda_cub::throw_on_error(status, "parallel_for failed");
+     status = cuda_cub::synchronize_optional(policy);
+     cuda_cub::throw_on_error(status, "parallel_for: failed to synchronize");),
     // CDP sequential impl:
     (for (Size idx = 0; idx != count; ++idx)
      {

thrust/system/cuda/detail/swap_ranges.h

@@ -92,11 +92,6 @@ swap_ranges(execution_policy<Derived> &policy,
                                                ItemsIt2>(first1, first2),
                          num_items);
 
-  cuda_cub::throw_on_error(
-    cuda_cub::synchronize_optional(policy)
-  , "swap_ranges: failed to synchronize"
-  );
-
   return first2 + num_items;
 }
 

thrust/system/cuda/detail/tabulate.h

@@ -76,11 +76,6 @@ tabulate(execution_policy<Derived>& policy,
   cuda_cub::parallel_for(policy,
                          functor_t(first, tabulate_op),
                          count);
-
-  cuda_cub::throw_on_error(
-    cuda_cub::synchronize_optional(policy)
-  , "tabulate: failed to synchronize"
-  );
 }
 
 }    // namespace cuda_cub

thrust/system/cuda/detail/transform.h

@@ -232,11 +232,6 @@ namespace __transform {
                                              predicate),
                            num_items);
 
-    cuda_cub::throw_on_error(
-      cuda_cub::synchronize_optional(policy)
-    , "transform: failed to synchronize"
-    );
-
     return result + num_items;
   }
 
@@ -278,11 +273,6 @@ namespace __transform {
                                               predicate),
                            num_items);
 
-    cuda_cub::throw_on_error(
-      cuda_cub::synchronize_optional(policy)
-    , "transform: failed to synchronize"
-    );
-
     return result + num_items;
   }
 

thrust/system/cuda/detail/uninitialized_copy.h

@@ -87,11 +87,6 @@ uninitialized_copy_n(execution_policy<Derived> &policy,
                          functor_t(first, result),
                          count);
 
-  cuda_cub::throw_on_error(
-    cuda_cub::synchronize_optional(policy)
-  , "uninitialized_copy_n: failed to synchronize"
-  );
-
   return result + count;
 }
 

thrust/system/cuda/detail/uninitialized_fill.h

@@ -85,11 +85,6 @@ uninitialized_fill_n(execution_policy<Derived>& policy,
                          functor_t(first, x),
                          count);
 
-  cuda_cub::throw_on_error(
-    cuda_cub::synchronize_optional(policy)
-  , "uninitialized_fill_n: failed to synchronize"
-  );
-
   return first + count;
 }