text:
//---------------------------------------------------------------------------//
//!
//! \file   Utility_SampleMomentCollectionSnapshots.hpp
//! \author Alex Robinson
//! \brief  The sample moment collection snapshots declaration
//!
//---------------------------------------------------------------------------//

#ifndef UTILITY_SAMPLE_MOMENT_COLLECTION_SNAPSHOTS_HPP
#define UTILITY_SAMPLE_MOMENT_COLLECTION_SNAPSHOTS_HPP

// Std Lib Includes
#include <type_traits>

// Boost Includes
#include <boost/serialization/split_member.hpp>
#include <boost/serialization/access.hpp>

// FRENSIE Includes
#include "Utility_SampleMomentCollection.hpp"
#include "Utility_List.hpp"
#include "Utility_SerializationHelpers.hpp"

namespace Utility{

namespace Details{

//! The sample moment collection snapshots data extractor class
template<size_t N, typename T, typename Enabled = void>
struct SampleMomentCollectionSnapshotsDataExtractor;

} // end Details namespace

/*! The sample moment collection snapshots
 * \details This represents the empty collection. It cannot be instantiated
 * directly - only the non-empty collections can instantiate it. Note that
 * this class is a variadic template class and is designed in a very similar
 * way to the std::tuple class.
 */
template<typename T, template<typename,typename...> class SnapshotContainer, size_t... Ns>
class SampleMomentCollectionSnapshots;

//! The specialization of the sample moment collection with snapshots for the desired moments
template<typename T, template<typename,typename...> class SnapshotContainer, size_t N, size_t... Ns>
class SampleMomentCollectionSnapshots<T,SnapshotContainer,N,Ns...> : public SampleMomentCollectionSnapshots<T,SnapshotContainer,Ns...>
{
  // The base type
  typedef SampleMomentCollectionSnapshots<T,SnapshotContainer,Ns...> BaseType;

public:

  // The snapshot moment container type
  typedef SnapshotContainer<typename Utility::SampleMoment<N,T>::ValueType> MomentSnapshotContainerType;

  //! The snapshot indices container type
  typedef SnapshotContainer<uint64_t> SummationIndexContainerType;

  //! The snapshot sampling time container type
  typedef SnapshotContainer<double> SamplingTimeContainerType;

  //! The moment value type
  typedef typename MomentSnapshotContainerType::value_type MomentValueType;

  //! The summation index type
  typedef typename SummationIndexContainerType::value_type SummationIndexType;

  //! Default Constructor
  SampleMomentCollectionSnapshots();

  //! Constructor
  SampleMomentCollectionSnapshots( const size_t i );

  //! Copy Constructor
  SampleMomentCollectionSnapshots( const SampleMomentCollectionSnapshots& other_collection_snapshots );

  //! Assignment Operator
  SampleMomentCollectionSnapshots& operator=( const SampleMomentCollectionSnapshots& other_collection_snapshots );

  //! Destructor
  virtual ~SampleMomentCollectionSnapshots()
  { /* ... */ }

  //! Clear the collection snapshots
  void clear();

  //! Reset the collection snapshots
  void reset();

  //! Resize the collection snapshots (number of bins)
  void resize( const size_t i );

  //! Get the size of the collection (number of bins)
  size_t size() const;

  //! Get the number of snapshots
  size_t getNumberOfSnapshots() const;

  //! Take a snapshot of a sample moment collection
  void takeSnapshot( const uint64_t number_of_additional_samples,
                     const double sampling_time_from_last_snapshot,
                     const SampleMomentCollection<T,N,Ns...>& collection );

  //! Merge the snapshots
  void mergeSnapshots( const SampleMomentCollectionSnapshots& collection );

  //! Get the snapshot indices (summation indices)
  const SummationIndexContainerType& getSnapshotIndices() const;

  //! Get the snapshot sampling times (sampling times)
  const SamplingTimeContainerType& getSnapshotSamplingTimes() const;

private:

  // Make the data extractor class a friend
  template<size_t M, typename Collection, typename Enabled>
  friend class Details::SampleMomentCollectionSnapshotsDataExtractor;

  // Make all moment collections friends
  template<typename U, template<typename,typename...> class V, size_t... Ms>
  friend class SampleMomentCollectionSnapshots;

  // Make the boost::serialization::access class a friend
  friend class boost::serialization::access;

  // Save the collection data to an archive
  template<typename Archive>
  void save( Archive& ar, const unsigned version ) const;

  // Load the collection data from an archive
  template<typename Archive>
  void load( Archive& ar, const unsigned version );

  BOOST_SERIALIZATION_SPLIT_MEMBER();

  // The current scores
  std::vector<MomentSnapshotContainerType> d_score_snapshots;
};

//! Get the moment snapshots from the collection
template<size_t N, typename T, template<typename,typename...> class Container, size_t... Ms>
const Container<typename Utility::SampleMoment<N,T>::ValueType>& getScoreSnapshots(
                      const SampleMomentCollectionSnapshots<T,Container,Ms...>&
                      collection_snapshots,
                      const size_t i );

//! Get a score snapshot from the collection
template<size_t N, typename T, template<typename,typename...> class Container, size_t... Ms>
const typename SampleMoment<N,T>::ValueType& getScoreSnapshot(
                      const SampleMomentCollectionSnapshots<T,Container,Ms...>&
                      collection_snapshots,
                      const size_t i,
                      const size_t snapshot_index );

//! Get a moment snapshot from the collection
template<size_t N, typename T, template<typename,typename...> class Container, size_t... Ms>
SampleMoment<N,T> getMomentSnapshot(
                      const SampleMomentCollectionSnapshots<T,Container,Ms...>&
                      collection_snapshots,
                      const size_t i,
                      const size_t snapshot_index );

} // end Utility namespace

#define BOOST_SERIALIZATION_SAMPLE_MOMENT_COLLECTION_SNAPSHOTS_VERSION( VERSION ) \
  BOOST_SERIALIZATION_TEMPLATE_CLASS_VERSION_IMPL( \
    SampleMomentCollectionSnapshots, Utility, VERSION, \
    __BOOST_SERIALIZATION_FORWARD_AS_SINGLE_ARG__( typename T, template<typename,typename...> class Container, size_t... Ns ), \
    __BOOST_SERIALIZATION_FORWARD_AS_SINGLE_ARG__( T, Container, Ns... ) )

BOOST_SERIALIZATION_SAMPLE_MOMENT_COLLECTION_SNAPSHOTS_VERSION( 0 );

//---------------------------------------------------------------------------//
// Template Includes.
//---------------------------------------------------------------------------//

#include "Utility_SampleMomentCollectionSnapshots_def.hpp"

//---------------------------------------------------------------------------//

#endif // end UTILITY_SAMPLE_MOMENT_COLLECTION_SNAPSHOTS_HPP

//---------------------------------------------------------------------------//
// end Utility_SampleMomentCollectionSnapshots.hpp
//---------------------------------------------------------------------------//


text:
#include <iostream>
#include <fstream>

// Boost includes
#include <boost/filesystem.hpp>

#include <traceConfiguration.h>
#include <warpSynchronousGenerator.h>
#include <traceEvent.h>

#include <debug.h>
#include <string.h>

// Debugging messages
#ifdef REPORT_BASE
#undef REPORT_BASE
#endif

#define REPORT_BASE 0
#define REPORT_KERNEL_INSTRUCTIONS 0

//! if 1, warp size is set to the number of threads in the CTA
#define UNIFORM_CONTROL_FLOW_ONLY 1

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::WarpCounter::WarpCounter() :
	synchronous(0), events(0)
{
}

void trace::WarpSynchronousGenerator::WarpCounter::event(int warpSize,
	const trace::TraceEvent &event) {

	int warps = (event.active.size() + warpSize - 1) / warpSize;
	for (int warp = 0; warp < warps; warp++) {
		int tid = warp * warpSize;
		int active = 0;
		for (int i = 0; i < warpSize; i++) {
			if (event.active[tid + i]) {
				++active;
			}
		}

		if (active) {
			++events;
			synchronous += (active == warpSize ? 1 : 0); //= (size_t)active;
		}
	}
}

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::InstructionCounter::InstructionCounter()
{

}

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::LSInstructionCounter::LSInstructionCounter():
	wordsLoaded(0), wordsStored(0)
{

}

void trace::WarpSynchronousGenerator::LSInstructionCounter::event(int warpSize,
	const trace::TraceEvent &event) {

	if (event.instruction->addressSpace != ir::PTXInstruction::Global) {
		return;
	}

	ir::PTXU64 mask = (~(4 * warpSize - 1));
	size_t *wordsAccessed = (event.instruction->opcode == ir::PTXInstruction::St ?
		&wordsStored : &wordsLoaded);

	std::map< ir::PTXU64, int > addresses;
	int warps = (event.active.size() + warpSize - 1) / warpSize;
	for (int warp = 0; warp < warps; warp++) {
		int tid = warp * warpSize;
		int active = 0;
		for (int i = 0; i < warpSize; i++) {
			if (event.active[tid + i]) {
				++active;
				++*wordsAccessed;
			}
		}

		if (active) {
			addresses.clear();
			bool aligned = false;
			if (event.memory_size == 4) {
				aligned = true;
				for (int i = 0; i < warpSize; i++) {
					if (event.active[tid+i]) {
						addresses[event.memory_addresses[tid + i] & mask] = (tid + i);
					}
				}
				if (addresses.size() == 1) {
					++synchronous;
				}
			}
			++events;
		}
	}
}

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::BranchCounter::BranchCounter()
{

}

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::TargetCounter::TargetCounter()
{

}

//////////////////////////////////////////////////////////////////////////////////////////////////

trace::WarpSynchronousGenerator::SynchronousInstructionCounter::SynchronousInstructionCounter() {
	// initialize the warp instruction counters

	// add, sub, mul, div, mad
	// f32, s32, u32

	InstructionCounter empty;

	counterVectorizable[ir::PTXInstruction::Add][ir::PTXOperand::f32] = empty;
	counterVectorizable[ir::PTXInstruction::Add][ir::PTXOperand::s32] = empty;
	counterVectorizable[ir::PTXInstruction::Add][ir::PTXOperand::u32] = empty;

	counterVectorizable[ir::PTXInstruction::Sub][ir::PTXOperand::f32] = empty;
	counterVectorizable[ir::PTXInstruction::Sub][ir::PTXOperand::s32] = empty;
	counterVectorizable[ir::PTXInstruction::Sub][ir::PTXOperand::u32] = empty;

	counterVectorizable[ir::PTXInstruction::Mul][ir::PTXOperand::f32] = empty;
	counterVectorizable[ir::PTXInstruction::Mul][ir::PTXOperand::s32] = empty;
	counterVectorizable[ir::PTXInstruction::Mul][ir::PTXOperand::u32] = empty;

	counterVectorizable[ir::PTXInstruction::Mad][ir::PTXOperand::f32] = empty;
	counterVectorizable[ir::PTXInstruction::Mad][ir::PTXOperand::s32] = empty;
	counterVectorizable[ir::PTXInstruction::Mad][ir::PTXOperand::u32] = empty;

	counterVectorizable[ir::PTXInstruction::Div][ir::PTXOperand::f32] = empty;
	counterVectorizable[ir::PTXInstruction::Div][ir::PTXOperand::s32] = empty;
	counterVectorizable[ir::PTXInstruction::Div][ir::PTXOperand::u32] = empty;
}

void trace::WarpSynchronousGenerator::SynchronousInstructionCounter::event(
	int warpSize, const trace::TraceEvent &event) {

	// count specific control flow targets
	std::map< int, TargetCounter >::iterator target_it = counterTargets.find(event.PC);
	if (target_it != counterTargets.end()) {
		target_it->second.event(warpSize, event);
	}

	// handle each specific opcode its own way
	switch (event.instruction->opcode) {
	case ir::PTXInstruction::Bra:
		// handle branches specially
		counterBranches[event.PC].event(warpSize, event);
		break;

	case ir::PTXInstruction::Ld:
	case ir::PTXInstruction::St:
		// handle loads and stores specially
		counterLoadStore.event(warpSize, event);
		break;

	default:
		// record inst
	{
		WarpInstructionCounter::iterator w_it = counterVectorizable.find(event.instruction->opcode);
		if (w_it != counterVectorizable.end()) {
			DataInstructionCounter::iterator inst_counter = w_it->second.find(event.instruction->type);
			if (inst_counter != w_it->second.end()) {
				inst_counter->second.event(warpSize, event);
			}
		}
	}
		break;
	}
}

static void write(std::ostream &out,
	const trace::WarpSynchronousGenerator::SynchronousInstructionCounter &counter) {

	out << "\t\t\"loadStore\": { \"events\": " << counter.counterLoadStore.events
		<< ", \"synchronous\": " << counter.counterLoadStore.synchronous
		<< ", \"accessed\": " << (counter.counterLoadStore.wordsLoaded + counter.counterLoadStore.wordsStored)
		<< " },\n";

	out << "\t\t\"arithmetic\": {\n";

	int someEvents = 0;
	trace::WarpSynchronousGenerator::WarpInstructionCounter::const_iterator warp_it = counter.counterVectorizable.begin();
	for (; warp_it != counter.counterVectorizable.end(); ++warp_it) {
		trace::WarpSynchronousGenerator::DataInstructionCounter::const_iterator inst_it = warp_it->second.begin();
		for (; inst_it != warp_it->second.end(); ++inst_it) {
			if (inst_it->second.events) {
				if (someEvents++) {
					out << ",\n";
				}
				out << "\t\t\t\"" << ir::PTXInstruction::toString(warp_it->first) << "." <<
					ir::PTXOperand::toString(inst_it->first) << "\": { ";
				out << " \"events\": " << inst_it->second.events << ", \"synchronous\": " << inst_it->second.synchronous;
				out << " }";
			}
		}
	}
	out << "\n";

	out << "\t\t},\n";
	out << "\t\t\"branches\": [\n";

	std::map< int, trace::WarpSynchronousGenerator::BranchCounter >::const_iterator bra_it;
	bra_it = counter.counterBranches.begin();
	for (; bra_it != counter.counterBranches.end();) {
		out << "\t\t\t{ \"pc\": " << bra_it->first << ", \"events\": " << bra_it->second.events
			<< ", \"synchronous\": " << bra_it->second.synchronous
			<< ", \"target\": " << bra_it->second.target;
		++bra_it;
		out << " }" << (bra_it != counter.counterBranches.end() ? "," :"") << "\n";
	}
	out << "\t\t],\n";

	out << "\t\t\"targets\": [\n";
	std::map< int, trace::WarpSynchronousGenerator::TargetCounter >::const_iterator targ_it;
	targ_it = counter.counterTargets.begin();
	for (; targ_it != counter.counterTargets.end();) {
		out << "\t\t\t{ \"pc\": " << targ_it->first << ", \"events\": " << targ_it->second.events
			<< ", \"synchronous\": " << targ_it->second.synchronous;
		++targ_it;
		out << " }" << (targ_it != counter.counterTargets.end() ? "," :"") << "\n";
	}
	out << "\t\t]\n";

}

//////////////////////////////////////////////////////////////////////////////////////////////////

unsigned int trace::WarpSynchronousGenerator::_counter = 0;

trace::WarpSynchronousGenerator::WarpSynchronousGenerator() {

	// warp sizes
//	int sizes[] = {2, 4, 8, 16, 32, 0};
	int sizes[] = {4, 0};
	for (int i = 0; sizes[i]; i++) {
		warpSizes.push_back(sizes[i]);
	}
}

trace::WarpSynchronousGenerator::~WarpSynchronousGenerator() {

}

/*!	\brief Initializes the trace generator when a new kernel is
		about to be launched.
	\param kernel The kernel used to initialize the generator
*/
void trace::WarpSynchronousGenerator::initialize(const executive::ExecutableKernel& kernel) {

	const executive::EmulatedKernel & emuKernel =
		static_cast<const executive::EmulatedKernel &>(kernel);

	//
	// initialize kernel header
	//
	_entry.name = kernel.name;
	_entry.module = kernel.module->path();
	_entry.format = TraceGenerator::WarpSynchronousTraceFormat;
	_entry.path = "";
	_entry.header = "";
	_entry.program = "";
	_entry.gridDim = kernel.gridDim();
	_entry.blockDim = kernel.blockDim();
		std::string name = kernel.name;

	// only look for complete uniformity
#if UNIFORM_CONTROL_FLOW_ONLY
	warpSizes.clear();
	warpSizes.push_back(_entry.blockDim.x * _entry.blockDim.y * _entry.blockDim.z);
#endif

	if( name.size() > 20 ) {
		name.resize( 20 );
	}

	std::stringstream stream;
	stream << _entry.format << "_" << trace::WarpSynchronousGenerator::_counter++;

	boost::filesystem::path path( database );
	path = path.parent_path();
	path /= _entry.program + "_" + name + "_" + stream.str() + ".json";
	path = boost::filesystem::system_complete( path );

	_entry.header = path.string();

	//
	// identify all branches and create counters for them
	//
	std::map< int, int > branchPCs;
	std::map< int, std::vector<int> > targetPCs; // map's a target PC to list of branches jumping to it
	if (kernel.ISA == ir::Instruction::Emulated) {
		const executive::EmulatedKernel & emuKernel =
			static_cast<const executive::EmulatedKernel &>(kernel);

		targetPCs[0].push_back(-1);	// first instruction gets a push

		executive::EmulatedKernel::PTXInstructionVector::const_iterator inst_it;

		int pc = 0;
		for (inst_it = emuKernel.instructions.begin(); inst_it != emuKernel.instructions.end();
			++inst_it, ++pc) {
			const ir::PTXInstruction &inst = *inst_it;

			// complete per-instruction analysis
			switch (inst.opcode) {
				case ir::PTXInstruction::Bra:
					branchPCs[pc] = inst.branchTargetInstruction;
					targetPCs[inst.branchTargetInstruction].push_back(pc);
					targetPCs[pc+1].push_back(pc);
					break;

				default:
					break;
			}
		}
	}

	// obtain control flow graph
	controlFlowGraph = emuKernel.cfg();
	branchTargetsToBlock = emuKernel.branchTargetsToBlock;
	for (std::map< int, std::string >::const_iterator it = branchTargetsToBlock.begin();
		it != branchTargetsToBlock.end(); ++it) {
		blockToBranchTarget[it->second] = it->first;
	}

	// for each warp size, initialize counters and create branch records
	warpCounters.clear();
	for (std::vector<int>::const_iterator ws_it = warpSizes.begin(); ws_it != warpSizes.end();
		++ws_it) {
		SynchronousInstructionCounter syncCounter;
		for (std::map<int,int>::const_iterator bra_it = branchPCs.begin();
			bra_it != branchPCs.end(); ++bra_it) {

			BranchCounter braCounter;
			braCounter.target = bra_it->second;
			syncCounter.counterBranches[bra_it->first] = braCounter;
		}
		for (std::map<int, std::vector<int> >::const_iterator targ_it = targetPCs.begin();
			targ_it != targetPCs.end(); ++targ_it) {
			TargetCounter targCounter;

			syncCounter.counterTargets[targ_it->first] = targCounter;
		}
		warpCounters[*ws_it] = syncCounter;
	}
}

/*!
	\brief Called whenever an event takes place.

	\param even The TraceEvent that just occurred
	Note, the const reference 'event' is only valid until event()
	returns
*/
void trace::WarpSynchronousGenerator::event( const trace::TraceEvent& event ) {
	for (std::map< int, SynchronousInstructionCounter >::iterator counter_it = warpCounters.begin();
		counter_it != warpCounters.end(); ++counter_it) {
		counter_it->second.event(counter_it->first, event);
	}
}

/*!	\brief Finalize the trace and dump the results to disk.

	Add a databse entry for the trace as well.
*/
void trace::WarpSynchronousGenerator::finish() {

	report("WarpSynchronousGenerator::finish()");

	// always update database and emit results as JSON
	outputSynchronousStatistics();

	if (TraceConfiguration::Singleton.warpSynchronous.emitHotPaths) {
		// emit DOT files
		for (std::vector<int>::const_iterator ws_it = warpSizes.begin();
			ws_it != warpSizes.end(); ++ws_it) {
			int warpSize = *ws_it;
			report("  outputting hot paths - ws: " << warpSize);
			outputHotPaths(warpSize);
		}
	}
}

void trace::WarpSynchronousGenerator::outputSynchronousStatistics() {

	// update database
	_entry.updateDatabase(database);

	std::ofstream traceLog(_entry.header.c_str());

	//
	// serialize results
	//
	traceLog << "{\n";
	traceLog << "\t\"kernel\": \"" << _entry.name << "\",\n";
	traceLog << "\t\"module\": \"" << _entry.module << "\",\n";
	traceLog << "\t\"threads\": "
		<< _entry.blockDim.x * _entry.blockDim.y * _entry.blockDim.z << ",\n";
	traceLog << "\t\"gridDim\": [" << _entry.gridDim.x << ", " << _entry.gridDim.y << ", "
		<< _entry.gridDim.z << "],\n";
	traceLog << "\t\"blockDim\": [" << _entry.blockDim.x << ", " << _entry.blockDim.y << ", "
		<< _entry.blockDim.z << "],\n";
	traceLog << "\t\"counters\":\n[\n";
	for (std::map< int, SynchronousInstructionCounter >::iterator
		counter_it = warpCounters.begin();
		counter_it != warpCounters.end(); ) {
		traceLog << "\t{\n\t\"warpSize\": " << counter_it->first << ",\n";
		write(traceLog, counter_it->second);
		++counter_it;
		traceLog << "\t}" << (counter_it != warpCounters.end() ? "," : "") << "\n";
	}
	traceLog << "]\n}\n\n";
}

//////////////////////////////////////////////////////////////////////////////////////////////////

/*
			class DotFormatter : public ir::ControlFlowGraph::BasicBlock::DotFormatter {
			public:

				const WarpSynchronousGenerator *traceGenerator;

				static std::string colorToString(unsigned int color);
			public:
				std::string toString(const ir::ControlFlowGraph::BasicBlock *block);
				std::string toString(const ir::ControlFlowGraph::Edge *edge);
			};
// trace::WarpSynchronousGenerator::DotFormatter::
*/

std::string trace::WarpSynchronousGenerator::DotFormatter::colorToString(unsigned int color) {

	unsigned short a = ((color >> 24) & 0x0ff);
	unsigned short r = ((color >> 16) & 0x0ff);
	unsigned short g = ((color >> 8) & 0x0ff);
	unsigned short b = ((color) & 0x0ff);

	return colorToString(r, g, b, a);
}

std::string trace::WarpSynchronousGenerator::DotFormatter::colorToString(
	unsigned int r, unsigned int g, unsigned int b, unsigned int a) {
	std::stringstream ss;

	ss << "#";
	ss << std::setfill('0') << std::setw(2) << std::hex << (r & 0x0ff);
	ss << std::setfill('0') << std::setw(2) << std::hex << (g & 0x0ff);
	ss << std::setfill('0') << std::setw(2) << std::hex << (b & 0x0ff);
//	ss << std::setfill('0') << std::setw(2) << std::hex << (a & 0x0ff);

	return ss.str();
}

std::string trace::WarpSynchronousGenerator::DotFormatter::toString(
	const ir::ControlFlowGraph::BasicBlock *block) {
	std::stringstream ss;

	std::string label = block->label();
	if (blockCounter.find(label) == blockCounter.end()) {
		ss << "[shape=record,label=\"" <<
			hydrazine::toGraphVizParsableLabel(block->label()) << "\"] // no block counter";
		return ss.str();
	}

	size_t events = blockCounter[block->label()].events,
		synchronous = blockCounter[block->label()].synchronous;
	float t = 0.0f;
	double activity = (events ? (double)synchronous / (double)events : 0);

	/*
	// log-scale color intensity based on number of times block has been reached by a warp
	if (events) {
		t = log((float)(events)) / (float)log((float)(maxEvents));
	}
	unsigned int r = (unsigned int)(t * 255.0);
	unsigned int g = 0, b = 0;
	*/
	if (events) {
		t = (float)activity;
		if (t < 0.95) { t = 0.0f; }
	}
	unsigned int r = (unsigned int)(t * 255.0);
	unsigned int g = 0, b = 0;

	ss << "[";
	ss << "fillcolor=\"" << colorToString(r,g,b) << "\",style=filled,";
	ss << "shape=record,label=\"{" << hydrazine::toGraphVizParsableLabel(block->label())
		<< " | events: " << events << " | synchronous: " << synchronous << " | activity: "
		<< activity << "}\"";
	if (t < 0.5) {
		ss << ",fontcolor=\"#aaaaaa\"";
	}
	ss << "]";
	return ss.str();
}

std::string trace::WarpSynchronousGenerator::DotFormatter::toString(
	const ir::ControlFlowGraph::BasicBlock::Edge *edge) {
	std::stringstream out;

	if (edge->type == ir::ControlFlowGraph::BasicBlock::Edge::Dummy) {
		out << "[style=dotted]";
	}
	else {
		// color it blue if the edge is entirely uniform or not
		if (counterEdges.find(edge) != counterEdges.end()) {
			if (counterEdges[edge].synchronous == counterEdges[edge].events) {
				out << "[color=blue,style=bold]";
			}
		}
	}

	return out.str();
}

//////////////////////////////////////////////////////////////////////////////////////////////////

/*!
	\brief emits DOT files for each kernel with blocks colored by hot path for a given warp size
*/
void trace::WarpSynchronousGenerator::outputHotPaths(int warpSize) {
	const SynchronousInstructionCounter & counter = warpCounters[warpSize];

	DotFormatter formatter;
	formatter.warpSize = warpSize;
	formatter.traceGenerator = this;

	std::map< int, TargetCounter >::const_iterator target_it;

	formatter.maxEvents = 0;
	for (target_it = counter.counterTargets.begin(); target_it != counter.counterTargets.end();
		++target_it) {
		//blockColors[branchTargetsToBlock[target_it->first]] = 0;
		if (formatter.maxEvents < target_it->second.events) {
			formatter.maxEvents = target_it->second.events;
		}
		if ( (!formatter.minEvents && target_it->second.events) || (target_it->second.events &&
			target_it->second.events < formatter.minEvents)) {
			formatter.minEvents = target_it->second.events;
		}
	}
	if (formatter.maxEvents) {
		for (target_it = counter.counterTargets.begin(); target_it != counter.counterTargets.end();
			++target_it) {
			//blockColors[branchTargetsToBlock[target_it->first]] = color;
			formatter.blockCounter[ branchTargetsToBlock[target_it->first] ] = target_it->second;
		}
	}

	std::stringstream ss;

	boost::filesystem::path path( database );
	path = path.parent_path();
	path = boost::filesystem::system_complete( path );
	ss << path.string() << "/" << _entry.program << "_" << _entry.name << "_"
		<< _counter << "_ws" << warpSize << "_hotpath.dot";

	std::string filename = ss.str();

	report(" outputHotPaths(ws: " << warpSize << ", filename = " << filename << ")");

	std::ofstream cfgFile(filename.c_str());
	cfgFile << "/*\n";
	cfgFile << "     kernel: " << " kernel-name " << "\n";
	cfgFile << "  warp size: " << warpSize << "\n";
	cfgFile << "*/\n";

	controlFlowGraph->write(cfgFile, formatter);
}

//////////////////////////////////////////////////////////////////////////////////////////////////

ext: cc

max_stars_repo_name: wahibium/KFF

max_issues_repo_name: wahibium/KFF

max_forks_repo_name: wahibium/KFF


text:
//
// Copyright (c) 2016-2020 Kris Jusiak (kris at jusiak dot net)
//
// Distributed under the Boost Software License, Version 1.0.
// (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
//
#include <boost/sml.hpp>
#include <memory>

namespace sml = boost::sml;

struct fwd_sm;

class c {
 public:
  c();
  void update();
  bool is_done() const;

 private:
  std::unique_ptr<sml::sm<fwd_sm>> sm;
};

ext: hpp

max_stars_repo_name: seigtm/sml

max_issues_repo_name: seigtm/sml

max_forks_repo_name: seigtm/sml


text:
/* =========================================================================
   Copyright (c) 2010-2014, Institute for Microelectronics,
                            Institute for Analysis and Scientific Computing,
                            TU Wien.
   Portions of this software are copyright by UChicago Argonne, LLC.

                            -----------------
                  ViennaCL - The Vienna Computing Library
                            -----------------

   Project Head:    Karl Rupp                   rupp@iue.tuwien.ac.at

   (A list of authors and contributors can be found in the PDF manual)

   License:         MIT (X11), see file LICENSE in the base directory
============================================================================= */

/** \file tests/src/spmdm.cpp  Tests sparse-matrix-dense-matrix products.
*   \test  Tests sparse-matrix-dense-matrix products.
**/

//
// include necessary system headers
//
#include <iostream>
#include <cmath>

//
// ublas includes
//
#include <boost/numeric/ublas/io.hpp>
#include <boost/numeric/ublas/triangular.hpp>
#include <boost/numeric/ublas/matrix_sparse.hpp>
#include <boost/numeric/ublas/matrix.hpp>
#include <boost/numeric/ublas/matrix_proxy.hpp>
#include <boost/numeric/ublas/operation.hpp>
#include <boost/numeric/ublas/operation_sparse.hpp>
#include <boost/numeric/ublas/io.hpp>
#include <boost/numeric/ublas/lu.hpp>

// Must be set if you want to use ViennaCL algorithms on ublas objects
#define VIENNACL_WITH_UBLAS 1
//#define VIENNACL_WITH_OPENCL 1
//#define VIENNACL_WITH_CUDA 1
//#define VIENNACL_DEBUG_KERNEL 1
//#define VIENNACL_BUILD_INFO 1

//
// ViennaCL includes
//
#include "viennacl/scalar.hpp"
#include "viennacl/vector.hpp"
#include "viennacl/matrix.hpp"
#include "viennacl/linalg/direct_solve.hpp"
#include "viennacl/compressed_matrix.hpp"
#include "viennacl/coordinate_matrix.hpp"
#include "viennacl/ell_matrix.hpp"
#include "viennacl/hyb_matrix.hpp"
#include "viennacl/linalg/prod.hpp"       //generic matrix-vector product
#include "viennacl/linalg/norm_2.hpp"     //generic l2-norm for vectors
#include "viennacl/io/matrix_market.hpp"

// Some helper functions for this tutorial:
#include "Random.hpp"

using namespace boost::numeric;

template< typename ScalarType >
int check_matrices(const ublas::matrix< ScalarType >& ref_mat, const ublas::matrix< ScalarType >& mat, ScalarType eps) {

  std::size_t size1, size2;
  size1 = ref_mat.size1(); size2 = ref_mat.size2();
  if ( (size1 != mat.size1()) || (size2 != mat.size2()) )
    return EXIT_FAILURE;

  for (unsigned int i = 0; i < size1; i++)
    for (unsigned int j = 0; j < size2; j++)
    {
      ScalarType rel_error = std::abs(ref_mat(i,j) - mat(i,j)) / std::max(std::abs(ref_mat(i,j)), std::abs(mat(i,j)));
      if ( rel_error > eps ) {
        std::cout << "ERROR: Verification failed at (" << i <<", "<< j << "): "
                  << " Expected: " << ref_mat(i,j) << ", got: " << mat(i,j) << " (relative error: " << rel_error << ")" << std::endl;
        return EXIT_FAILURE;
      }
    }

  std::cout << "Everything went well!" << std::endl;
  return EXIT_SUCCESS;
}

template<typename NumericT, typename ResultLayoutT, typename FactorLayoutT>
int test(NumericT epsilon)
{
  int retVal = EXIT_SUCCESS;

  ublas::compressed_matrix<NumericT>    ublas_lhs;

  if (viennacl::io::read_matrix_market_file(ublas_lhs, "../examples/testdata/mat65k.mtx") == EXIT_FAILURE)
  {
    std::cout << "Error reading Matrix file" << std::endl;
    return EXIT_FAILURE;
  }

  // add some extra weight to diagonal in order to avoid issues with round-off errors
  for (std::size_t i=0; i<ublas_lhs.size1(); ++i)
    ublas_lhs(i,i) *= NumericT(1.5);

  std::size_t cols_rhs = 1;

  viennacl::compressed_matrix<NumericT> compressed_lhs;
  viennacl::ell_matrix<NumericT>        ell_lhs;
  viennacl::coordinate_matrix<NumericT> coo_lhs;
  viennacl::hyb_matrix<NumericT>     hyb_lhs;

  ublas::matrix<NumericT> ublas_result;
  viennacl::matrix<NumericT, ResultLayoutT> result;

  viennacl::copy( ublas_lhs, compressed_lhs);
  viennacl::copy( ublas_lhs, ell_lhs);
  viennacl::copy( ublas_lhs, coo_lhs);
  viennacl::copy( ublas_lhs, hyb_lhs);

  ublas::matrix<NumericT> ublas_rhs1(ublas_lhs.size2(), cols_rhs);
  viennacl::matrix<NumericT, FactorLayoutT> rhs1(ublas_lhs.size2(), cols_rhs);

  ublas::matrix<NumericT> ublas_rhs2;
  viennacl::matrix<NumericT, FactorLayoutT> rhs2;

  ublas::matrix<NumericT> temp(ublas_rhs1.size1(), cols_rhs);

  for (unsigned int i = 0; i < ublas_rhs1.size1(); i++)
    for (unsigned int j = 0; j < ublas_rhs1.size2(); j++)
      ublas_rhs1(i,j) = NumericT(0.5) + NumericT(0.1) * random<NumericT>();
  viennacl::copy( ublas_rhs1, rhs1);

  ublas_rhs2 = ublas::trans( ublas_rhs1);
  viennacl::copy( ublas_rhs2, rhs2);

  /* gold result */
  ublas_result = ublas::prod( ublas_lhs, ublas_rhs1);

  /******************************************************************/
  std::cout << "Testing compressed(CSR) lhs * dense rhs" << std::endl;
  result = viennacl::linalg::prod( compressed_lhs, rhs1);

  temp.clear();
  viennacl::copy( result, temp);
  retVal = check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/
  std::cout << "Testing compressed(ELL) lhs * dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( ell_lhs, rhs1);

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/

  std::cout << "Testing compressed(COO) lhs * dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( coo_lhs, rhs1);

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/

  std::cout << "Testing compressed(HYB) lhs * dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( hyb_lhs, rhs1);

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/

  /* gold result */
  ublas_result = ublas::prod( ublas_lhs, ublas::trans(ublas_rhs2));

  /******************************************************************/
  std::cout << std::endl << "Testing compressed(CSR) lhs * transposed dense rhs:" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( compressed_lhs, viennacl::trans(rhs2));

  temp.clear();
  viennacl::copy( result, temp);
  retVal = check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/
  std::cout << "Testing compressed(ELL) lhs * transposed dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( ell_lhs, viennacl::trans(rhs2));

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/
  std::cout << "Testing compressed(COO) lhs * transposed dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( coo_lhs, viennacl::trans(rhs2));

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/

  std::cout << "Testing compressed(HYB) lhs * transposed dense rhs" << std::endl;
  result.clear();
  result = viennacl::linalg::prod( hyb_lhs, viennacl::trans(rhs2));

  temp.clear();
  viennacl::copy( result, temp);
  check_matrices(ublas_result, temp, epsilon);

  /******************************************************************/
  if (retVal == EXIT_SUCCESS) {
    std::cout << "Tests passed successfully" << std::endl;
  }

  return retVal;
}

//
// -------------------------------------------------------------
//
int main()
{
  std::cout << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << "## Test :: Sparse-Dense Matrix Multiplication" << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << std::endl;

  int retval = EXIT_SUCCESS;

  std::cout << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << std::endl;
  {
    typedef float NumericT;
    NumericT epsilon = static_cast<NumericT>(1E-4);
    std::cout << "# Testing setup:" << std::endl;
    std::cout << "  eps:     " << epsilon << std::endl;
    std::cout << "  numeric: float" << std::endl;
    std::cout << "  layout:  row-major, row-major" << std::endl;
    retval = test<NumericT, viennacl::row_major, viennacl::row_major>(epsilon);
    if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
    else
        return retval;

    std::cout << "# Testing setup:" << std::endl;
    std::cout << "  eps:     " << epsilon << std::endl;
    std::cout << "  numeric: float" << std::endl;
    std::cout << "  layout:  row-major, column-major" << std::endl;
    retval = test<NumericT, viennacl::row_major, viennacl::column_major>(epsilon);
    if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
    else
        return retval;

    std::cout << "# Testing setup:" << std::endl;
    std::cout << "  eps:     " << epsilon << std::endl;
    std::cout << "  numeric: float" << std::endl;
    std::cout << "  layout:  column-major, row-major" << std::endl;
    retval = test<NumericT, viennacl::column_major, viennacl::row_major>(epsilon);
    if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
    else
        return retval;

    std::cout << "# Testing setup:" << std::endl;
    std::cout << "  eps:     " << epsilon << std::endl;
    std::cout << "  numeric: float" << std::endl;
    std::cout << "  layout:  column-major, column-major" << std::endl;
    retval = test<NumericT, viennacl::column_major, viennacl::column_major>(epsilon);
    if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
    else
        return retval;

  }
  std::cout << std::endl;
  std::cout << "----------------------------------------------" << std::endl;
  std::cout << std::endl;

#ifdef VIENNACL_WITH_OPENCL
  if ( viennacl::ocl::current_device().double_support() )
#endif
  {
    {
      typedef double NumericT;
      NumericT epsilon = 1.0E-12;
      std::cout << "# Testing setup:" << std::endl;
      std::cout << "  eps:     " << epsilon << std::endl;
      std::cout << "  numeric: double" << std::endl;
      std::cout << "  layout:  row-major, row-major" << std::endl;
      retval = test<NumericT, viennacl::row_major, viennacl::row_major>(epsilon);
      if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
      else
        return retval;

      std::cout << "# Testing setup:" << std::endl;
      std::cout << "  eps:     " << epsilon << std::endl;
      std::cout << "  numeric: double" << std::endl;
      std::cout << "  layout:  row-major, column-major" << std::endl;
      retval = test<NumericT, viennacl::row_major, viennacl::column_major>(epsilon);
      if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
      else
        return retval;

      std::cout << "# Testing setup:" << std::endl;
      std::cout << "  eps:     " << epsilon << std::endl;
      std::cout << "  numeric: double" << std::endl;
      std::cout << "  layout:  column-major, row-major" << std::endl;
      retval = test<NumericT, viennacl::column_major, viennacl::row_major>(epsilon);
      if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
      else
        return retval;

      std::cout << "# Testing setup:" << std::endl;
      std::cout << "  eps:     " << epsilon << std::endl;
      std::cout << "  numeric: double" << std::endl;
      std::cout << "  layout:  column-major, column-major" << std::endl;
      retval = test<NumericT, viennacl::column_major, viennacl::column_major>(epsilon);
      if ( retval == EXIT_SUCCESS )
        std::cout << "# Test passed" << std::endl;
      else
        return retval;
    }
    std::cout << std::endl;
    std::cout << "----------------------------------------------" << std::endl;
    std::cout << std::endl;
  }
#ifdef VIENNACL_WITH_OPENCL
  else
    std::cout << "No double precision support, skipping test..." << std::endl;
#endif

  std::cout << std::endl;
  std::cout << "------- Test completed --------" << std::endl;
  std::cout << std::endl;

  return retval;
}

ext: cpp

max_stars_repo_name: ddemidov/viennacl-dev

max_issues_repo_name: ddemidov/viennacl-dev

max_forks_repo_name: ddemidov/viennacl-dev


text:
// http://turtle.sourceforge.net
//
// Copyright Mathieu Champlon 2015
//
// Distributed under the Boost Software License, Version 1.0.
// (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)

//[ static_objects_problem
#define BOOST_AUTO_TEST_MAIN
#include <boost/test/auto_unit_test.hpp>
#include <turtle/mock.hpp>
#include <ostream>

namespace
{
    struct my_class
    {
        my_class( int i )
            : i_( i )
        {}

        int i_;
    };

    std::ostream& operator<<( std::ostream& os, const my_class* c )
    {
        return os << "my_class " << c->i_; // the 'c' pointer must be valid when logging
    }

    MOCK_FUNCTION( f, 1, void( my_class* ) ) // being static 'f' outlive the test case
}

BOOST_AUTO_TEST_CASE( static_objects_problem )
{
    my_class c( 42 );
    MOCK_EXPECT( f ).once().with( &c ); // the set expectation will also outlive the test case and leak into other test cases using 'f'
} // the 'c' instance goes out of scope and the '&c' pointer becomes dangling
//]

//[ static_objects_partial_solution
struct fixture
{
    ~fixture()
    {
        mock::reset(); // the use of a fixture ensures the reset will prevent the expectations from leaking into other test cases
    }
};

BOOST_FIXTURE_TEST_CASE( static_object_partial_solution, fixture )
{
    my_class c( 42 );
    MOCK_EXPECT( f ).once().with( &c );
    f( &c );
    mock::verify(); // verify the expectations before local objects are destroyed and before the fixture resets them
}
//]

//[ static_objects_solution
BOOST_FIXTURE_TEST_CASE( static_objects_solution, mock::cleanup ) // actually the library includes a ready to use fixture just like the one described
{
    my_class c( 42 );
    MOCK_EXPECT( f ).once().with( &c );
    f( &c );
    mock::verify();
}
//]


text:
/***************************************************************************
 * Copyright 1998-2020 by authors (see AUTHORS.txt)                        *
 *                                                                         *
 *   This file is part of LuxCoreRender.                                   *
 *                                                                         *
 * Licensed under the Apache License, Version 2.0 (the "License");         *
 * you may not use this file except in compliance with the License.        *
 * You may obtain a copy of the License at                                 *
 *                                                                         *
 *     http://www.apache.org/licenses/LICENSE-2.0                          *
 *                                                                         *
 * Unless required by applicable law or agreed to in writing, software     *
 * distributed under the License is distributed on an "AS IS" BASIS,       *
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.*
 * See the License for the specific language governing permissions and     *
 * limitations under the License.                                          *
 ***************************************************************************/

#if !defined(LUXRAYS_DISABLE_OPENCL)

#include <boost/lexical_cast.hpp>
#include <boost/algorithm/string/replace.hpp>

#include "luxrays/core/geometry/transform.h"
#include "luxrays/utils/ocl.h"
#include "luxrays/devices/ocldevice.h"
#include "luxrays/kernels/kernels.h"

#include "luxcore/cfg.h"

#include "slg/slg.h"
#include "slg/kernels/kernels.h"
#include "slg/renderconfig.h"
#include "slg/engines/pathoclbase/pathoclbase.h"
#include "slg/samplers/sobol.h"

using namespace std;
using namespace luxrays;
using namespace slg;

//------------------------------------------------------------------------------
// PathOCLBaseOCLRenderThread kernels related methods
//------------------------------------------------------------------------------

void PathOCLBaseOCLRenderThread::CompileKernel(HardwareIntersectionDevice *device,
			HardwareDeviceProgram *program,
			HardwareDeviceKernel **kernel,
			size_t *workGroupSize, const std::string &name) {
	delete *kernel;
	SLG_LOG("[PathOCLBaseRenderThread::" << threadIndex << "] Compiling " << name << " Kernel");
	device->GetKernel(program, kernel, name.c_str());

	if (device->GetDeviceDesc()->GetForceWorkGroupSize() > 0)
		*workGroupSize = device->GetDeviceDesc()->GetForceWorkGroupSize();
	else {
		*workGroupSize = device->GetKernelWorkGroupSize(*kernel);
		SLG_LOG("[PathOCLBaseRenderThread::" << threadIndex << "] " << name << " workgroup size: " << *workGroupSize);
	}
}

void PathOCLBaseOCLRenderThread::GetKernelParamters(
		vector<string> &params,
		HardwareIntersectionDevice *intersectionDevice,
		const string renderEngineType,
		const float epsilonMin, const float epsilonMax) {
	params.push_back("-D LUXRAYS_OPENCL_KERNEL");
	params.push_back("-D SLG_OPENCL_KERNEL");
	params.push_back("-D RENDER_ENGINE_" + renderEngineType);
	params.push_back("-D PARAM_RAY_EPSILON_MIN=" + ToString(epsilonMin) + "f");
	params.push_back("-D PARAM_RAY_EPSILON_MAX=" + ToString(epsilonMax) + "f");

	const OpenCLDeviceDescription *oclDeviceDesc = dynamic_cast<const OpenCLDeviceDescription *>(intersectionDevice->GetDeviceDesc());
	if (oclDeviceDesc) {
		if (oclDeviceDesc->IsAMDPlatform())
			params.push_back("-D LUXCORE_AMD_OPENCL");
		else if (oclDeviceDesc->IsNVIDIAPlatform())
			params.push_back("-D LUXCORE_NVIDIA_OPENCL");
		else
			params.push_back("-D LUXCORE_GENERIC_OPENCL");
	}
}

string PathOCLBaseOCLRenderThread::GetKernelSources() {
	// Compile sources
	stringstream ssKernel;
	ssKernel <<
			// OpenCL LuxRays Types
			luxrays::ocl::KernelSource_luxrays_types <<
			luxrays::ocl::KernelSource_randomgen_types <<
			luxrays::ocl::KernelSource_uv_types <<
			luxrays::ocl::KernelSource_point_types <<
			luxrays::ocl::KernelSource_vector_types <<
			luxrays::ocl::KernelSource_normal_types <<
			luxrays::ocl::KernelSource_triangle_types <<
			luxrays::ocl::KernelSource_ray_types <<
			luxrays::ocl::KernelSource_bbox_types <<
			luxrays::ocl::KernelSource_epsilon_types <<
			luxrays::ocl::KernelSource_color_types <<
			luxrays::ocl::KernelSource_frame_types <<
			luxrays::ocl::KernelSource_matrix4x4_types <<
			luxrays::ocl::KernelSource_quaternion_types <<
			luxrays::ocl::KernelSource_transform_types <<
			luxrays::ocl::KernelSource_motionsystem_types <<
			luxrays::ocl::KernelSource_trianglemesh_types <<
			luxrays::ocl::KernelSource_exttrianglemesh_types <<
			// OpenCL LuxRays Funcs
			luxrays::ocl::KernelSource_randomgen_funcs <<
			luxrays::ocl::KernelSource_atomic_funcs <<
			luxrays::ocl::KernelSource_epsilon_funcs <<
			luxrays::ocl::KernelSource_utils_funcs <<
			luxrays::ocl::KernelSource_mc_funcs <<
			luxrays::ocl::KernelSource_vector_funcs <<
			luxrays::ocl::KernelSource_ray_funcs <<
			luxrays::ocl::KernelSource_bbox_funcs <<
			luxrays::ocl::KernelSource_color_funcs <<
			luxrays::ocl::KernelSource_frame_funcs <<
			luxrays::ocl::KernelSource_matrix4x4_funcs <<
			luxrays::ocl::KernelSource_quaternion_funcs <<
			luxrays::ocl::KernelSource_transform_funcs <<
			luxrays::ocl::KernelSource_motionsystem_funcs <<
			luxrays::ocl::KernelSource_triangle_funcs <<
			luxrays::ocl::KernelSource_exttrianglemesh_funcs <<
			// OpenCL SLG Types
			slg::ocl::KernelSource_sceneobject_types <<
			slg::ocl::KernelSource_scene_types <<
			slg::ocl::KernelSource_hitpoint_types <<
			slg::ocl::KernelSource_imagemap_types <<
			slg::ocl::KernelSource_mapping_types <<
			slg::ocl::KernelSource_texture_types <<
			slg::ocl::KernelSource_bsdf_types <<
			slg::ocl::KernelSource_material_types <<
			slg::ocl::KernelSource_volume_types <<
			slg::ocl::KernelSource_sampleresult_types <<
			slg::ocl::KernelSource_film_types <<
			slg::ocl::KernelSource_filter_types <<
			slg::ocl::KernelSource_sampler_types <<
			slg::ocl::KernelSource_camera_types <<
			slg::ocl::KernelSource_light_types <<
			slg::ocl::KernelSource_indexbvh_types <<
			slg::ocl::KernelSource_dlsc_types <<
			slg::ocl::KernelSource_elvc_types <<
			slg::ocl::KernelSource_pgic_types <<
			// OpenCL SLG Funcs
			slg::ocl::KernelSource_mortoncurve_funcs <<
			slg::ocl::KernelSource_evalstack_funcs <<
			slg::ocl::KernelSource_hitpoint_funcs << // Required by mapping funcs
			slg::ocl::KernelSource_mapping_funcs <<
			slg::ocl::KernelSource_imagemap_funcs <<
			slg::ocl::KernelSource_texture_bump_funcs <<
			slg::ocl::KernelSource_texture_noise_funcs <<
			slg::ocl::KernelSource_texture_blender_noise_funcs <<
			slg::ocl::KernelSource_texture_blender_noise_funcs2 <<
			slg::ocl::KernelSource_texture_blender_funcs <<
			slg::ocl::KernelSource_texture_abs_funcs <<
			slg::ocl::KernelSource_texture_bilerp_funcs <<
			slg::ocl::KernelSource_texture_blackbody_funcs <<
			slg::ocl::KernelSource_texture_brick_funcs <<
			slg::ocl::KernelSource_texture_clamp_funcs <<
			slg::ocl::KernelSource_texture_colordepth_funcs <<
			slg::ocl::KernelSource_texture_densitygrid_funcs <<
			slg::ocl::KernelSource_texture_fresnelcolor_funcs <<
			slg::ocl::KernelSource_texture_fresnelconst_funcs <<
			slg::ocl::KernelSource_texture_hitpoint_funcs <<
			slg::ocl::KernelSource_texture_hsv_funcs <<
			slg::ocl::KernelSource_texture_irregulardata_funcs <<
			slg::ocl::KernelSource_texture_triplanar_funcs <<
			slg::ocl::KernelSource_texture_others_funcs <<
			slg::ocl::KernelSource_texture_random_funcs <<
			slg::ocl::KernelSource_texture_funcs_evalops <<
			slg::ocl::KernelSource_texture_funcs;

	ssKernel <<
			slg::ocl::KernelSource_materialdefs_funcs_generic <<
			slg::ocl::KernelSource_materialdefs_funcs_default <<
			slg::ocl::KernelSource_materialdefs_funcs_thinfilmcoating <<
			slg::ocl::KernelSource_materialdefs_funcs_archglass <<
			slg::ocl::KernelSource_materialdefs_funcs_carpaint <<
			slg::ocl::KernelSource_materialdefs_funcs_clearvol <<
			slg::ocl::KernelSource_materialdefs_funcs_cloth <<
			slg::ocl::KernelSource_materialdefs_funcs_disney <<
			slg::ocl::KernelSource_materialdefs_funcs_glass <<
			slg::ocl::KernelSource_materialdefs_funcs_glossy2 <<
			slg::ocl::KernelSource_materialdefs_funcs_glossycoating <<
			slg::ocl::KernelSource_materialdefs_funcs_glossytranslucent <<
			slg::ocl::KernelSource_materialdefs_funcs_heterogeneousvol <<
			slg::ocl::KernelSource_materialdefs_funcs_homogeneousvol <<
			slg::ocl::KernelSource_materialdefs_funcs_matte <<
			slg::ocl::KernelSource_materialdefs_funcs_matte_translucent <<
			slg::ocl::KernelSource_materialdefs_funcs_metal2 <<
			slg::ocl::KernelSource_materialdefs_funcs_mirror <<
			slg::ocl::KernelSource_materialdefs_funcs_mix <<
			slg::ocl::KernelSource_materialdefs_funcs_null <<
			slg::ocl::KernelSource_materialdefs_funcs_roughglass <<
			slg::ocl::KernelSource_materialdefs_funcs_roughmatte_translucent <<
			slg::ocl::KernelSource_materialdefs_funcs_velvet <<
			slg::ocl::KernelSource_material_funcs_evalops <<
			slg::ocl::KernelSource_material_funcs;

	ssKernel <<
			slg::ocl::KernelSource_pathdepthinfo_types <<
			slg::ocl::KernelSource_pathvolumeinfo_types <<
			slg::ocl::KernelSource_pathinfo_types <<
			slg::ocl::KernelSource_pathtracer_types <<
			// PathOCL types
			slg::ocl::KernelSource_pathoclbase_datatypes;

	ssKernel <<
			slg::ocl::KernelSource_bsdfutils_funcs << // Must be before volumeinfo_funcs
			slg::ocl::KernelSource_volume_funcs <<
			slg::ocl::KernelSource_pathdepthinfo_funcs <<
			slg::ocl::KernelSource_pathvolumeinfo_funcs <<
			slg::ocl::KernelSource_pathinfo_funcs <<
			slg::ocl::KernelSource_camera_funcs <<
			slg::ocl::KernelSource_dlsc_funcs <<
			slg::ocl::KernelSource_elvc_funcs <<
			slg::ocl::KernelSource_lightstrategy_funcs <<
			slg::ocl::KernelSource_light_funcs <<
			slg::ocl::KernelSource_filter_funcs <<
			slg::ocl::KernelSource_sampleresult_funcs <<
			slg::ocl::KernelSource_filmdenoiser_funcs <<
			slg::ocl::KernelSource_film_funcs <<
			slg::ocl::KernelSource_varianceclamping_funcs <<
			slg::ocl::KernelSource_sampler_random_funcs <<
			slg::ocl::KernelSource_sampler_sobol_funcs <<
			slg::ocl::KernelSource_sampler_metropolis_funcs <<
			slg::ocl::KernelSource_sampler_tilepath_funcs <<
			slg::ocl::KernelSource_sampler_funcs <<
			slg::ocl::KernelSource_bsdf_funcs <<
			slg::ocl::KernelSource_scene_funcs <<
			slg::ocl::KernelSource_pgic_funcs <<
			// PathOCL Funcs
			slg::ocl::KernelSource_pathoclbase_funcs <<
			slg::ocl::KernelSource_pathoclbase_kernels_micro;

	return ssKernel.str();
}

void PathOCLBaseOCLRenderThread::InitKernels() {
	//--------------------------------------------------------------------------
	// Compile kernels
	//--------------------------------------------------------------------------

	const double tStart = WallClockTime();

	// A safety check
	switch (intersectionDevice->GetAccelerator()->GetType()) {
		case ACCEL_BVH:
			break;
		case ACCEL_MBVH:
			break;
		case ACCEL_EMBREE:
			throw runtime_error("EMBREE accelerator is not supported in PathOCLBaseRenderThread::InitKernels()");
		default:
			throw runtime_error("Unknown accelerator in PathOCLBaseRenderThread::InitKernels()");
	}

	vector<string> kernelsParameters;
	GetKernelParamters(kernelsParameters, intersectionDevice,
			RenderEngine::RenderEngineType2String(renderEngine->GetType()),
			MachineEpsilon::GetMin(), MachineEpsilon::GetMax());

	const string kernelSource = GetKernelSources();

	if (renderEngine->writeKernelsToFile) {
		// Some debug code to write the OpenCL kernel source to a file
		const string kernelFileName = "kernel_source_device_" + ToString(threadIndex) + ".cl";
		ofstream kernelFile(kernelFileName.c_str());
		string kernelDefs = oclKernelPersistentCache::ToOptsString(kernelsParameters);
		boost::replace_all(kernelDefs, "-D", "\n#define");
		boost::replace_all(kernelDefs, "=", " ");
		kernelFile << kernelDefs << endl << endl << kernelSource << endl;
		kernelFile.close();
	}

	if ((renderEngine->additionalOpenCLKernelOptions.size() > 0) &&
			(intersectionDevice->GetDeviceDesc()->GetType() & DEVICE_TYPE_OPENCL_ALL))
		kernelsParameters.insert(kernelsParameters.end(), renderEngine->additionalOpenCLKernelOptions.begin(), renderEngine->additionalOpenCLKernelOptions.end());
	if ((renderEngine->additionalCUDAKernelOptions.size() > 0) &&
			(intersectionDevice->GetDeviceDesc()->GetType() & DEVICE_TYPE_CUDA_ALL))
		kernelsParameters.insert(kernelsParameters.end(), renderEngine->additionalCUDAKernelOptions.begin(), renderEngine->additionalCUDAKernelOptions.end());

	// Build the kernel source/parameters hash
	const string newKernelSrcHash = oclKernelPersistentCache::HashString(oclKernelPersistentCache::ToOptsString(kernelsParameters))
			+ "-" +
			oclKernelPersistentCache::HashString(kernelSource);
	if (newKernelSrcHash == kernelSrcHash) {
		// There is no need to re-compile the kernel
		return;
	} else
		kernelSrcHash = newKernelSrcHash;

	SLG_LOG("[PathOCLBaseRenderThread::" << threadIndex << "] Compiling kernels ");

	HardwareDeviceProgram *program = nullptr;
	intersectionDevice->CompileProgram(&program, kernelsParameters, kernelSource, "PathOCL kernel");

	// Film clear kernel
	CompileKernel(intersectionDevice, program, &filmClearKernel, &filmClearWorkGroupSize, "Film_Clear");

	// Init kernel

	CompileKernel(intersectionDevice, program, &initSeedKernel, &initWorkGroupSize, "InitSeed");
	CompileKernel(intersectionDevice, program, &initKernel, &initWorkGroupSize, "Init");

	// AdvancePaths kernel (Micro-Kernels)

	size_t workGroupSize;
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_RT_NEXT_VERTEX, &advancePathsWorkGroupSize,
			"AdvancePaths_MK_RT_NEXT_VERTEX");
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_HIT_NOTHING, &workGroupSize,
			"AdvancePaths_MK_HIT_NOTHING");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_HIT_OBJECT, &workGroupSize,
			"AdvancePaths_MK_HIT_OBJECT");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_RT_DL, &workGroupSize,
			"AdvancePaths_MK_RT_DL");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_DL_ILLUMINATE, &workGroupSize,
			"AdvancePaths_MK_DL_ILLUMINATE");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_DL_SAMPLE_BSDF, &workGroupSize,
			"AdvancePaths_MK_DL_SAMPLE_BSDF");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_GENERATE_NEXT_VERTEX_RAY, &workGroupSize,
			"AdvancePaths_MK_GENERATE_NEXT_VERTEX_RAY");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_SPLAT_SAMPLE, &workGroupSize,
			"AdvancePaths_MK_SPLAT_SAMPLE");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_NEXT_SAMPLE, &workGroupSize,
			"AdvancePaths_MK_NEXT_SAMPLE");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	CompileKernel(intersectionDevice, program, &advancePathsKernel_MK_GENERATE_CAMERA_RAY, &workGroupSize,
			"AdvancePaths_MK_GENERATE_CAMERA_RAY");
	advancePathsWorkGroupSize = Min(advancePathsWorkGroupSize, workGroupSize);
	SLG_LOG("[PathOCLBaseRenderThread::" << threadIndex << "] AdvancePaths_MK_* workgroup size: " << advancePathsWorkGroupSize);

	const double tEnd = WallClockTime();
	SLG_LOG("[PathOCLBaseRenderThread::" << threadIndex << "] Kernels compilation time: " << int((tEnd - tStart) * 1000.0) << "ms");

	delete program;
}

void PathOCLBaseOCLRenderThread::SetInitKernelArgs(const u_int filmIndex) {
	// initSeedKernel kernel
	u_int argIndex = 0;
	intersectionDevice->SetKernelArg(initSeedKernel, argIndex++, tasksBuff);
	intersectionDevice->SetKernelArg(initSeedKernel, argIndex++, renderEngine->seedBase + threadIndex * renderEngine->taskCount);

	// initKernel kernel
	argIndex = 0;
	intersectionDevice->SetKernelArg(initKernel, argIndex++, taskConfigBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, tasksBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, tasksDirectLightBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, tasksStateBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, taskStatsBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, samplerSharedDataBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, samplesBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, sampleDataBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, sampleResultsBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, eyePathInfosBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, pixelFilterBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, raysBuff);
	intersectionDevice->SetKernelArg(initKernel, argIndex++, cameraBuff);

	// Film parameters
	argIndex = threadFilms[filmIndex]->SetFilmKernelArgs(intersectionDevice, initKernel, argIndex);

	initKernelArgsCount = argIndex;
}

void PathOCLBaseOCLRenderThread::SetAdvancePathsKernelArgs(HardwareDeviceKernel *advancePathsKernel, const u_int filmIndex) {
	CompiledScene *cscene = renderEngine->compiledScene;

	u_int argIndex = 0;
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, taskConfigBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, tasksBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, tasksDirectLightBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, tasksStateBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, taskStatsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pixelFilterBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, samplerSharedDataBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, samplesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, sampleDataBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, sampleResultsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, eyePathInfosBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, directLightVolInfosBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, raysBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, hitsBuff);

	// Film parameters
	argIndex = threadFilms[filmIndex]->SetFilmKernelArgs(intersectionDevice, advancePathsKernel, argIndex);

	// Scene parameters
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->worldBSphere.center.x);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->worldBSphere.center.y);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->worldBSphere.center.z);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->worldBSphere.rad);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, materialsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, materialEvalOpsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, materialEvalStackBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->maxMaterialEvalStackSize);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, texturesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, textureEvalOpsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, textureEvalStackBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->maxTextureEvalStackSize);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, scnObjsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, meshDescsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, vertsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, normalsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, triNormalsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, uvsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, colsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, alphasBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, vertexAOVBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, triAOVBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, trianglesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, interpolatedTransformsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cameraBuff);
	// Lights
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, lightsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, envLightIndicesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, (u_int)cscene->envLightIndices.size());
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, lightIndexOffsetByMeshIndexBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, lightIndexByTriIndexBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, envLightDistributionsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, lightsDistributionBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, infiniteLightSourcesDistributionBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, dlscAllEntriesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, dlscDistributionsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, dlscBVHNodesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->dlscRadius2);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->dlscNormalCosAngle);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, elvcAllEntriesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, elvcDistributionsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, elvcTileDistributionOffsetsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, elvcBVHNodesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->elvcRadius2);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->elvcNormalCosAngle);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->elvcTilesXCount);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->elvcTilesYCount);

	// Images
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, imageMapDescsBuff);
	for (u_int i = 0; i < 8; ++i) {
		if (i < imageMapsBuff.size())
			intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, imageMapsBuff[i]);
		else
			intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, nullptr);
	}

	// PhotonGI cache
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pgicRadiancePhotonsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, cscene->pgicLightGroupCounts);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pgicRadiancePhotonsValuesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pgicRadiancePhotonsBVHNodesBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pgicCausticPhotonsBuff);
	intersectionDevice->SetKernelArg(advancePathsKernel, argIndex++, pgicCausticPhotonsBVHNodesBuff);
}

void PathOCLBaseOCLRenderThread::SetAllAdvancePathsKernelArgs(const u_int filmIndex) {
	if (advancePathsKernel_MK_RT_NEXT_VERTEX)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_RT_NEXT_VERTEX, filmIndex);
	if (advancePathsKernel_MK_HIT_NOTHING)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_HIT_NOTHING, filmIndex);
	if (advancePathsKernel_MK_HIT_OBJECT)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_HIT_OBJECT, filmIndex);
	if (advancePathsKernel_MK_RT_DL)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_RT_DL, filmIndex);
	if (advancePathsKernel_MK_DL_ILLUMINATE)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_DL_ILLUMINATE, filmIndex);
	if (advancePathsKernel_MK_DL_SAMPLE_BSDF)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_DL_SAMPLE_BSDF, filmIndex);
	if (advancePathsKernel_MK_GENERATE_NEXT_VERTEX_RAY)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_GENERATE_NEXT_VERTEX_RAY, filmIndex);
	if (advancePathsKernel_MK_SPLAT_SAMPLE)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_SPLAT_SAMPLE, filmIndex);
	if (advancePathsKernel_MK_NEXT_SAMPLE)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_NEXT_SAMPLE, filmIndex);
	if (advancePathsKernel_MK_GENERATE_CAMERA_RAY)
		SetAdvancePathsKernelArgs(advancePathsKernel_MK_GENERATE_CAMERA_RAY, filmIndex);
}

void PathOCLBaseOCLRenderThread::SetKernelArgs() {
	// Set OpenCL kernel arguments

	// OpenCL kernel setArg() is the only non thread safe function in OpenCL 1.1 so
	// I need to use a mutex here

	boost::unique_lock<boost::mutex> lock(renderEngine->setKernelArgsMutex);

	//--------------------------------------------------------------------------
	// advancePathsKernels
	//--------------------------------------------------------------------------

	SetAllAdvancePathsKernelArgs(0);

	//--------------------------------------------------------------------------
	// initKernel
	//--------------------------------------------------------------------------

	SetInitKernelArgs(0);
}

void PathOCLBaseOCLRenderThread::EnqueueAdvancePathsKernel() {
	const u_int taskCount = renderEngine->taskCount;

	// Micro kernels version
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_RT_NEXT_VERTEX,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_HIT_NOTHING,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_HIT_OBJECT,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_RT_DL,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_DL_ILLUMINATE,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_DL_SAMPLE_BSDF,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_GENERATE_NEXT_VERTEX_RAY,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_SPLAT_SAMPLE,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_NEXT_SAMPLE,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
	intersectionDevice->EnqueueKernel(advancePathsKernel_MK_GENERATE_CAMERA_RAY,
			HardwareDeviceRange(taskCount), HardwareDeviceRange(advancePathsWorkGroupSize));
}

#endif

max_stars_repo_name: mbrukman/LuxCore

max_issues_repo_name: mbrukman/LuxCore

max_forks_repo_name: mbrukman/LuxCore


text:
/** @file
    @brief Implementation

    @date 2015

    @author
    Sensics, Inc.
    <http://sensics.com/osvr>
*/

// Copyright 2015 Sensics, Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//        http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#define OSVR_HAVE_BOOST

// Internal Includes
#include "ImageSources/ImageSource.h"
#include "ImageSources/ImageSourceFactories.h"
#include <osvr/Util/MiniArgsHandling.h>

// Library/third-party includes
#include <boost/lexical_cast.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <vrpn_SerialPort.h>

// Standard includes
#include <chrono>
#include <cstddef> // for std::size_t
#include <cstdint>
#include <fstream>
#include <functional>
#include <iostream>
#include <memory>
#include <random>
#include <sstream>
#include <thread>

using SerialBufType = const unsigned char;

/// Utility function to take in a (typically) string literal and write it,
/// without its null terminator, to the serial port, optionally repeatedly.
/// (Size deduced by template)
template <typename T, std::size_t N>
inline void writeStringWithoutNullTerminator(vrpn_SerialPort &port, T (&buf)[N],
                                             std::size_t repetitions = 1) {
    static_assert(sizeof(T) == 1, "Can only pass character/byte arrays");
    std::size_t len = N;
    if ('\0' == buf[len - 1]) {
        len--;
    }
    for (std::size_t i = 0; i < repetitions; ++i) {
        port.write(reinterpret_cast<SerialBufType *>(&(buf[0])), len);
    }
}

/// Utility function to take in a string literal and write it,
/// without its null terminator, followed by a newline, to the serial port,
/// optionally repeatedly. (Size deduced by template)
template <typename T, std::size_t N>
inline void writeCommand(vrpn_SerialPort &port, T (&buf)[N],
                         std::size_t repetitions = 1) {
    for (std::size_t i = 0; i < repetitions; ++i) {
        writeStringWithoutNullTerminator(port, buf);
        writeStringWithoutNullTerminator(port, "\n");
    }
}

/// Utility function to take in a binary array and write it verbatim to the
/// serial port, optionally repeatedly. (Size deduced by template)
template <std::size_t N>
inline void writeBinaryArray(vrpn_SerialPort &port, SerialBufType (&buf)[N],
                             std::size_t repetitions = 1) {
    std::size_t len = N;
    for (std::size_t i = 0; i < repetitions; ++i) {
        port.write(reinterpret_cast<SerialBufType *>(&(buf[0])), N);
    }
}

/// Abstract base class for ways of controlling an IR LED through a serial port.
class SerialIRLED {
  public:
    /// Call to cause an IR LED pulse to be emitted.
    void trigger() { handle_trigger(getPort()); }
    virtual ~SerialIRLED() {}

  protected:
    SerialIRLED(vrpn_SerialPort &port) : port_(std::ref(port)) {}
    virtual void handle_trigger(vrpn_SerialPort &port) = 0;

    vrpn_SerialPort &getPort() { return port_.get(); }

  private:
    std::reference_wrapper<vrpn_SerialPort> port_;
};
using SerialIRLEDPtr = std::unique_ptr<SerialIRLED>;
using SerialLEDControllerFactory =
    std::function<SerialIRLEDPtr(vrpn_SerialPort &)>;

/// A class for an LED connected to an Arduino running a sketch that pulses an
/// LED on then off on receiving a single character over serial.
class ArduinoLED : public SerialIRLED {
  public:
    static SerialIRLEDPtr make(vrpn_SerialPort &port) {
        SerialIRLEDPtr ret(new ArduinoLED(port));
        return ret;
    }

  private:
    ArduinoLED(vrpn_SerialPort &port) : SerialIRLED(port) {}
    void handle_trigger(vrpn_SerialPort &port) override {
        // write any byte to trigger a pulse
        writeStringWithoutNullTerminator(port, "a");
        port.drain_output_buffer();
    }
};

/// A class for an LED connected to a Bus Pirate's AUX pin (with resistor),
/// using the binary bit-bang mode to turn it on then off. Bus pirate should be
/// in normal "terminal" mode before starting.
class BusPirateAuxLED : public SerialIRLED {
  public:
    static SerialIRLEDPtr make(vrpn_SerialPort &port) {
        SerialIRLEDPtr ret;
        try {
            ret.reset(new BusPirateAuxLED(port));
        } catch (std::exception &e) {
            std::cerr << "Error setting up bus pirate: " << e.what()
                      << std::endl;
        }
        return ret;
    }

    virtual ~BusPirateAuxLED() {
        const std::uint8_t resetToTerminal[] = {RESET_TO_TERMINAL};
        writeBinaryArray(getPort(), resetToTerminal);
        getPort().drain_output_buffer();
    }

  private:
    BusPirateAuxLED(vrpn_SerialPort &port) : SerialIRLED(port) {
        bool success = enterBitBangMode(port);
        if (success) {
            std::cout << "Bus Pirate successfully entered binary bit-bang mode!"
                      << std::endl;
        } else {
            std::cout << "Couldn't get bus pirate into binary bit-bang mode: "
                         "check that nothing else is using the port and try "
                         "resetting it."
                      << std::endl;
            throw std::runtime_error("Could not transition bus pirate to "
                                     "binary bit-bang mode: check that nothing "
                                     "else is using the port and try resetting "
                                     "it.");
        }

        const std::uint8_t configPortModes[] = {PORT_MODE, ONLY_POWER_ON};
        writeBinaryArray(port, configPortModes);
        // wait until everything has been sent.
        port.drain_output_buffer();
        // throw away replies.
        port.flush_input_buffer();
    }
    static const std::uint8_t ONE = 0x01;
    static const std::uint8_t ZERO = 0x00;

    // 01001111: aka 0x4f configure everything except AUX as input.
    static const std::uint8_t PORT_MODE = 0x4f;
    // 11000000: aka 0xc0 turn on the power only
    static const std::uint8_t ONLY_POWER_ON = 0xc0;
    // 11010000: aka 0xd0 turn on the power and AUX
    static const std::uint8_t POWER_AND_AUX_ON = 0xd0;
    // 00001111: aka 0x0f reset bus pirate to terminal mode.
    static const std::uint8_t RESET_TO_TERMINAL = 0x0f;
    bool enterBitBangMode(vrpn_SerialPort &port) {

        SerialBufType enter[] = "\n";
        // throw away anything we got to read.
        port.flush_input_buffer();
        // write at least 10 enters to get out of any menus.
        writeStringWithoutNullTerminator(port, enter, 10);
        // send reset command
        writeCommand(port, "#");
        // wait until everything has been sent.
        port.drain_output_buffer();
        std::cout << "Resetting bus pirate" << std::endl;
        std::this_thread::sleep_for(std::chrono::seconds(1));
        // throw away anything we got to read.
        port.flush_input_buffer();

        bool success = false;
        const std::uint8_t bitBangReset[] = {ZERO};
        static const auto desiredReply = "BBIO1";
        writeBinaryArray(port, bitBangReset, 10);
        for (int i = 0; i < 40; ++i) {
            writeBinaryArray(port, bitBangReset);
            std::this_thread::sleep_for(std::chrono::milliseconds(5));
            // try more times to get BBIO1
            auto response = port.read_available_characters(5);
            if (desiredReply == response) {
                return true;
            } else if (!response.empty()) {
                std::cout << "When waiting for " << desiredReply
                          << " Bus Pirate sent us " << response << std::endl;
            }
        }
        return false;
    }
    void handle_trigger(vrpn_SerialPort &port) override {

        const std::uint8_t turnOn[] = {POWER_AND_AUX_ON};

        writeBinaryArray(port, turnOn);
        // wait until everything has been sent.
        port.drain_output_buffer();
        // throw away replies.
        port.flush_input_buffer();

        std::this_thread::sleep_for(std::chrono::milliseconds(5));

        const std::uint8_t turnOff[] = {ONLY_POWER_ON};

        writeBinaryArray(port, turnOff);
        // wait until everything has been sent.
        port.drain_output_buffer();
        // throw away replies.
        port.flush_input_buffer();
    }
};

static const std::string windowNameAndInstructions(
    "OSVR tracking camera preview | q or esc to quit");
static const auto BUS_PIRATE_ARGS = {"--bp", "--buspirate", "-b"};
static const auto ARDUINO_ARGS = {"--arduino", "-a"};
static const auto DEVICE_ARG = "--dev";
static const auto DEFAULT_DEVICE = "COM1";
static const auto RATE_ARG = "--rate";
static const auto DEFAULT_RATE = 115200;
static const auto MEASUREMENTS_ARG = "--measurements";
static const auto DEFAULT_MEASUREMENTS = 20;

static const double CUTOFF_VALUE = 200.;

struct Config {
    SerialLEDControllerFactory controllerFactory;
    std::string device = DEFAULT_DEVICE;
    std::uint32_t rate = DEFAULT_RATE;
    std::uint32_t measurements = DEFAULT_MEASUREMENTS;
} g_config;

int main(int argc, char *argv[]) {
    {
        auto withUsageError = [] {
            /// @todo add usage
            return 1;
        };
        using namespace osvr::util::args;
        auto args = makeArgList(argc, argv);
        if (handle_has_any_iswitch_of(args, BUS_PIRATE_ARGS)) {
            std::cout << "Will use the Bus Pirate, IR LED on AUX, controller."
                      << std::endl;
            g_config.controllerFactory = BusPirateAuxLED::make;
        }
        if (handle_has_any_iswitch_of(args, ARDUINO_ARGS)) {
            if (g_config.controllerFactory) {
                std::cerr << "Got an arduino arg, but already was specified to "
                             "use bus pirate!"
                          << std::endl;
                return withUsageError();
            }
            g_config.controllerFactory = ArduinoLED::make;
        }

        if (!g_config.controllerFactory) {
            std::cerr << "No controller type specified! Assuming bus pirate."
                      << std::endl;
            g_config.controllerFactory = BusPirateAuxLED::make;
        }

        handle_value_arg(
            args, [](std::string const &a) { return a == DEVICE_ARG; },
            [&](std::string const &val) {
                std::cout << "Setting serial device to " << val << std::endl;
                g_config.device = val;
            });

        handle_value_arg(
            args, [](std::string const &a) { return a == RATE_ARG; },
            [&](std::string const &val) {
                g_config.rate = boost::lexical_cast<std::uint32_t>(val);
                std::cout << "Setting serial rate to " << g_config.rate
                          << std::endl;
            });

        handle_value_arg(
            args, [](std::string const &a) { return a == MEASUREMENTS_ARG; },
            [&](std::string const &val) {
                g_config.measurements = boost::lexical_cast<std::uint32_t>(val);
                std::cout << "Setting number of measurements to "
                          << g_config.measurements << std::endl;
            });
    }
#ifdef _WIN32
    auto cam = osvr::vbtracker::openHDKCameraDirectShow();
#else
    std::cerr << "Warning: Just using OpenCV to open Camera #0, which may not "
                 "be the tracker camera."
              << std::endl;
    auto cam = osvr::vbtracker::openOpenCVCamera(0);
#endif
    if (!cam || !cam->ok()) {
        std::cerr << "Couldn't find, open, or read from the OSVR HDK tracking "
                     "camera.\n"
                  << "Press enter to exit." << std::endl;
        std::cin.ignore();
        return -1;
    }

    auto FRAME_DISPLAY_STRIDE = 3u;
    cam->grab();

    auto frame = cv::Mat{};
    auto grayFrame = cv::Mat{};

    vrpn_SerialPort comPort(g_config.device.c_str(), g_config.rate);
    if (!comPort.is_open()) {
        std::cerr << "Couldn't open com port!" << std::endl;
        return -1;
    }

    std::unique_ptr<SerialIRLED> LEDController(
        g_config.controllerFactory(comPort));

    if (!LEDController) {
        std::cerr << "Could not start LEDController, exiting." << std::endl;
        return 1;
    }

    enum class State { StartedTrigger, SawFlash, AwaitingNewTrigger };

    using osvr::util::time::TimeValue;
    TimeValue triggerTime = {};
    double peakVal = 0.;

    State s = State::AwaitingNewTrigger;

    std::random_device rd;
    std::mt19937 mt(rd());
    // distribution of how many frames we'll wait before triggering a new pulse.
    std::uniform_int_distribution<int> dist(5, 50);
    auto generateNewCountdown = [&mt, &dist] { return dist(mt); };
    int countdownToTrigger = generateNewCountdown();

    std::uint32_t samples = 0;
    std::ofstream os("latency.csv");
    do {
        TimeValue tv;
        cam->retrieve(frame, grayFrame, tv);
        switch (s) {
        case State::AwaitingNewTrigger: {
            if (countdownToTrigger == 0) {
                // trigger timing
                s = State::StartedTrigger;
                peakVal = 0.;
                LEDController->trigger();
                triggerTime = osvr::util::time::getNow();
            } else {
                countdownToTrigger--;
            }
            break;
        }
        case State::StartedTrigger: {
            auto now = osvr::util::time::getNow();
            double minVal, maxVal;
            cv::minMaxIdx(grayFrame, &minVal, &maxVal);
            // if (maxVal < peakVal) {
            if (maxVal > CUTOFF_VALUE) {
                osvrTimeValueDifference(&now, &triggerTime);
                osvrTimeValueDifference(&tv, &triggerTime);
                if (tv.microseconds > 0) {
                    // we got it
                    s = State::SawFlash;
                    std::cout << "   Current: " << maxVal << "\n";
                    std::cout << "Latency from trigger to sample time: "
                              << tv.microseconds << "us\n";
                    std::cout << "Latency from trigger to retrieval time: "
                              << now.microseconds << "us\n";
                    cv::imwrite("triggered.png", frame);
                    samples++;
                    os << tv.microseconds << std::endl;
                } else {
                    std::cout << "Got a spurious flash, negative trigger to "
                                 "sample duration"
                              << std::endl;
                }
            }
            break;
        }
        case State::SawFlash: {
            // wait for the flash to disappear
            double minVal, maxVal;
            cv::minMaxIdx(grayFrame, &minVal, &maxVal);
            if (maxVal < CUTOFF_VALUE) {
                // OK, ready to go for a new trigger at some point.
                s = State::AwaitingNewTrigger;
                countdownToTrigger = generateNewCountdown();
                std::cout << "Will wait " << countdownToTrigger
                          << " frames before triggering again\n";
            }
            break;
        }
        }
    } while (cam->grab() && samples < g_config.measurements);

    LEDController.reset();

    return 0;
}

max_stars_repo_name: ethanpeng/OSVR-Core

max_issues_repo_name: ethanpeng/OSVR-Core

max_forks_repo_name: ethanpeng/OSVR-Core


text:
/* CirKit: A circuit toolkit
 * Copyright (C) 2009-2015  University of Bremen
 * Copyright (C) 2015-2017  EPFL
 *
 * Permission is hereby granted, free of charge, to any person
 * obtaining a copy of this software and associated documentation
 * files (the "Software"), to deal in the Software without
 * restriction, including without limitation the rights to use,
 * copy, modify, merge, publish, distribute, sublicense, and/or sell
 * copies of the Software, and to permit persons to whom the
 * Software is furnished to do so, subject to the following
 * conditions:
 *
 * The above copyright notice and this permission notice shall be
 * included in all copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
 * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
 * OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
 * NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
 * HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
 * WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
 * OTHER DEALINGS IN THE SOFTWARE.
 */

/**
 * @file cut_enumeration.hpp
 *
 * @brief Enumerates cuts
 *
 * @author Mathias Soeken
 * @since  2.0
 */

#ifndef CUT_ENUMERATION_HPP
#define CUT_ENUMERATION_HPP

#include <map>
#include <stack>
#include <vector>

#include <boost/assign/std/vector.hpp>
#include <boost/range/iterator_range.hpp>
#include <core/properties.hpp>
#include <core/utils/bitset_utils.hpp>
#include <core/utils/timer.hpp>
#include <classical/utils/cut_enumeration_traits.hpp>
#include <classical/utils/truth_table_utils.hpp>

using namespace boost::assign;

namespace cirkit
{

/******************************************************************************
 * Cut implementations                                                        *
 ******************************************************************************/

template<typename CutType>
CutType cut_from_index( unsigned index, unsigned n )
{
  assert( false );
}

template<typename CutType>
CutType empty_cut( unsigned n )
{
  assert( false );
}

template<typename CutType>
unsigned cut_size( const CutType& cut )
{
  assert( false );
}

template<typename CutType>
CutType cut_union( const CutType& cut1, const CutType& cut2 )
{
  assert( false );
}

template<typename CutType>
CutType& cut_insert( CutType& cut, unsigned index )
{
  assert( false );
}

template<typename CutType>
void foreach_node_in_cut( const CutType& cut, const std::function<void(unsigned)>&& f )
{
  assert( false );
}

template<typename CutType>
std::ostream& print_cut( std::ostream& os, const CutType& cut )
{
  assert( false );
}

template<typename CutsType>
CutsType cuts_from_index( unsigned index, unsigned n )
{
  return {cut_from_index<typename CutsType::value_type>( index, n )};
}

template<typename CutsType>
CutsType empty_cuts( unsigned n )
{
  return {empty_cut<typename CutsType::value_type>( n )};
}

/* structural_cut */
using structural_cut = std::vector<boost::dynamic_bitset<>>;

template<>
boost::dynamic_bitset<> cut_from_index( unsigned index, unsigned n );

template<>
boost::dynamic_bitset<> empty_cut( unsigned n );

template<>
unsigned cut_size<>( const boost::dynamic_bitset<>& cut );

template<>
boost::dynamic_bitset<> cut_union<>( const boost::dynamic_bitset<>& cut1, const boost::dynamic_bitset<>& cut2 );

template<>
boost::dynamic_bitset<>& cut_insert<>( boost::dynamic_bitset<>& cut, unsigned index );

template<>
void foreach_node_in_cut<>( const boost::dynamic_bitset<>& cut, const std::function<void(unsigned)>&& f );

template<>
std::ostream& print_cut<>( std::ostream& os, const boost::dynamic_bitset<>& cut );

/* std::set */
template<>
std::set<unsigned> cut_from_index( unsigned index, unsigned n );

template<>
std::set<unsigned> empty_cut( unsigned n );

template<>
unsigned cut_size<>( const std::set<unsigned>& cut );

template<>
std::set<unsigned> cut_union<>( const std::set<unsigned>& cut1, const std::set<unsigned>& cut2 );

template<>
std::set<unsigned>& cut_insert<>( std::set<unsigned>& cut, unsigned index );

template<>
void foreach_node_in_cut<>( const std::set<unsigned>& cut, const std::function<void(unsigned)>&& f );

template<>
std::ostream& print_cut<>( std::ostream& os, const std::set<unsigned>& cut );

/* std::vector */
template<>
std::vector<unsigned> cut_from_index( unsigned index, unsigned n );

template<>
std::vector<unsigned> empty_cut( unsigned n );

template<>
unsigned cut_size<>( const std::vector<unsigned>& cut );

template<>
std::vector<unsigned> cut_union<>( const std::vector<unsigned>& cut1, const std::vector<unsigned>& cut2 );

template<>
std::vector<unsigned>& cut_insert<>( std::vector<unsigned>& cut, unsigned index );

template<>
void foreach_node_in_cut<>( const std::vector<unsigned>& cut, const std::function<void(unsigned)>&& f );

template<>
std::ostream& print_cut<>( std::ostream& os, const std::vector<unsigned>& cut );

namespace detail
{

/******************************************************************************
 * Simulation algorithm for cut enumeration                                   *
 ******************************************************************************/

template<typename T, typename CutsType>
class structural_cut_enumeration_simulator : public cut_enumeration_circuit_traits<T>::template simulator<CutsType>
{
public:
  using node_t = typename circuit_traits<T>::node;
  using cut_t  = typename CutsType::value_type;

public:
  structural_cut_enumeration_simulator( const T& circ, unsigned k, bool include_constant, const boost::dynamic_bitset<>& cut_boundary )
    : circ( circ ),
      n( boost::num_vertices( circ ) ),
      include_constant( include_constant ),
      k( k ),
      cut_boundary( cut_boundary )
  {
  }

  structural_cut_enumeration_simulator( const T& circ, unsigned k, bool include_constant )
    : structural_cut_enumeration_simulator( circ, k, include_constant, boost::dynamic_bitset<>( boost::num_vertices( circ ) ) )
  {
  }

  CutsType get_input( const node_t& node, const std::string& name, unsigned pos, const T& circ ) const
  {
    return cuts_from_index<CutsType>( node, n );
  }

  CutsType get_constant() const
  {
    if ( include_constant )
    {
      return cuts_from_index<CutsType>( boost::get_property( circ, boost::graph_name ).constant, n );
    }
    else
    {
      return empty_cuts<CutsType>( n );
    }
  }

  CutsType invert( const CutsType& v ) const { return v; }

  CutsType and_op( const node_t& node, const CutsType& v1, const CutsType& v2 ) const
  {
    CutsType cuts;

    if ( !cut_boundary.test( node ) )
    {
      for ( const auto& c1 : v1 )
      {
        for ( const auto& c2 : v2 )
        {
          auto new_cut = cut_union( c1, c2 );
          if ( cut_size( new_cut ) > k ) continue;
          cuts += new_cut;
        }
      }
    }

    cuts += cut_from_index<cut_t>( node, n );

    return cuts;
  }

  CutsType maj_op( const node_t& node, const CutsType& v1, const CutsType& v2, const CutsType& v3 ) const
  {
    CutsType cuts;

    if ( !cut_boundary.test( node ) )
    {
      for ( const auto& c1 : v1 )
      {
        for ( const auto& c2 : v2 )
        {
          for ( const auto& c3 : v3 )
          {
            auto new_cut = cut_union( c1, cut_union( c2, c3 ) );
            if ( cut_size( new_cut ) > k ) continue;
            cuts += new_cut;
          }
        }
      }
    }

    cuts += cut_from_index<cut_t>( node, n );

    return cuts;
  }

  bool terminate( const node_t& node, const T& circ ) const
  {
    return cut_boundary.test( node );
  }

private:
  const T&                circ;
  unsigned                n;
  bool                    include_constant;
  unsigned                k;
  boost::dynamic_bitset<> cut_boundary;
};

template<typename T, typename CutsType>
class structural_cut_cone_simulator : public cut_enumeration_circuit_traits<T>::template simulator<typename CutsType::value_type>
{
public:
  using node_t = typename circuit_traits<T>::node;
  using cut_t  = typename CutsType::value_type;

public:
  structural_cut_cone_simulator( unsigned n ) : n( n ) {}

  cut_t get_input( const node_t& node, const std::string& name, unsigned pos, const T& circ ) const
  {
    return cut_from_index<cut_t>( node, n );
  }

  cut_t get_constant() const
  {
    return cut_from_index<cut_t>( 0u, n );
  }

  cut_t invert( const cut_t& v ) const { return v; }

  cut_t and_op( const node_t& node, const cut_t& v1, const cut_t& v2 ) const
  {
    auto cut = cut_union( v1, v2 );
    return cut_insert( cut, node );
  }

  cut_t maj_op( const node_t& node, const cut_t& v1, const cut_t& v2, const cut_t& v3 ) const
  {
    auto cut = cut_union( v1, cut_union( v2, v3 ) );
    return cut_insert( cut, node );
  }

private:
  unsigned n;
};

template<typename T>
class structural_cut_depth_simulator : public cut_enumeration_circuit_traits<T>::template simulator<unsigned>
{
public:
  using node_t = typename circuit_traits<T>::node;

public:
  unsigned get_input( const node_t& node, const std::string& name, unsigned pos, const T& circ ) const
  {
    return 0u;
  }

  unsigned get_constant() const
  {
    return 0u;
  }

  unsigned invert( const unsigned& v ) const { return v; }

  unsigned and_op( const node_t& node, const unsigned& v1, const unsigned& v2 ) const
  {
    return std::max( v1, v2 ) + 1u;
  }

  unsigned maj_op( const node_t& node, const unsigned& v1, const unsigned& v2, const unsigned& v3 ) const
  {
    return std::max( v1, std::max( v2, v3 ) ) + 1u;
  }
};

template<typename T>
inline void copy_outputs_from_info( const T& circ, std::vector<typename circuit_traits<T>::node>& nodes )
{
  for ( const auto& o : boost::get_property( circ, boost::graph_name ).outputs )
  {
    nodes += o.first.node;
  }
}

}

/******************************************************************************
 * Public functions                                                           *
 ******************************************************************************/

template<typename T, typename CutsType = structural_cut>
std::map<typename circuit_traits<T>::node, CutsType> structural_cut_enumeration( const T& circ, unsigned k,
                                                                                const properties::ptr& settings = properties::ptr(),
                                                                                const properties::ptr& statistics = properties::ptr() )
{
  using node_t = typename circuit_traits<T>::node;

  /* settings */
  const auto boundary         = get( settings, "boundary",         boost::dynamic_bitset<>( boost::num_vertices( circ ) ) );
        auto start_nodes      = get( settings, "start_nodes",      std::vector<node_t>() );
  const auto include_constant = get( settings, "include_constant", false );

  /* timer */
  properties_timer t( statistics );

  std::map<node_t, CutsType> cut_map;

  /* nodes */
  if ( start_nodes.empty() )
  {
    detail::copy_outputs_from_info( circ, start_nodes );
  }

  /* simulation */
  typename circuit_traits<T>::node_color_map colors;
  detail::structural_cut_enumeration_simulator<T, CutsType> simulator( circ, k, include_constant, boundary );

  for ( const auto& node : start_nodes )
  {
    detail::cut_enumeration_circuit_traits<T>::simulate_node( circ, node, simulator, colors, cut_map );
  }

  return cut_map;
}

template<typename T, typename CutsType = structural_cut>
tt simulate_cut_tt( const typename circuit_traits<T>::node& node,
                    const typename CutsType::value_type& cut,
                    const T& circ,
                    const properties::ptr& settings = properties::ptr(),
                    const properties::ptr& statistics = properties::ptr() )
{
  /* timer */
  properties_timer t( statistics );

  /* simulation */
  std::map<typename circuit_traits<T>::node, tt> assignment;
  auto i = 0u;

  foreach_node_in_cut( cut, [&]( unsigned node ) {
      if ( node == 0u ) { return; }

      assignment.insert( {node, tt_nth_var( i++ )} );
    } );

  typename detail::cut_enumeration_circuit_traits<T>::tt_simulator tt_sim;
  typename detail::cut_enumeration_circuit_traits<T>::template partial_node_assignment_simulator<tt> sim( tt_sim, assignment, tt_const0() );

  return detail::cut_enumeration_circuit_traits<T>::template simulate_node<tt>( circ, node, sim );
}

template<typename T, typename CutsType = structural_cut>
typename CutsType::value_type cut_cone( const typename circuit_traits<T>::node& node, const typename CutsType::value_type& cut, const T& circ )
{
  using cut_t = typename CutsType::value_type;

  std::map<typename circuit_traits<T>::node, cut_t> assignment;
  const auto n = boost::num_vertices( circ );

  foreach_node_in_cut( cut, [&]( unsigned node ) {
      assignment.insert( {node, cut_from_index<cut_t>( node, n )} );
    } );

  detail::structural_cut_cone_simulator<T, CutsType> count_sim( n );
  typename detail::cut_enumeration_circuit_traits<T>::template partial_node_assignment_simulator<cut_t> sim( count_sim, assignment, empty_cut<cut_t>( n ) );

  return detail::cut_enumeration_circuit_traits<T>::template simulate_node<cut_t>( circ, node, sim );
}

template<typename T, typename CutsType = structural_cut>
unsigned cut_cone_size( const typename circuit_traits<T>::node& node, const typename CutsType::value_type& cut, const T& circ )
{
  return cut_size( cut_cone<T, CutsType>( node, cut, circ ) );
}

template<typename T, typename CutsType = structural_cut>
unsigned cut_cone_depth( const typename circuit_traits<T>::node& node, const typename CutsType::value_type& cut, const T& circ )
{
  std::map<typename circuit_traits<T>::node, unsigned> assignment;

  foreach_node_in_cut( cut, [&]( unsigned node ) {
      assignment.insert( {node, 0u} );
    } );

  detail::structural_cut_depth_simulator<T> depth_sim;
  typename detail::cut_enumeration_circuit_traits<T>::template partial_node_assignment_simulator<unsigned> sim( depth_sim, assignment, 0u );

  return detail::cut_enumeration_circuit_traits<T>::template simulate_node<unsigned>( circ, node, sim );
}

// template<typename T, typename CutType = structural_cut>
// std::map<typename circuit_traits<T>::node, unsigned> cut_cone_children_depth( const typename circuit_traits<T>::node& node, const typename CutType::value_type& cut, const T& circ )
// {
//   std::map<typename circuit_traits<T>::node, unsigned> depths;
//   std::stack<std::pair<typename circuit_traits<T>::node, unsigned>> stack;
//   stack.push( {node, 0u} );

//   while( !stack.empty() )
//   {
//     auto top = stack.top(); stack.pop();

//     if ( cut[top.first] )
//     {
//       depths.insert( {top.first, top.second} );
//     }
//     else
//     {
//       for ( const auto& child : boost::make_iterator_range( boost::adjacent_vertices( top.first, circ ) ) )
//       {
//         stack.push( {child, top.second + 1u} );
//       }
//     }
//   }
// }

}

#endif

// Local Variables:
// c-basic-offset: 2
// eval: (c-set-offset 'substatement-open 0)
// eval: (c-set-offset 'innamespace 0)
// End:

max_stars_repo_name: eletesta/cirkit

max_issues_repo_name: eletesta/cirkit

max_forks_repo_name: eletesta/cirkit