Add patch reverting CCCL PR 211. (#511)

Something is going wrong in `thrust::copy`, so I am adding a patch to revert NVIDIA/cccl#211. This will fix the failing tests for rapidsai/cuspatial#1315.

This is where it's failing, with `thrust::copy`:
https://github.com/rapidsai/cuspatial/blob/8e8fa0689339c0288746b01a8af66fb599f5c770/cpp/tests/range/multipolygon_range_test.cu#L800

CI log: https://github.com/rapidsai/cuspatial/actions/runs/7229405991/job/19700477218?pr=1315#step:7:3280

```
[ RUN      ] MultipolygonRangeOneTest/1.OneMultipolygonRange
unknown file: Failure
C++ exception with description "__copy:: D->D: failed: cudaErrorInvalidValue: invalid argument" thrown in the test body.

[  FAILED  ] MultipolygonRangeOneTest/1.OneMultipolygonRange, where TypeParam = double (0 ms)
```

I traced it down to this CCCL PR - note that `thrust::transform(..., thrust::identity{})` works, but the new `thrust::copy` code path in NVIDIA/cccl#211 that calls `cudaMemcpyAsync` doesn't work for some reason. I was not able to identify the exact reason why that PR causes the failure. Patching this out will unblock RAPIDS and let us migrate forward with CCCL 2.2.0.

This helps with #502.

Authors:
  - Bradley Dice (https://github.com/bdice)

Approvers:
  - Michael Schellenberger Costa (https://github.com/miscco)
  - Robert Maynard (https://github.com/robertmaynard)

URL: #511

rapids-cmake/cpm/patches/cccl/revert_pr_211.diff

@@ -0,0 +1,79 @@
+diff --git a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
+index a0d9d492d..69c4e20df 100644
+--- a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
++++ b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
+@@ -33,38 +33,12 @@
+ #include <thrust/system/cuda/config.h>
+ #include <thrust/system/cuda/detail/execution_policy.h>
+ #include <thrust/system/cuda/detail/transform.h>
+-#include <thrust/system/cuda/detail/util.h>
+-#include <thrust/distance.h>
+ #include <thrust/functional.h>
+-#include <thrust/type_traits/is_trivially_relocatable.h>
+
+ THRUST_NAMESPACE_BEGIN
+ namespace cuda_cub {
+
+ namespace __copy {
+-  template <class Derived,
+-            class InputIt,
+-            class OutputIt>
+-  OutputIt THRUST_RUNTIME_FUNCTION
+-  device_to_device(execution_policy<Derived>& policy,
+-                   InputIt                    first,
+-                   InputIt                    last,
+-                   OutputIt                   result,
+-                   thrust::detail::true_type)
+-  {
+-    typedef typename thrust::iterator_traits<InputIt>::value_type InputTy;
+-    const auto n = thrust::distance(first, last);
+-    if (n > 0) {
+-      cudaError status;
+-      status = trivial_copy_device_to_device(policy,
+-                                             reinterpret_cast<InputTy*>(thrust::raw_pointer_cast(&*result)),
+-                                             reinterpret_cast<InputTy const*>(thrust::raw_pointer_cast(&*first)),
+-                                             n);
+-      cuda_cub::throw_on_error(status, "__copy:: D->D: failed");
+-    }
+-
+-    return result + n;
+-  }
+
+   template <class Derived,
+             class InputIt,
+@@ -73,32 +47,16 @@ namespace __copy {
+   device_to_device(execution_policy<Derived>& policy,
+                    InputIt                    first,
+                    InputIt                    last,
+-                   OutputIt                   result,
+-                   thrust::detail::false_type)
++                   OutputIt                   result)
+   {
+     typedef typename thrust::iterator_traits<InputIt>::value_type InputTy;
+     return cuda_cub::transform(policy,
+-                              first,
+-                              last,
+-                              result,
+-                              thrust::identity<InputTy>());
+-  }
+-
+-  template <class Derived,
+-            class InputIt,
+-            class OutputIt>
+-  OutputIt THRUST_RUNTIME_FUNCTION
+-  device_to_device(execution_policy<Derived>& policy,
+-                   InputIt                    first,
+-                   InputIt                    last,
+-                   OutputIt                   result)
+-  {
+-    return device_to_device(policy,
+                             first,
+                             last,
+                             result,
+-                            typename is_indirectly_trivially_relocatable_to<InputIt, OutputIt>::type());
++                            thrust::identity<InputTy>());
+   }
++
+ }    // namespace __copy
+
+ }    // namespace cuda_cub

rapids-cmake/cpm/versions.json

@@ -14,6 +14,11 @@
           "file" : "cccl/bug_fixes.diff",
           "issue" : "CCCL installs header-search.cmake files in nondeterministic order and has a typo in checking target creation that leads to duplicates",
           "fixed_in" : "2.3"
+        },
+        {
+          "file" : "cccl/revert_pr_211.diff",
+          "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue.",
+          "fixed_in" : ""
         }
       ]
     },