Автоматизация ввода-вывода CUDA/Пример внедрения автоматизации

#include <vector>
#include <cuda/vector.cuh>
#include <cuda/automation>

class SomeClass;

#ifdef __CUDACC__

class SomeClassCuda;

template <class T>
__device__ auto read_as(cuda_input_buffer_stream_d& is)->std::enable_if_t<std::is_same<SomeClassCuda, T>::value, T>;
template <class T>
__device__ auto read_as(cuda_input_buffer_stream_d& is)->std::enable_if_t<std::is_same<ComplexObject, T>::value, T>;

#endif  //__CUDACC__

struct ComplexObject
{
	int first;
	int second;
	#ifdef __CUDACC__
	template <class T>
	friend __device__ auto read_as(cuda_input_buffer_stream_d& is)->std::enable_if_t<std::is_same<ComplexObject, T>::value, T>;
	
	friend cuda_input_buffer_stream& operator<<(cuda_input_buffer_stream& is, const ComplexObject& obj);
	#endif  //__CUDACC__
};

#ifdef __CUDACC__

__host__ cuda_input_buffer_stream& operator<<(cuda_input_buffer_stream& is, const ComplexObject& obj)
{
	is << first;
	is << second;
	
	return is;
}

template <class T>
__device__ auto read_as(cuda_input_buffer_stream_d& is) -> std::enable_if_t<std::is_same<ComplexObject, T>::value, T>
{
	ComplexObject result;

	result.first = read_as<int>(is);
	result.first = read_as<int>(is);
	
	return result;
}

class SomeClassCuda
{
	cu::vector<ComplexObject> container;
public:

	typedef std::true_type custom_cuda_device_read;
	
	template <class T>
	friend __device__ auto read_as(cuda_input_buffer_stream_d& is) -> std::enable_if_t<std::is_same<SomeClassCuda, T>::value, T>;
};

template <class T>
__device__ auto read_as(cuda_input_buffer_stream_d& is) -> std::enable_if_t<std::is_same<SomeClassCuda, T>::value, T>
{
	SomeClassCuda result;

	std::size_t containerSize = read_as<std::size_t>(is);
	for (std::size_t i = 0; i < containerSize; ++i)
		result.container.emplace_back(read_as<ComplexObject>(is));
	
	return result;
}
#endif //__CUDACC__

class SomeClass
{
	std::vector<ComplexObject> container;
#ifdef __CUDACC__
	friend cuda_input_buffer_stream& operator<<(cuda_input_buffer_stream& is, const SomeClass& obj);
#endif //__CUDACC__
};

#ifdef __CUDACC__
__host__ cuda_input_buffer_stream& operator<<(cuda_input_buffer_stream& is, const SomeClass& obj)
{
	is << obj.container.size();
	for (std::size_t i = 0; i < obj.container.size(); ++i)
			is << obj.container[i];
	return is;
}
#endif //__CUDACC__