Fix memcopy direction for concatenate (#18058)

`cudaMemcpyAsync` can return a `cudaError_t` value which should be checked for runtime errors. 

We should preserve the behavior of `thrust::copy` which got replaced with the `cudaMemcpyAsync` call in #17584. The driver may do the right thing and infer the source and destination pointer location instead of using the `cudaMemcpyKind`, but this still leads to weird circumstances where the copy type in code is DtoD while the actual copy at runtime is HtoD

Authors:
  - Tanmay Gujar (https://github.com/tgujar)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Karthikeyan (https://github.com/karthikeyann)
  - David Wendt (https://github.com/davidwendt)

URL: #18058

cpp/src/copying/concatenate.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -308,11 +308,11 @@ std::unique_ptr<column> for_each_concatenate(host_span<column_view const> views,
 
   auto count = 0;
   for (auto& v : views) {
-    cudaMemcpyAsync(m_view.begin<T>() + count,
-                    v.begin<T>(),
-                    v.size() * sizeof(T),
-                    cudaMemcpyDeviceToDevice,
-                    stream.value());
+    CUDF_CUDA_TRY(cudaMemcpyAsync(m_view.begin<T>() + count,
+                                  v.begin<T>(),
+                                  v.size() * sizeof(T),
+                                  cudaMemcpyDefault,
+                                  stream.value()));
     count += v.size();
   }