From 09d392f09de384d7df6e04cdb46d1a469549e3a4 Mon Sep 17 00:00:00 2001
From: Dan <dtmoodie>
Date: Tue, 15 Sep 2015 12:17:30 -0400
Subject: [PATCH] Added thrust tutorial.

---
 .../gpu/gpu-thrust-interop/main.cu            | 38 +++++++++++++++++--
 1 file changed, 34 insertions(+), 4 deletions(-)

diff --git a/samples/cpp/tutorial_code/gpu/gpu-thrust-interop/main.cu b/samples/cpp/tutorial_code/gpu/gpu-thrust-interop/main.cu
index c6784fef0b..e138fcc6fa 100644
--- a/samples/cpp/tutorial_code/gpu/gpu-thrust-interop/main.cu
+++ b/samples/cpp/tutorial_code/gpu/gpu-thrust-interop/main.cu
@@ -1,12 +1,14 @@
 #include "Thrust_interop.hpp"
+#include <opencv2/core/cuda_stream_accessor.hpp>
 
 #include <thrust/transform.h>
 #include <thrust/random.h>
 #include <thrust/sort.h>
+#include <thrust/system/cuda/execution_policy.h>
 struct prg
 {
 	float a, b;
-	
+
 	__host__ __device__
 		prg(float _a = 0.f, float _b = 1.f) : a(_a), b(_b) {};
 
@@ -39,6 +41,25 @@ template<typename T> struct pred_eq
 			return val.val[channel] == value;
 		return false;
 	}
+
+	__host__ __device__ bool operator()( const thrust::tuple<T, T, T>& val)
+	{
+		if (channel == 0)
+			return thrust::get<0>(val) == value;
+		if (channel == 1)
+			return thrust::get<1>(val) == value;
+		if (channel == 2)
+			return thrust::get<2>(val) == value;
+	}
+};
+template<typename T> struct pred_greater
+{
+	T value;
+	__host__ __device__ pred_greater(T value_) : value(value_){}
+	__host__ __device__ bool operator()(const T& val) const
+	{
+		return val > value;
+	}
 };
 
 
@@ -80,9 +101,18 @@ int main(void)
 		auto count = thrust::count(GpuMatBeginItr<int>(d_value), GpuMatEndItr<int>(d_value), 15);
 		std::cout << count << std::endl;
 	}
-
+	// Randomly fill an array then copy only values greater than 0.  Perform these tasks on a stream.
+	{
+		cv::cuda::GpuMat d_value(1, 100, CV_32F);
+		auto valueBegin = GpuMatBeginItr<float>(d_value);
+		auto valueEnd = GpuMatEndItr<float>(d_value);
+		cv::cuda::Stream stream;
+		thrust::transform(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), thrust::make_counting_iterator(0), thrust::make_counting_iterator(d_value.cols), valueBegin, prg(-1, 1));
+		int count = thrust::count_if(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), valueBegin, valueEnd, pred_greater<float>(0.0));
+		cv::cuda::GpuMat d_valueGreater(1, count, CV_32F);
+		thrust::copy_if(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), valueBegin, valueEnd, GpuMatBeginItr<float>(d_valueGreater), pred_greater<float>(0.0));
+		cv::Mat h_greater(d_valueGreater);
+	}
 	
-
-
 	return 0;
 }