 #include <thrust/transform.h>
 #include <thrust/functional.h>
 #include <thrust/execution_policy.h>
+#include <thrust/iterator/iterator_adaptor.h>
 #include <thrust/device_vector.h>
 
 


1528935590000
'2018-06-14 00:19:50'
*/

struct gdf_extract_datetime_hour_date64_op : public thrust::unary_function<int64_t, int16_t>
{

	gdf_size_type num_chars_bitmask = ( ( input->size +( GDF_VALID_BITSIZE - 1)) / GDF_VALID_BITSIZE );
	thrust::copy(thrust::cuda::par.on(stream), input->valid, input->valid + num_chars_bitmask, output->valid); // copy over valid bitmask

	if ( input->dtype == GDF_DATE64 ) {
		thrust::device_ptr<int64_t> input_ptr((int64_t *) input->data);
		thrust::device_ptr<int16_t> output_ptr((int16_t *) output->data);
		gdf_extract_datetime_hour_date64_op op;
		thrust::transform(thrust::cuda::par.on(stream), thrust::detail::make_normal_iterator(input_ptr),
				thrust::detail::make_normal_iterator(input_ptr) + input->size, thrust::detail::make_normal_iterator(output_ptr), op);

	} else {
		return GDF_UNSUPPORTED_DTYPE;
	}

	cudaStreamSynchronize(stream);
