Merge pull request #1696 from senior-zero/fix-main/github/scan_by_key_docs

add tests for scan_by_key in-place execution

Author: gevtushenko

testing/cuda/scan_by_key.cu

@@ -78,7 +78,7 @@ void TestScanByKeyDevice(ExecutionPolicy exec)
   }
   ASSERT_EQUAL(d_output, h_output);
 
-  // in-place scans
+  // in-place scans: in/out values aliasing
   h_output = h_vals;
   d_output = d_vals;
   thrust::inclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_output.begin(), h_output.begin());
@@ -98,6 +98,24 @@ void TestScanByKeyDevice(ExecutionPolicy exec)
     ASSERT_EQUAL(cudaSuccess, err);
   }
   ASSERT_EQUAL(d_output, h_output);
+
+  // in-place scans: keys/values aliasing
+  thrust::inclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin());
+  inclusive_scan_by_key_kernel<<<1,1>>>(exec, d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys.begin());
+  {
+    cudaError_t const err = cudaDeviceSynchronize();
+    ASSERT_EQUAL(cudaSuccess, err);
+  }
+  ASSERT_EQUAL(d_keys, h_output);
+
+  d_keys = h_keys;
+  thrust::exclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin(), 11);
+  exclusive_scan_by_key_kernel<<<1,1>>>(exec, d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys.begin(), 11);
+  {
+    cudaError_t const err = cudaDeviceSynchronize();
+    ASSERT_EQUAL(cudaSuccess, err);
+  }
+  ASSERT_EQUAL(d_keys, h_output);
 }
 
 

testing/scan_by_key.exclusive.cu

@@ -309,6 +309,7 @@ void TestExclusiveScanByKeyInPlace(const size_t n)
   }
   thrust::device_vector<T> d_vals = h_vals;
 
+  // in-place scans: in/out values aliasing
   thrust::host_vector<T> h_output   = h_vals;
   thrust::device_vector<T> d_output = d_vals;
   thrust::exclusive_scan_by_key(h_keys.begin(),
@@ -322,6 +323,19 @@ void TestExclusiveScanByKeyInPlace(const size_t n)
                                 d_output.begin(),
                                 (T)11);
   ASSERT_EQUAL(d_output, h_output);
+
+  // in-place scans: in/out keys aliasing
+  thrust::exclusive_scan_by_key(h_keys.begin(),
+                                h_keys.end(),
+                                h_vals.begin(),
+                                h_keys.begin(),
+                                (T)11);
+  thrust::exclusive_scan_by_key(d_keys.begin(),
+                                d_keys.end(),
+                                d_vals.begin(),
+                                d_keys.begin(),
+                                (T)11);
+  ASSERT_EQUAL(d_keys, h_keys);
 }
 DECLARE_VARIABLE_UNITTEST(TestExclusiveScanByKeyInPlace);
 

testing/scan_by_key.inclusive.cu

@@ -309,7 +309,7 @@ void TestInclusiveScanByKeyInPlace(const size_t n)
   thrust::host_vector<T> h_output(n);
   thrust::device_vector<T> d_output(n);
 
-  // in-place scans
+  // in-place scans: in/out values aliasing
   h_output = h_vals;
   d_output = d_vals;
   thrust::inclusive_scan_by_key(h_keys.begin(),
@@ -321,6 +321,17 @@ void TestInclusiveScanByKeyInPlace(const size_t n)
                                 d_output.begin(),
                                 d_output.begin());
   ASSERT_EQUAL(d_output, h_output);
+
+  // in-place scans: in/out keys aliasing
+  thrust::inclusive_scan_by_key(h_keys.begin(),
+                                h_keys.end(),
+                                h_vals.begin(),
+                                h_keys.begin());
+  thrust::inclusive_scan_by_key(d_keys.begin(),
+                                d_keys.end(),
+                                d_vals.begin(),
+                                d_keys.begin());
+  ASSERT_EQUAL(d_keys, h_keys);
 }
 DECLARE_VARIABLE_UNITTEST(TestInclusiveScanByKeyInPlace);