#include #include #include #include #include template __global__ void adjacent_difference_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result) { thrust::adjacent_difference(exec, first, last, result); } template __global__ void adjacent_difference_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result, BinaryFunction f) { thrust::adjacent_difference(exec, first, last, result, f); } template void TestAdjacentDifferenceDevice(ExecutionPolicy exec, const size_t n) { thrust::host_vector h_input = unittest::random_samples(n); thrust::device_vector d_input = h_input; thrust::host_vector h_output(n); thrust::device_vector d_output(n); thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin()); adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_output.begin()); { cudaError_t const err = cudaDeviceSynchronize(); ASSERT_EQUAL(cudaSuccess, err); } ASSERT_EQUAL(h_output, d_output); thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), thrust::plus()); adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_output.begin(), thrust::plus()); { cudaError_t const err = cudaDeviceSynchronize(); ASSERT_EQUAL(cudaSuccess, err); } ASSERT_EQUAL(h_output, d_output); // in-place operation thrust::adjacent_difference(h_input.begin(), h_input.end(), h_input.begin(), thrust::plus()); adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_input.begin(), thrust::plus()); { cudaError_t const err = cudaDeviceSynchronize(); ASSERT_EQUAL(cudaSuccess, err); } ASSERT_EQUAL(h_input, h_output); //computed previously ASSERT_EQUAL(d_input, d_output); //computed previously } template void TestAdjacentDifferenceDeviceSeq(const size_t n) { TestAdjacentDifferenceDevice(thrust::seq, n); } DECLARE_VARIABLE_UNITTEST(TestAdjacentDifferenceDeviceSeq); template void TestAdjacentDifferenceDeviceDevice(const size_t n) { TestAdjacentDifferenceDevice(thrust::device, n); } DECLARE_VARIABLE_UNITTEST(TestAdjacentDifferenceDeviceDevice); void TestAdjacentDifferenceCudaStreams() { cudaStream_t s; cudaStreamCreate(&s); thrust::device_vector input(3); thrust::device_vector output(3); input[0] = 1; input[1] = 4; input[2] = 6; thrust::adjacent_difference(thrust::cuda::par.on(s), input.begin(), input.end(), output.begin()); cudaStreamSynchronize(s); ASSERT_EQUAL(output[0], 1); ASSERT_EQUAL(output[1], 3); ASSERT_EQUAL(output[2], 2); cudaStreamDestroy(s); } DECLARE_UNITTEST(TestAdjacentDifferenceCudaStreams); struct detect_wrong_difference { bool * flag; __host__ __device__ detect_wrong_difference operator++() const { return *this; } __host__ __device__ detect_wrong_difference operator*() const { return *this; } template __host__ __device__ detect_wrong_difference operator+(Difference) const { return *this; } template __host__ __device__ detect_wrong_difference operator[](Index) const { return *this; } __device__ void operator=(long long difference) const { if (difference != 1) { *flag = false; } } }; void TestAdjacentDifferenceWithBigIndexesHelper(int magnitude) { thrust::counting_iterator begin(1); thrust::counting_iterator end = begin + (1ll << magnitude); ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude); thrust::device_ptr all_differences_correct = thrust::device_malloc(1); *all_differences_correct = true; detect_wrong_difference out = { thrust::raw_pointer_cast(all_differences_correct) }; thrust::adjacent_difference(thrust::device, begin, end, out); bool all_differences_correct_h = *all_differences_correct; thrust::device_free(all_differences_correct); ASSERT_EQUAL(all_differences_correct_h, true); } void TestAdjacentDifferenceWithBigIndexes() { TestAdjacentDifferenceWithBigIndexesHelper(30); TestAdjacentDifferenceWithBigIndexesHelper(31); TestAdjacentDifferenceWithBigIndexesHelper(32); TestAdjacentDifferenceWithBigIndexesHelper(33); } DECLARE_UNITTEST(TestAdjacentDifferenceWithBigIndexes);