Merge pull request #209 from elstehle/fix/universal_vector

Fixes universal_vector test failure on CTK 11.1 & gcc-6
NVIDIA · Jul 12, 2023 · 5161b5e · 5161b5e
2 parents 84fc205 + 0063850
commit 5161b5e
Showing 1 changed file with 27 additions and 13 deletions.
diff --git a/thrust/testing/cuda/device_side_universal_vector.cu b/thrust/testing/cuda/device_side_universal_vector.cu
@@ -5,8 +5,8 @@
 template <class VecT>
 __host__ __device__ void universal_vector_access(VecT &in, thrust::universal_vector<bool> &out)
 {
-  const int expected_front  = 4;
-  const int expected_back   = 2;
+  const int expected_front = 4;
+  const int expected_back  = 2;
 
   out[0] = in.size() == 2 &&               //
            in[0] == expected_front &&      //
@@ -43,15 +43,23 @@ void test_universal_vector_access(VecT &vec, thrust::universal_vector<bool> &out
 
 void TestUniversalVectorDeviceAccess()
 {
-  thrust::universal_vector<thrust::universal_vector<int>> in_storage(1);
-  thrust::universal_vector<int> &in = *thrust::raw_pointer_cast(in_storage.data());
+  using in_vector_t  = thrust::universal_vector<int>;
+  using out_vector_t = thrust::universal_vector<bool>;
 
+  in_vector_t *in_ptr{};
+  cudaMallocManaged(&in_ptr, sizeof(*in_ptr));
+  new (in_ptr) in_vector_t(1);
+
+  auto &in = *in_ptr;
   in.resize(2);
   in[0] = 4;
   in[1] = 2;
 
-  thrust::universal_vector<thrust::universal_vector<bool>> out_storage(1);
-  thrust::universal_vector<bool> &out = *thrust::raw_pointer_cast(out_storage.data());
+  out_vector_t *out_ptr{};
+  cudaMallocManaged(&out_ptr, sizeof(*out_ptr));
+  new (out_ptr) out_vector_t(1);
+  auto &out = *out_ptr;
+
   out.resize(1);
   out[0] = false;
 
@@ -61,24 +69,30 @@ DECLARE_UNITTEST(TestUniversalVectorDeviceAccess);
 
 void TestConstUniversalVectorDeviceAccess()
 {
-  thrust::universal_vector<thrust::universal_vector<int>> in_storage(1);
+  using in_vector_t  = thrust::universal_vector<int>;
+  using out_vector_t = thrust::universal_vector<bool>;
 
-  {
-    thrust::universal_vector<int> &in = *thrust::raw_pointer_cast(in_storage.data());
+  in_vector_t *in_ptr{};
+  cudaMallocManaged(&in_ptr, sizeof(*in_ptr));
+  new (in_ptr) in_vector_t(1);
 
+  {
+    auto &in = *in_ptr;
     in.resize(2);
     in[0] = 4;
     in[1] = 2;
   }
 
-  const thrust::universal_vector<int> &in = *thrust::raw_pointer_cast(in_storage.data());
+  const auto &const_in = *in_ptr;
 
-  thrust::universal_vector<thrust::universal_vector<bool>> out_storage(1);
-  thrust::universal_vector<bool> &out = *thrust::raw_pointer_cast(out_storage.data());
+  out_vector_t *out_ptr{};
+  cudaMallocManaged(&out_ptr, sizeof(*out_ptr));
+  new (out_ptr) out_vector_t(1);
+  auto &out = *out_ptr;
 
   out.resize(1);
   out[0] = false;
 
-  test_universal_vector_access(in, out);
+  test_universal_vector_access(const_in, out);
 }
 DECLARE_UNITTEST(TestConstUniversalVectorDeviceAccess);