diff --git a/netbeans_build.sh b/netbeans_build.sh
new file mode 100755
index 0000000000000000000000000000000000000000..dc7cce984098242a560571bc49394ee021783ba3
--- /dev/null
+++ b/netbeans_build.sh
@@ -0,0 +1,5 @@
+#!/bin/bash
+
+module load gcc-5.3.0 cmake-3.4.3 intel_parallel_studio_ex-2016.1
+make -j 6 tnlSatanMICVectorExperimentalTest-dbg
+#make -j 6 tnl-image-converter-dbg
\ No newline at end of file
diff --git a/netbeans_run_debug_satanexperimental.sh b/netbeans_run_debug_satanexperimental.sh
new file mode 100755
index 0000000000000000000000000000000000000000..14d2858f09cd95c1e17f960aa84f69d50735b977
--- /dev/null
+++ b/netbeans_run_debug_satanexperimental.sh
@@ -0,0 +1,7 @@
+#!/bin/bash
+
+module load gcc-5.3.0 cmake-3.4.3 intel_parallel_studio_ex-2016.1
+cd Debug/bin
+#OFFLOAD_REPORT=2 ./tnlSatanExperimentalTest-dbg
+#./tnlSatanExperimentalTest-dbg
+./tnlSatanMICVectorExperimentalTest-dbg
\ No newline at end of file
diff --git a/src/core/arrays/tnlArrayMIC_impl.h b/src/core/arrays/tnlArrayMIC_impl.h
new file mode 100644
index 0000000000000000000000000000000000000000..7f8e902306e5cc579a2722d59b88e54c82e90506
--- /dev/null
+++ b/src/core/arrays/tnlArrayMIC_impl.h
@@ -0,0 +1,934 @@
+/*
+ * To change this license header, choose License Headers in Project Properties.
+ * To change this template file, choose Tools | Templates
+ * and open the template in the editor.
+ */
+
+/*
+ * File: tnlarrayMIC_impl.h
+ * Author: hanouvit
+ *
+ * Created on 19. dubna 2016, 12:55
+ */
+
+/*
+ * Because MIC Array differs from Host or Cuda Arrays, there is specialization for whole tnlArray
+ *
+ * It differ, becouse Im not able get pointer to MIC memory on processor, for funny pointer work with partially shared arrays.
+ * So Im storing base adress and offset.
+ *
+ * */
+
+#ifndef TNLARRAYMIC_IMPL_H
+#define TNLARRAYMIC_IMPL_H
+
+#ifdef ENABLE_MIC_SATAN
+
+#include <core/tnlMIC.h>
+
+template< typename Element,
+ typename Index >
+struct tnlArrayParams {
+ mutable Index size;
+ mutable Element* allocationPointer;
+ mutable Element* data;
+ mutable int* referenceCounter;
+};
+
+
+template< typename Element,
+ typename Index >
+class tnlArray <Element, tnlMIC, Index> : public virtual tnlObject
+{
+
+ protected:
+
+ void releaseData() const; ///proč je tohle konst, když to měnà objekt, a dost zásandě?
+
+ //!Number of elements in array
+ mutable Index size;
+
+ mutable Element* allocationPointer;
+ mutable Element* data;
+ // mutable Index offset;
+
+
+ //! Pointer to data on MIC -- MUST be SET UP at beginign of PRAGMA traget
+// mutable Element* data;
+
+ /****
+ * Pointer to the originally allocated data. They might differ if one
+ * long array is partitioned into more shorter arrays. Each of them
+ * must know the pointer on allocated data because the last one must
+ * deallocate the array. If outer data (not allocated by TNL) are bind
+ * then this pointer is zero since no deallocation is necessary.
+ */
+ // mutable Element* allocationPointer;
+
+ /****
+ * Counter of objects sharing this array or some parts of it. The reference counter is
+ * allocated after first sharing of the data between more arrays. This is to avoid
+ * unnecessary dynamic memory allocation.
+ */
+ mutable int* referenceCounter;
+
+
+
+ public:
+ typedef Element ElementType;
+ typedef tnlMIC DeviceType;
+ typedef Index IndexType;
+ typedef tnlArray< Element, tnlHost, Index > HostType;
+ typedef tnlArray< Element, tnlCuda, Index > CudaType;
+ typedef tnlArray< Element, tnlMIC, Index > ThisType;
+
+ tnlArray();
+
+ tnlArray( const IndexType& size );
+
+ tnlArray( Element* data,
+ const IndexType& size );
+
+ tnlArray( tnlArray< Element, tnlMIC, Index >& array,
+ const IndexType& begin = 0,
+ const IndexType& size = 0 );
+
+ static tnlString getType();
+
+ tnlString getTypeVirtual() const;
+
+ static tnlString getSerializationType();
+
+ virtual tnlString getSerializationTypeVirtual() const;
+
+ /*Holy fuck HACK*/
+ __device_callable__ tnlArrayParams<Element,Index> getParams(void) const;
+ __device_callable__ void setParams(tnlArrayParams<Element,Index> p);
+
+ /****
+ * This sets size of the array. If the array shares data with other arrays
+ * these data are released. If the current data are not shared and the current
+ * size is the same as the new one, nothing happens.
+ */
+ bool setSize( Index size );
+
+ template< typename Array >
+ bool setLike( const Array& array );
+
+ void bind( Element* _data,
+ const Index _size );
+
+ void bind( const tnlArray< Element, tnlMIC, Index >& array,
+ const IndexType& begin = 0,
+ const IndexType& size = 0 );
+
+ template< int Size >
+ void bind( tnlStaticArray< Size, Element >& array );
+
+ void swap( tnlArray< Element, tnlMIC, Index >& array );
+
+ void reset();
+
+ __device_callable__ Index getSize() const;
+
+ void setElement( const Index& i, const Element& x );
+
+ Element getElement( const Index& i ) const;
+
+ __device_callable__ inline Element& operator[] ( const Index& i );
+
+ __device_callable__ inline const Element& operator[] ( const Index& i ) const;
+
+ tnlArray< Element, tnlMIC, Index >& operator = ( const tnlArray< Element, tnlMIC, Index >& array );
+
+ template< typename Array >
+ tnlArray< Element, tnlMIC, Index >& operator = ( const Array& array );
+
+ template< typename Array >
+ bool operator == ( const Array& array ) const;
+
+ template< typename Array >
+ bool operator != ( const Array& array ) const;
+
+ void setValue( const Element& e ); //nastaveni hodnoty celymu poly
+
+ __device_callable__ const Element* getData() const;
+
+ __device_callable__ Element* getData();
+
+ /*!
+ * Returns true if non-zero size is set.
+ */
+ operator bool() const;
+
+ //! This method measures data transfers done by this vector.
+ /*!
+ * Every time one touches this grid touches * size * sizeof( Real ) bytes are added
+ * to transfered bytes in tnlStatistics.
+ */
+ #ifdef HAVE_NOT_CXX11
+ template< typename IndexType2 >
+ void touch( IndexType2 touches = 1 ) const;
+ #else
+ // template< typename IndexType2 = Index >
+ // void touch( IndexType2 touches = 1 ) const;
+ #endif
+
+ //! Method for saving the object to a file as a binary data.
+ bool save( tnlFile& file ) const;
+
+ //! Method for loading the object from a file as a binary data.
+ bool load( tnlFile& file );
+
+ //! This method loads data without reallocation.
+ /****
+ * This is useful for loading data into shared arrays.
+ * If the array was not initialize yet, common load is
+ * performed. Otherwise, the array size must fit with
+ * the size of array being loaded.
+ */
+ bool boundLoad( tnlFile& file );
+
+ bool boundLoad( const tnlString& fileName );
+
+ using tnlObject::load;
+
+ using tnlObject::save;
+
+ ~tnlArray();
+
+
+};
+
+/*template< typename Element, typename Index >
+ostream& operator <<<Element,tnlMIC,Index> ( ostream& str, const tnlArray< Element, tnlMIC, Index >& v );
+*/
+/////////////////////////////////////////////////////////////////////////////////////////////
+//-----------------------------IMPLEMENTATION------------------------------------------------
+/////////////////////////////////////////////////////////////////////////////////////////////
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >::
+tnlArray()
+: size( 0 ),
+ // offset( 0 ),
+ data(0),
+ allocationPointer( 0 ),
+ referenceCounter( 0 )
+{
+};
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >::
+tnlArray( const IndexType& size )
+: size( 0 ),
+// offset( 0 ),
+ data( 0 ),
+ allocationPointer( 0 ),
+ referenceCounter( 0 )
+{
+ this->setSize( size );
+}
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >::
+tnlArray( Element* data,
+ const IndexType& size )
+{
+ tnlAssert(false,cerr << "This constructor is NOT implemnted yet on MIC" <<endl;);
+}
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >::
+tnlArray( tnlArray< Element, tnlMIC, Index >& array,
+ const IndexType& begin,
+ const IndexType& size )
+: size( size ),
+// offset( begin ),
+ data( 0 ),
+ allocationPointer( array.allocationPointer ),
+ referenceCounter( 0 )
+{
+ tnlAssert( begin < array.getSize(),
+ std::cerr << " begin = " << begin << " array.getSize() = " << array.getSize() );
+ tnlAssert( begin + size < array.getSize(),
+ std::cerr << " begin = " << begin << " size = " << size << " array.getSize() = " << array.getSize() );
+ if( ! this->size )
+ this->size = array.getSize() - begin;
+ if( array.allocationPointer )
+ {
+ if( array.referenceCounter )
+ {
+ this->referenceCounter = array.referenceCounter;
+ *this->referenceCounter++;
+ }
+ else
+ {
+ this->referenceCounter = array.referenceCounter = new int;
+ *this->referenceCounter = 2;
+ }
+ }
+}
+
+
+template< typename Element,typename Index >
+tnlString
+tnlArray< Element, tnlMIC, Index >::
+getType()
+{
+ return tnlString( "tnlArray< " ) +
+ ::getType< Element >() + ", " +
+ DeviceType :: getDeviceType() + ", " +
+ ::getType< Index >() +
+ " >";
+};
+
+template< typename Element,
+ typename Index >
+tnlString
+tnlArray< Element, tnlMIC, Index >::
+getTypeVirtual() const
+{
+ return this->getType();
+};
+
+template< typename Element,
+ typename Index >
+tnlString
+tnlArray< Element, tnlMIC, Index >::
+getSerializationType()
+{
+ return HostType::getType();
+};
+
+template< typename Element,
+ typename Index >
+tnlString
+tnlArray< Element, tnlMIC, Index >::
+getSerializationTypeVirtual() const
+{
+ return this->getSerializationType();
+};
+
+
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+template< typename Element,
+ typename Index >
+__device_callable__
+tnlArrayParams<Element,Index>
+tnlArray< Element, tnlMIC, Index >::
+getParams(void) const
+{
+ tnlArrayParams<Element,Index> p;
+ p.size=this->size;
+ p.allocationPointer=this->allocationPointer;
+ p.data=this->data;
+ p.referenceCounter=this->referenceCounter;
+ return p;
+}
+
+template< typename Element,
+ typename Index >
+__device_callable__
+void
+tnlArray< Element, tnlMIC, Index >::
+setParams(tnlArrayParams<Element,Index> p)
+{
+ this->size=p.size;
+ this->allocationPointer=p.allocationPointer;
+ this->data=p.data;
+ this->referenceCounter=p.referenceCounter;
+
+}
+
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+releaseData() const
+{
+ if( this->referenceCounter )
+ {
+ if( --*this->referenceCounter == 0 )
+ {
+ //tnlArrayOperations< Device >::freeMemory( this->allocationPointer );
+ //free memory on MIC
+ #pragma offload target(mic) //it may work -- look at tnlSatanExperimentalTest.cpp
+ {
+ free(allocationPointer);
+ } ;
+ delete this->referenceCounter;
+ //std::cerr << "Deallocating reference counter " << this->referenceCounter << std::endl;
+ }
+ }
+ else
+ if( allocationPointer )
+ {
+ //free memory on MIC
+ #pragma offload target(mic) //it may work -- look at tnlSatanExperimentalTest.cpp
+ {
+ free(allocationPointer);
+ };
+ }
+ this->allocationPointer = 0;
+ this->size = 0;
+ this->data=0;
+ this->referenceCounter = 0;
+}
+
+
+template< typename Element,
+ typename Index >
+bool
+tnlArray< Element, tnlMIC, Index >::
+setSize( const Index size )
+{
+ tnlAssert( size >= 0,
+ cerr << "You try to set size of tnlArray to negative value."
+ << "New size: " << size << endl );
+ if( this->size == size && allocationPointer && ! referenceCounter ) return true;
+ this->releaseData();
+ //tnlArrayOperations< Device >::allocateMemory( this->allocationPointer, size );
+ //alloc memory on MIC
+ #pragma offload target(mic) //it may work -- look at tnlSatanExperimentalTest.cpp
+ {
+ allocationPointer=(Element*)malloc(size*sizeof(Element));
+ };
+ this->size = size;
+ // this->offset=0;
+ this->data=this->allocationPointer;
+ return true;
+};
+
+template< typename Element,
+ typename Index >
+ template< typename Array >
+bool
+tnlArray< Element, tnlMIC, Index >::
+setLike( const Array& array )
+{
+ tnlAssert( array. getSize() >= 0,
+ cerr << "You try to set size of tnlArray to negative value."
+ << "Array size: " << array. getSize() << endl );
+ return setSize( array.getSize() );
+};
+
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+reset()
+{
+ this->releaseData();
+};
+
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+bind( Element* data,
+ const Index size )
+{
+ this->releaseData();
+ this->data = data;
+ this->size = size;
+}
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+bind( const tnlArray< Element, tnlMIC, Index >& array,
+ const IndexType& begin,
+ const IndexType& size )
+{
+ tnlAssert( begin <= array.getSize(),
+ std::cerr << " begin = " << begin << " array.getSize() = " << array.getSize() );
+ tnlAssert( begin + size <= array.getSize(),
+ std::cerr << " begin = " << begin << " size = " << size << " array.getSize() = " << array.getSize() );
+
+ this->releaseData();
+ if( size )
+ this->size = size;
+ else
+ this->size = array.getSize() - begin;
+ //this->data = const_cast< Element* >( &array.getData()[ begin ] );
+ this->data = const_cast< Element* >( array.getData()+begin );
+ this->allocationPointer = array.allocationPointer;
+ if( array.allocationPointer )
+ {
+ if( array.referenceCounter )
+ {
+ this->referenceCounter = array.referenceCounter;
+ ( *this->referenceCounter )++;
+ }
+ else
+ {
+ this->referenceCounter = array.referenceCounter = new int;
+ *this->referenceCounter = 2;
+ //std::cerr << "Allocating reference counter " << this->referenceCounter << std::endl;
+ }
+ }
+}
+
+template< typename Element,
+ typename Index >
+ template< int Size >
+void
+tnlArray< Element, tnlMIC, Index >::
+bind( tnlStaticArray< Size, Element >& array )
+{
+ /*this->releaseData();
+ this->size = Size;
+ this->data = array.getData();*/
+
+ tnlAssert(false, cerr << "Not implemented jet for MIC :-(" <<endl );
+}
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+swap( tnlArray< Element, tnlMIC, Index >& array )
+{
+ ::swap( this->size, array.size );
+ ::swap( this->data, array.data );
+ ::swap( this->allocationPointer, array.allocationPointer );
+ ::swap( this->referenceCounter, array.referenceCounter );
+};
+
+template< typename Element,
+ typename Index >
+__device_callable__
+Index
+tnlArray< Element, tnlMIC, Index >::
+getSize() const
+{
+ return this->size;
+}
+
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+setValue( const Element& e )
+{
+ Element ee=e;
+ #pragma offload target(mic) in(ee)
+ {
+ for(int i=0;i<this->size;i++)
+ data[i]=ee;
+ }
+}
+
+template< typename Element,
+ typename Index >
+void
+tnlArray< Element, tnlMIC, Index >::
+setElement( const Index& i, const Element& x )
+{
+ Element xx=x;
+ Index ii=i;
+ #pragma offload target(mic) in(xx,ii)
+ {
+ this->data[ii]=xx;
+ }
+}
+
+template< typename Element,
+ typename Index >
+Element
+tnlArray< Element, tnlMIC, Index >::
+getElement( const Index& i ) const
+{
+ Element x;
+ Index ii=i;
+ #pragma offload target(mic) in(ii) out(x)
+ {
+ x=this->data[i];
+ }
+ return x;
+}
+
+template< typename Element,
+ typename Index >
+__device_callable__
+Element*
+tnlArray< Element, tnlMIC, Index > ::
+getData()
+{
+ return this->data;
+}
+
+template< typename Element,
+ typename Index >
+__device_callable__
+const Element*
+tnlArray< Element, tnlMIC, Index > ::
+getData() const
+{
+ return this->data;
+}
+
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index > ::
+operator bool() const
+{
+ return data != 0;
+};
+
+
+template< typename Element,
+ typename Index >
+__device_callable__
+inline Element&
+tnlArray< Element, tnlMIC, Index >::
+operator[] ( const Index& i )
+{
+ tnlAssert( 0 <= i && i < this->getSize(),
+ cerr << "Wrong index for operator[] in tnlArray "
+ << " index is " << i
+ << " and array size is " << this->getSize() );
+ return this->data[ i ];
+};
+
+template< typename Element,
+ typename Index >
+__device_callable__
+inline const Element&
+tnlArray< Element, tnlMIC, Index >::
+operator[] ( const Index& i ) const
+{
+ tnlAssert( 0 <= i && i < this->getSize(),
+ cerr << "Wrong index for operator[] in tnlArray "
+ << " index is " << i
+ << " and array size is " << this->getSize() );
+ return this->data[ i ];
+};
+
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >&
+tnlArray< Element, tnlMIC, Index >::
+operator = ( const tnlArray< Element, tnlMIC, Index >& array )
+{
+ tnlAssert( array. getSize() == this->getSize(),
+ cerr << "Source size: " << array. getSize() << endl
+ << "Target size: " << this->getSize() << endl );
+
+ tnlArrayParams<Element,Index>params= array.getParams();
+ #pragma offload target(mic) in(params)
+ {
+ for(int i=0;i<params.size;i++)
+ {
+ data[i]=params.data[i];
+ }
+ }
+ return ( *this );
+};
+
+template< typename Element,
+ typename Index >
+ template< typename Array >
+tnlArray< Element, tnlMIC, Index >&
+tnlArray< Element, tnlMIC, Index >::
+operator = ( const Array& array )
+{
+
+ tnlAssert( array. getSize() == this->getSize(),
+ cerr << "Source size: " << array. getSize() << endl
+ << "Target size: " << this->getSize() << endl );
+
+
+
+ if( std::is_same< typename Array::DeviceType, tnlMIC >::value )
+ {
+ //MIC -> MIC
+ cout << "MIC to MIC transfer.." <<endl;
+ cout << "ZajĂmalo by mne zda se tento kĂłd nÄ›kdy spustĂ... SNAD NE, protoĹľe nic nedÄ›lá " << endl;
+ cout << "S implementaci je trochu háčk" <<endl;
+
+ }
+ else
+ {
+ if(std::is_same< typename Array::DeviceType, tnlHost >::value)
+ {
+ //Host -> MIC
+ //cout << "Host to MIC" <<endl;
+ const typename Array::ElementType * sdata = array.getData();
+ #pragma offload target(mic) in(sdata:length(this->size))
+ {
+ for(int i=0;i<this->size;i++) //MIC má memcpy, ale tohle je KISS
+ {
+ data[i]=sdata[i];
+ }
+ }
+ }
+ else
+ {
+ //?? ->MIC
+ cout << "Extremly slow transfer at = of tnlArray to MIC count:" << this->size <<endl;
+ for(int i=0;i<this->size;i++)
+ {
+ this->setElement(i,array.getElement(i));
+ }
+ }
+ }
+
+ return ( *this );
+};
+
+//Obdoba TNL ARRAY OPERATIONS... možná by se to tam mělo pak přepsat...
+template <typename Device>
+class satanArrayCompare{};
+
+template <>
+class satanArrayCompare< tnlHost >{
+
+ public:
+ //expected ArrayA is tnlMIC
+ template< typename ArrayA, typename ArrayB>
+ static bool compare( ArrayA& A, ArrayB& B )
+ {
+ if(A.getSize()!=B.getSize())
+ {
+ return false;
+ }
+ else
+ {
+ tnlArrayParams<typename ArrayA::ElementType,typename ArrayA::IndexType> A_params=A.getParams();
+ const typename ArrayB::ElementType * B_data=B.getData();
+ typename ArrayB::IndexType B_size=B.getSize();
+ bool result=true;
+ typename ArrayA::ThisType AA;
+ #pragma offload target(mic) nocopy(AA) in(A_params) in(B_data:length(B_size)) inout(result)
+ {
+
+ AA.setParams(A_params);
+ for(int i=0;i<B_size;i++)
+ {
+ if(AA[i]!=B_data[i])
+ {
+ result=false;
+ break;
+ }
+ }
+ }
+ return result;
+ }
+ };
+};
+
+template <>
+class satanArrayCompare< tnlMIC >{
+
+ public:
+ //expected ArrayA is on MIC
+ template< typename ArrayA, typename ArrayB>
+ static bool compare( ArrayA& A, ArrayB& B )
+ {
+ if(A.getSize()!=B.getSize())
+ {
+ return false;
+ }
+ else
+ {
+ tnlArrayParams<typename ArrayA::ElementType,typename ArrayA::IndexType> A_params=A.getParams();
+ tnlArrayParams<typename ArrayB::ElementType,typename ArrayB::IndexType> B_params=B.getParams();
+ typename ArrayA::ThisType AA;
+ typename ArrayB::ThisType BB;
+ bool result=true;
+ #pragma offload target(mic) nocopy(AA,BB) in(A_params,B_params) inout(result)
+ {
+ AA.setParams(A_params);
+ BB.setParams(B_params);
+ for(int i=0;i<AA.getSize();i++)
+ {
+ if(AA[i]!=BB[i])
+ {
+ result=false;
+ break;
+ }
+ }
+ }
+ return result;
+ }
+ };
+};
+
+
+template< typename Element,
+ typename Index >
+ template< typename Array >
+bool
+tnlArray< Element, tnlMIC, Index >::
+operator == ( const Array& array ) const
+{
+ //cout << "Comparing: " << this->getType() << ": " <<this->getSize() << " and " << array.getType() <<": "<< array.getSize() <<endl;
+ return satanArrayCompare< typename Array :: DeviceType> :: compare((*this),array);
+
+}
+
+
+template< typename Element,
+ typename Index >
+ template< typename Array >
+bool tnlArray< Element, tnlMIC, Index > ::
+operator != ( const Array& array ) const
+{
+ return ! ( ( *this ) == array );
+}
+
+
+template< typename Element,
+ typename Index >
+tnlArray< Element, tnlMIC, Index >::
+~tnlArray()
+{
+ this->releaseData();
+}
+
+template< typename Element,
+ typename Index >
+bool tnlArray< Element, tnlMIC, Index > ::
+save( tnlFile& file ) const
+{
+ if( ! tnlObject :: save( file ) )
+ return false;
+#ifdef HAVE_NOT_CXX11
+ if( ! file. write< const Index, tnlHost >( &this->size ) )
+ return false;
+#else
+ if( ! file. write( &this->size ) )
+ return false;
+#endif
+ if( this->size != 0 && ! tnlArrayIO< Element, tnlMIC, Index >::save( file, this->data, this->size ) )
+ {
+ cerr << "I was not able to save " << this->getType()
+ << " with size " << this->getSize() << endl;
+ return false;
+ }
+ return true;
+};
+
+template< typename Element,
+ typename Index >
+bool
+tnlArray< Element, tnlMIC, Index >::
+load( tnlFile& file )
+{
+ if( ! tnlObject :: load( file ) )
+ return false;
+ Index _size;
+#ifdef HAVE_NOT_CXX11
+ if( ! file. read< Index, tnlHost >( &_size ) )
+ return false;
+#else
+ if( ! file. read( &_size ) )
+ return false;
+#endif
+ if( _size < 0 )
+ {
+ cerr << "Error: The size " << _size << " of the file is not a positive number or zero." << endl;
+ return false;
+ }
+ setSize( _size );
+ if( _size )
+ {
+ if( ! tnlArrayIO< Element, tnlMIC, Index >::load( file, this->data, this->size ) )
+ {
+ cerr << "I was not able to load " << this->getType()
+ << " with size " << this->getSize() << endl;
+ return false;
+ }
+ }
+ return true;
+}
+
+template< typename Element,
+ typename Index >
+bool
+tnlArray< Element, tnlMIC, Index >::
+boundLoad( tnlFile& file )
+{
+ if( ! tnlObject :: load( file ) )
+ return false;
+ Index _size;
+#ifdef HAVE_NOT_CXX11
+ if( ! file. read< Index, tnlHost >( &_size ) )
+ return false;
+#else
+ if( ! file. read( &_size ) )
+ return false;
+#endif
+ if( _size < 0 )
+ {
+ cerr << "Error: The size " << _size << " of the file is not a positive number or zero." << endl;
+ return false;
+ }
+ if( this->getSize() != 0 )
+ {
+ if( this->getSize() != _size )
+ {
+ std::cerr << "Error: The current array size is not zero and it is different from the size of" << std::endl
+ << "the array being loaded. This is not possible. Call method reset() before." << std::endl;
+ return false;
+ }
+ }
+ else setSize( _size );
+ if( _size )
+ {
+ if( ! tnlArrayIO< Element, tnlMIC, Index >::load( file, this->data, this->size ) )
+ {
+ cerr << "I was not able to load " << this->getType()
+ << " with size " << this->getSize() << endl;
+ return false;
+ }
+ }
+ return true;
+}
+
+template< typename Element,
+ typename Index >
+bool
+tnlArray< Element, tnlMIC, Index >::
+boundLoad( const tnlString& fileName )
+{
+ tnlFile file;
+ if( ! file. open( fileName, tnlReadMode ) )
+ {
+ cerr << "I am not bale to open the file " << fileName << " for reading." << endl;
+ return false;
+ }
+ if( ! this->boundLoad( file ) )
+ return false;
+ if( ! file. close() )
+ {
+ cerr << "An error occurred when I was closing the file " << fileName << "." << endl;
+ return false;
+ }
+ return true;
+}
+
+#endif
+
+#endif /* TNLARRAYMIC_IMPL_H */
+
+
diff --git a/src/core/tnlDevice_Callable.h b/src/core/tnlDevice_Callable.h
new file mode 100644
index 0000000000000000000000000000000000000000..a95454d9508931ee34720c3c9c17ddd8f68fdf6d
--- /dev/null
+++ b/src/core/tnlDevice_Callable.h
@@ -0,0 +1,34 @@
+/*
+ * To change this license header, choose License Headers in Project Properties.
+ * To change this template file, choose Tools | Templates
+ * and open the template in the editor.
+ */
+
+/*
+ * File: tnldevice_callable.h
+ * Author: hanouvit
+ *
+ * Created on 18. dubna 2016, 15:49
+ */
+
+#ifndef TNLDEVICE_CALLABLE_H
+#define TNLDEVICE_CALLABLE_H
+
+//deprecated __cuda_callable__
+#ifdef HAVE_CUDA
+ #define __cuda_callable__ __device__ __host__
+#else
+ #define __cuda_callable__
+#endif
+
+//NEW, better __device_callable__ --used only with MIC touch code
+#ifdef HAVE_ICPC
+ #define __device_callable__ __attribute__((target(mic)))
+#elif HAVE_CUDA
+ #define __device_callable__ __device__ __host__
+#else
+ #define __device_callable__
+#endif
+
+#endif /* TNLDEVICE_CALLABLE_H */
+
diff --git a/src/core/tnlMIC.h b/src/core/tnlMIC.h
new file mode 100644
index 0000000000000000000000000000000000000000..0906728f1cadce4ed748cdec2c12d831bc6747c8
--- /dev/null
+++ b/src/core/tnlMIC.h
@@ -0,0 +1,56 @@
+/*
+ * To change this license header, choose License Headers in Project Properties.
+ * To change this template file, choose Tools | Templates
+ * and open the template in the editor.
+ */
+
+/*
+ * File: tnlMIC.h
+ * Author: hanouvit
+ *
+ * Created on 18. dubna 2016, 12:38
+ */
+
+#include <iostream>
+#include <core/tnlString.h>
+#include <core/tnlDevice.h>
+#include <core/tnlDevice_Callable.h>
+
+
+#ifndef TNLMIC_H
+#define TNLMIC_H
+
+#define ALLOC alloc_if(1) //naalokuj promenou na zacatku offload bloku -- default
+#define FREE free_if(1) // smaz promenou na konci offload bloku -- default
+#define RETAIN free_if(0) //nesmaz promenou na konci bloku
+#define REUSE alloc_if(0) //nealokuj proměnnou na zacatku
+
+
+class tnlMIC
+{
+ public:
+ //useful debuging -- but produce warning
+ __device_callable__ static inline void CheckMIC(void)
+ {
+ #ifdef __MIC__
+ std::cout<<"ON MIC"<<std::endl;
+ #else
+ std::cout<<"ON CPU" <<std::endl;
+ #endif
+ }
+
+
+ static tnlString getDeviceType()
+ {
+ return tnlString( "tnlMIC" );
+ }
+
+ __device_callable__ static inline tnlDeviceEnum getDevice()
+ {
+ return tnlMICDevice;
+ }
+};
+
+
+#endif /* TNLMIC_H */
+
diff --git a/src/core/vectors/tnlVectorOperationsMIC_impl.h b/src/core/vectors/tnlVectorOperationsMIC_impl.h
new file mode 100644
index 0000000000000000000000000000000000000000..d138a929b27a480f522f6486ed249b8825e6b125
--- /dev/null
+++ b/src/core/vectors/tnlVectorOperationsMIC_impl.h
@@ -0,0 +1,512 @@
+/*
+ * To change this license header, choose License Headers in Project Properties.
+ * To change this template file, choose Tools | Templates
+ * and open the template in the editor.
+ */
+
+/*
+ * File: tnlVectorOperationsMIC_impl.h
+ * Author: hanouvit
+ *
+ * Created on 2. května 2016, 12:57
+ */
+
+#ifndef TNLVECTOROPERATIONSMIC_IMPL_H
+#define TNLVECTOROPERATIONSMIC_IMPL_H
+
+//static const int OpenMPVectorOperationsThreshold = 65536; // TODO: check this threshold
+
+template< typename Vector >
+void tnlVectorOperations< tnlMIC >::addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value )
+{
+ // v[ i ] += value;
+ //cout << "Errorous function, not clear wher should be called (device or Host)" << endl;
+ v.setElement(i,v.getElemet(i)+value);
+}
+
+template< typename Vector >
+void tnlVectorOperations< tnlMIC >::addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value,
+ const typename Vector::RealType& thisElementMultiplicator )
+{
+ //v[ i ] = thisElementMultiplicator * v[ i ] + value;
+ // cout << "Errorous function, not clear wher should be called (device or Host)" << endl;
+ v.setElement(i,thisElementMultiplicator*v.getElemet(i)+value);
+}
+
+template< typename Vector >
+typename Vector::RealType tnlVectorOperations< tnlMIC >::getVectorMax( const Vector& v )
+{
+ //tady je moĹľnost paralelizace
+ typename Vector :: RealType result;
+ #pragma offload target(mic) out(result)
+ {
+ result=v[0];
+ for(typename Vector :: Index i=1;i<v.getSize();i++)
+ {
+ if(result<v[i])
+ result=v[i];
+ }
+ }
+ return result;
+}
+
+template< typename Vector >
+typename Vector :: RealType tnlVectorOperations< tnlMIC > :: getVectorMin( const Vector& v )
+{
+ //tady je moĹľnost paralelizace
+ typename Vector :: RealType result;
+ #pragma offload target(mic) out(result)
+ {
+ result=v[0];
+ for(typename Vector :: Index i=1;i<v.getSize();i++)
+ {
+ if(result>v[i])
+ result=v[i];
+ }
+ }
+ return result;
+}
+
+template< typename Vector >
+typename Vector :: RealType tnlVectorOperations< tnlMIC > :: getVectorAbsMax( const Vector& v )
+{
+ //tady je moĹľnost paralelizace
+ typename Vector :: RealType result;
+ /* #pragma offload target(mic) out(result)
+ {
+ result=fabs(v[0]);
+ for(typename Vector :: Index i=1;i<v.getSize();i++)
+ {
+ if(result<fabs(v[i]))
+ result=fabs(v[i]);
+ }
+ }*/
+ return result;
+}
+
+
+template< typename Vector >
+typename Vector :: RealType tnlVectorOperations< tnlMIC > :: getVectorAbsMin( const Vector& v )
+{
+ /*typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+ tnlAssert( v. getSize() > 0, );
+ Real result = fabs( v. getElement( 0 ) );
+ const Index n = v. getSize();
+ for( Index i = 1; i < n; i ++ )
+ result = Min( result, ( Real ) fabs( v. getElement( i ) ) );
+ return result;*/
+ return 0;
+}
+
+template< typename Vector >
+typename Vector::RealType
+tnlVectorOperations< tnlMIC >::getVectorL1Norm( const Vector& v )
+{
+/* typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+ tnlAssert( v. getSize() > 0, );
+
+ Real result( 0.0 );
+ const Index n = v. getSize();
+#ifdef HAVE_OPENMP
+#pragma omp parallel for reduction(+:result) if( tnlHost::isOMPEnabled() && n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
+#endif
+ for( Index i = 0; i < n; i ++ )
+ result += fabs( v[ i ] );
+ return result;*/
+ return 0;
+}
+
+template< typename Vector >
+typename Vector::RealType
+tnlVectorOperations< tnlMIC >::getVectorL2Norm( const Vector& v )
+{
+ /* typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+ tnlAssert( v. getSize() > 0, );
+ Real result( 0.0 );
+ const Index n = v. getSize();
+#ifdef HAVE_OPENMP
+#pragma omp parallel for reduction(+:result) if( tnlHost::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
+#endif
+ for( Index i = 0; i < n; i ++ )
+ {
+ const Real& aux = v[ i ];
+ result += aux * aux;
+ }
+ return sqrt( result );*/
+ return 0;
+}
+
+template< typename Vector >
+typename Vector::RealType
+tnlVectorOperations< tnlMIC >:: getVectorLpNorm( const Vector& v,
+ const typename Vector :: RealType& p )
+{
+ /*typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+ tnlAssert( v. getSize() > 0, );
+ tnlAssert( p > 0.0,
+ cerr << " p = " << p );
+ if( p == 1.0 )
+ return getVectorL1Norm( v );
+ if( p == 2.0 )
+ return getVectorL2Norm( v );
+
+ Real result( 0.0 );
+ const Index n = v. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ result += pow( fabs( v[ i ] ), p );
+ return pow( result, 1.0 / p );*/
+ return 0;
+}
+
+template< typename Vector >
+typename Vector :: RealType tnlVectorOperations< tnlMIC > :: getVectorSum( const Vector& v )
+{
+ /*
+ typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+ tnlAssert( v. getSize() > 0, );
+
+ Real result( 0.0 );
+ const Index n = v. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ result += v[ i ];
+ return result;*/
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1 :: RealType tnlVectorOperations< tnlMIC> :: getVectorDifferenceMax( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+ Real result = v1. getElement( 0 ) - v2. getElement( 0 );
+ const Index n = v1. getSize();
+ for( Index i = 1; i < n; i ++ )
+ result = Max( result, v1. getElement( i ) - v2. getElement( i ) );
+ return result;*/
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1 :: RealType tnlVectorOperations< tnlMIC > :: getVectorDifferenceMin( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result = v1. getElement( 0 ) - v2. getElement( 0 );
+ const Index n = v1. getSize();
+ for( Index i = 1; i < n; i ++ )
+ result = Min( result, v1. getElement( i ) - v2. getElement( i ) );
+ return result;
+ */
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1 :: RealType tnlVectorOperations< tnlMIC > :: getVectorDifferenceAbsMax( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result = fabs( v1. getElement( 0 ) - v2. getElement( 0 ) );
+ const Index n = v1. getSize();
+ for( Index i = 1; i < n; i ++ )
+ result = Max( result, ( Real ) fabs( v1. getElement( i ) - v2. getElement( i ) ) );
+ return result;
+ */
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1 :: RealType tnlVectorOperations< tnlMIC > :: getVectorDifferenceAbsMin( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result = fabs( v1[ 0 ] - v2[ 0 ] );
+ const Index n = v1. getSize();
+ for( Index i = 1; i < n; i ++ )
+ result = Min( result, ( Real ) fabs( v1[ i ] - v2[ i ] ) );
+ return result;
+ */
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1::RealType
+tnlVectorOperations< tnlMIC >::
+getVectorDifferenceL1Norm( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result( 0.0 );
+ const Index n = v1. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ result += fabs( v1[ i ] - v2[ i ] );
+ return result;
+ */
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1::RealType
+tnlVectorOperations< tnlMIC >::
+getVectorDifferenceL2Norm( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result( 0.0 );
+ const Index n = v1. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ {
+ Real aux = fabs( v1[ i ] - v2[ i ] );
+ result += aux * aux;
+ }
+ return sqrt( result );
+ */
+ return 0;
+}
+
+
+template< typename Vector1, typename Vector2 >
+typename Vector1::RealType
+tnlVectorOperations< tnlMIC >::
+getVectorDifferenceLpNorm( const Vector1& v1,
+ const Vector2& v2,
+ const typename Vector1::RealType& p )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( p > 0.0,
+ cerr << " p = " << p );
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ if( p == 1.0 )
+ return getVectorDifferenceL1Norm( v1, v2 );
+ if( p == 2.0 )
+ return getVectorDifferenceL2Norm( v1, v2 );
+
+ Real result( 0.0 );
+ const Index n = v1. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ result += pow( fabs( v1. getElement( i ) - v2. getElement( i ) ), p );
+ return pow( result, 1.0 / p );
+ */
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+typename Vector1::RealType tnlVectorOperations< tnlMIC > :: getVectorDifferenceSum( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result( 0.0 );
+ const Index n = v1. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ result += v1. getElement( i ) - v2. getElement( i );
+ return result;
+ */
+ return 0;
+}
+
+
+template< typename Vector >
+void tnlVectorOperations< tnlMIC > :: vectorScalarMultiplication( Vector& v,
+ const typename Vector :: RealType& alpha )
+{
+ /*
+ typedef typename Vector :: RealType Real;
+ typedef typename Vector :: IndexType Index;
+
+ tnlAssert( v. getSize() > 0, );
+
+ const Index n = v. getSize();
+
+ for( Index i = 0; i < n; i ++ )
+ v[ i ] *= alpha;
+ */
+
+}
+
+
+template< typename Vector1, typename Vector2 >
+typename Vector1 :: RealType tnlVectorOperations< tnlMIC > :: getScalarProduct( const Vector1& v1,
+ const Vector2& v2 )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v1. getSize() > 0, );
+ tnlAssert( v1. getSize() == v2. getSize(), );
+
+ Real result( 0.0 );
+ const Index n = v1. getSize();
+
+ for( Index i = 0; i < n; i++ )
+ result += v1[ i ] * v2[ i ];
+ /*Real result1( 0.0 ), result2( 0.0 ), result3( 0.0 ), result4( 0.0 ),
+ result5( 0.0 ), result6( 0.0 ), result7( 0.0 ), result8( 0.0 );
+ Index i( 0 );
+ while( i + 8 < n )
+ {
+ result1 += v1[ i ] * v2[ i ];
+ result2 += v1[ i + 1 ] * v2[ i + 1 ];
+ result3 += v1[ i + 2 ] * v2[ i + 2 ];
+ result4 += v1[ i + 3 ] * v2[ i + 3 ];
+ result5 += v1[ i + 4 ] * v2[ i + 4 ];
+ result6 += v1[ i + 5 ] * v2[ i + 5 ];
+ result7 += v1[ i + 6 ] * v2[ i + 6 ];
+ result8 += v1[ i + 7 ] * v2[ i + 7 ];
+ i += 8;
+ }
+ Real result = result1 + result2 + result3 + result4 + result5 +result6 +result7 +result8;
+ while( i < n )
+ result += v1[ i ] * v2[ i++ ];*/
+ /*return result;*/
+ return 0;
+}
+
+template< typename Vector1, typename Vector2 >
+void tnlVectorOperations< tnlMIC > :: addVector( Vector1& y,
+ const Vector2& x,
+ const typename Vector2::RealType& alpha,
+ const typename Vector1::RealType& thisMultiplicator )
+{
+ /*typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( x. getSize() > 0, );
+ tnlAssert( x. getSize() == y. getSize(), );
+
+ const Index n = y. getSize();
+ if( thisMultiplicator == 1.0 )
+
+ for( Index i = 0; i < n; i ++ )
+ y[ i ] += alpha * x[ i ];
+ else
+
+ for( Index i = 0; i < n; i ++ )
+ y[ i ] = thisMultiplicator * y[ i ] + alpha * x[ i ];
+ */
+ return 0;
+}
+
+template< typename Vector1,
+ typename Vector2,
+ typename Vector3 >
+void
+tnlVectorOperations< tnlMIC >::
+addVectors( Vector1& v,
+ const Vector2& v1,
+ const typename Vector2::RealType& multiplicator1,
+ const Vector3& v2,
+ const typename Vector3::RealType& multiplicator2,
+ const typename Vector1::RealType& thisMultiplicator )
+{
+ /*
+ typedef typename Vector1 :: RealType Real;
+ typedef typename Vector1 :: IndexType Index;
+
+ tnlAssert( v.getSize() > 0, );
+ tnlAssert( v.getSize() == v1.getSize(), );
+ tnlAssert( v.getSize() == v2.getSize(), );
+
+ const Index n = v.getSize();
+ if( thisMultiplicator == 1.0 )
+
+ for( Index i = 0; i < n; i ++ )
+ v[ i ] += multiplicator1 * v1[ i ] + multiplicator2 * v2[ i ];
+ else
+
+ for( Index i = 0; i < n; i ++ )
+ v[ i ] = thisMultiplicator * v[ i ] * multiplicator1 * v1[ i ] + multiplicator2 * v2[ i ];
+ */
+}
+
+template< typename Vector >
+void tnlVectorOperations< tnlMIC >::computePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
+{
+ /*
+ typedef typename Vector::IndexType Index;
+ for( Index i = begin + 1; i < end; i++ )
+ v[ i ] += v[ i - 1 ];
+ */
+}
+
+template< typename Vector >
+void tnlVectorOperations< tnlMIC >::computeExclusivePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
+{
+ /*
+ typedef typename Vector::IndexType Index;
+ typedef typename Vector::RealType Real;
+ Real aux( v[ begin ] );
+ v[ begin ] = 0.0;
+ for( Index i = begin + 1; i < end; i++ )
+ {
+ Real x = v[ i ];
+ v[ i ] = aux;
+ aux += x;
+ }
+ */
+}
+
+#endif /* TNLVECTOROPERATIONSMIC_IMPL_H */
+
diff --git a/tests/unit-tests/core/arrays/tnlSatanExperimentalTest.cpp b/tests/unit-tests/core/arrays/tnlSatanExperimentalTest.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..a8eb530c8ea5c37b632fd43897ec929205d99583
--- /dev/null
+++ b/tests/unit-tests/core/arrays/tnlSatanExperimentalTest.cpp
@@ -0,0 +1,241 @@
+/**
+ * ExperimentalnĂ test pro zĂskánĂ zkušenostĂ s TNL a (a strarĂ˝m dobrĂ˝m MIC, intel offloadingem a podobnÄ›)
+ *
+ * 1) Přidat tento Hello world Test do projektu, přeložit a spustit
+ * 2) Write SATANTEST makro for DUMMY testing
+ * 3) Hello from MIC (number of threads)
+ *
+ * 4) Transfer DATA to MIC from MIC.... mno...
+ * 5) Persistant alloc ...... řešà to nocopy, veškeré alokace musà být řešeny "takto"
+ * 6) shared array... no snadne reseni nevidim, omocime to
+ * 7) MIC_callable
+ *
+ * 8) Implement Basic MIC in TNL
+ * 9) tnlMIC.h -- very basic things
+ * 10) tnlArrayMIC_impl.h -->
+ * -->a) search for how to hack Pointer translation (hacked by AUTOMAGIC acces to class memberers)
+ * -->b) search for how to hack moving of non-bitwise copyable object to FUCKING device
+ *
+ * -->c) waiting for other problems...
+ * (It may be usefull to have CUDA version for experiments, how TNL works on coprocessor)
+ * 11) tnlArrayMIC_impl.h --> support most functions, expect saving and loading.
+ * 12) tnlArrayMIC_impl.h --> probably complete.
+ *
+ */
+
+
+#include <iostream>
+#include <core/arrays/tnlArray.h>
+#include <omp.h>
+
+#include <core/tnlMIC.h>
+
+//TUNE MACROS FOR YOUR FUNKY OUTPUT
+#define SATANVERBOSE
+
+#define SATANTEST(a) if((a)){cout << __LINE__ <<":\t OK" <<endl;}else{cout << __LINE__<<":\t FAIL" <<endl;}
+inline void SatanSay( const char * message)
+{
+#ifdef SATANVERBOSE
+ cout << message <<endl;
+#endif
+}
+
+using namespace std;
+
+/*
+ * ICPC offload překládá out of box, pro vypnutà přidat parametr -offload=none
+ *
+ * Variables on MIC are maped to variables on CPU (hash table: hash is base adress on CPU, then contain MIC Address and length on MIC)
+ * https://software.intel.com/en-us/blogs/2013/03/27/behind-the-scenes-offload-memory-management-on-the-intel-xeon-phi-coprocessor
+ *
+ * KEY table for transfers:
+ * https://software.intel.com/en-us/articles/effective-use-of-the-intel-compilers-offload-features
+ *
+ *
+ * It seems that class members can be accesed in offload regionst in class methods without any in/out/inout.
+ * It seems to be copied AUTOMAGICLY and without any pointer translations. (no verification in literature)
+ * -->This allow me to stole pointer from MIC, but doesnt allow me to pass it as parametr of other function for offload
+ * -->Only class members can use this pointer directly, I hope.
+ *
+ * Fucking offload doesnt allow me to copy object with constructor to MIC. But I can alloc it With NOCOPY clauslue,
+ * and upload simple sturcture containing all data and set up object on MIC (HOLY FUCK HACK)
+ * I.E.:
+ * param p =a.getparam();
+ * #pragma offload target(mic) nocopy(a) in(p)
+ * {
+ * a.setparam(p);
+ * ...
+ * }
+ *-->When MIC change object a, then use INOUT and getparam() setparam at the end of OFFLOAD
+ *
+ * Other fucking hacks were written: 1) template <typename Device> class satanArrayCompare {}; in tnlArrayMIC_impl.h
+ * -->very similar to tnlArrayOperations, but expect first array to be on MIC and only for compare
+ * 2) template <typename Type> struct satanHider{ Type * pointer} in tnlFile_impl.h
+ * -->use for passing untranslated pointer into MIC.
+ *
+ */
+
+int main(void)
+{
+ cout << "tnlArray on MIC test by Satan:" <<endl; //LOL
+
+ #ifdef HAVE_ICPC
+ cout << "ICPC in USE" <<endl; //LOL
+ #endif
+
+//prepare arrays with data
+ tnlArray<double,tnlMIC,int> aa(10);
+ tnlArray<double,tnlMIC,int> ee(6);
+ tnlArray<double,tnlHost,int> cc(5);
+
+//fill it UP
+tnlArrayParams<double,int> aa_params=aa.getParams();
+#pragma offload target(mic) nocopy(aa) in(aa_params)
+{
+ aa.setParams(aa_params);
+ for(int i=0;i<aa.getSize();i++)
+ {
+ aa[i]=i;
+ }
+}
+
+for(int i=0;i<5;i++)
+{
+ cc[i]=10+i;
+}
+
+tnlFile soubor;
+soubor.open("/home/hanousek/tnlArrayExperimentSave.bin",tnlWriteMode);
+cc.save(soubor);
+soubor.close();
+
+ee.bind(aa,5,5);
+ee.boundLoad("/home/hanousek/tnlArrayExperimentSave.bin");
+
+
+/*
+tnlFile soubor;
+soubor.open("/home/hanousek/tnlArrayExperimentSave.bin",tnlWriteMode);
+cc.save(soubor);
+soubor.close();
+
+cout << aa.getSize() << endl;
+for(int i=0;i<aa.getSize();i++)
+{
+ cout << aa.getElement(i)<<endl;
+}
+cout <<endl;*/
+/*
+tnlFile soubor;
+soubor.open("/home/hanousek/tnlArrayExperimentSave.bin",tnlWriteMode);
+aa.save(soubor);
+soubor.close();*/
+
+cout << aa.getSize() << endl;
+for(int i=0;i<aa.getSize();i++)
+{
+ cout << aa.getElement(i)<<endl;
+}
+cout <<endl;
+
+cout << ee.getSize() << endl;
+for(int i=0;i<ee.getSize();i++)
+{
+ cout << ee.getElement(i)<<endl;
+}
+cout <<endl;
+
+/*
+soubor.open("/home/hanousek/tnlArrayExperimentSave.bin",tnlReadMode);
+aa.load(soubor);
+soubor.close();
+
+cout << aa.getSize() << endl;
+for(int i=0;i<aa.getSize();i++)
+{
+ cout << aa.getElement(i)<<endl;
+}
+*/
+
+soubor.open("/home/hanousek/tnlArrayExperimentSave.bin",tnlReadMode);
+cc.load(soubor);
+soubor.close();
+
+cout << cc.getSize() << endl;
+for(int i=0;i<cc.getSize();i++)
+{
+ cout << cc.getElement(i)<<endl;
+}
+
+/*
+//prepare arrays for funky tests
+tnlArray<double,tnlMIC,int> bb(10);
+tnlArray<double,tnlMIC,int> dd(0);
+
+//TEST IT!
+SatanSay("aa.getSize():");
+SATANTEST(aa.getSize()==10);
+
+SatanSay("Is aa filled correctly? (aa.getElement):");
+for(int i=0;i<10;i++)
+ SATANTEST(aa.getElement(i)==i);
+
+SatanSay("Copy to bb(MIC->MIC) (=):");
+bb=aa;
+SATANTEST(aa.getSize()==bb.getSize());
+for(int i=0;i<bb.getSize();i++)
+ SATANTEST(bb.getElement(i)==i);
+
+SatanSay("setLike:");
+bb.setLike(cc);
+SATANTEST(bb.getSize()==cc.getSize());
+SatanSay("Copy (Host -> MIC) (=)");
+bb=cc;
+for(int i=0;i<bb.getSize();i++)
+ SATANTEST(bb.getElement(i)==i+10);
+
+SatanSay("setValue:");
+bb.setValue(5);
+for(int i=0;i<bb.getSize();i++)
+ SATANTEST(bb.getElement(i)==5);
+
+SatanSay("swap:");
+aa.swap(bb);
+SATANTEST(aa.getSize()==5||bb.getSize()==10);
+for(int i=0;i<aa.getSize();i++)
+{
+ SATANTEST(aa.getElement(i)==5);
+}
+for(int i=0;i<bb.getSize();i++)
+{
+ SATANTEST(bb.getElement(i)==i);
+}
+
+SatanSay("(MIC -> MIC) ==");
+aa.setLike(bb);
+aa=bb;
+SATANTEST(aa==bb);
+SATANTEST(!(aa!=bb));
+SATANTEST(aa!=ee);
+bb.setElement(5,66);
+SATANTEST(aa!=bb);
+
+SatanSay("(Host -> MIC) !=");
+aa.setLike(cc);
+aa=cc;
+SATANTEST(aa==cc);
+aa.setElement(3,66);
+SATANTEST(aa!=cc);
+
+SatanSay("bidn (light test)");
+dd.bind(bb,5);
+SATANTEST(dd.getSize()==5);
+SATANTEST(dd.getElement(1)==6);
+
+//MylsĂm, Ĺľe nenĂ zdaleka testováno vše...
+
+SATANTEST(false);
+*/
+ return 0;
+}
diff --git a/tests/unit-tests/core/vectors/tnlSatanMICVectorExperimentalTest.cpp b/tests/unit-tests/core/vectors/tnlSatanMICVectorExperimentalTest.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..02b18801dc9c018e3c38fcd5ea9d763377f54b1f
--- /dev/null
+++ b/tests/unit-tests/core/vectors/tnlSatanMICVectorExperimentalTest.cpp
@@ -0,0 +1,68 @@
+/**
+ * Test develop by Satan for implementation of tnlVector on MIC
+ *
+ * 1) Edit tnlVectorOperations.h with specialization for MIC (same prototypes of methods )
+ * 2) Create tnlVectorOperationsMIC_impl.h (dummy, do nothing or retrun 0)
+ * 3) Create test for this (this file)
+ *
+ * 4) Lots of Troubles occured.... some problems in file read/write int tnlFile_impl.h (see it for details, soved by memcpy on MIC)
+ * there is no png.h for MIC (okey, but wo dont need it I thing)
+ *
+ */
+
+
+#include <iostream>
+#include <omp.h>
+
+#include <core/tnlMIC.h>
+//#include <core/tnlFile.h>
+#include <core/arrays/tnlArray.h>
+
+#include <png.h>
+
+
+
+//TUNE MACROS FOR YOUR FUNKY OUTPUT
+#define SATANVERBOSE
+
+#define SATANTEST(a) if((a)){cout << __LINE__ <<":\t OK" <<endl;}else{cout << __LINE__<<":\t FAIL" <<endl;}
+inline void SatanSay( const char * message)
+{
+#ifdef SATANVERBOSE
+ cout << message <<endl;
+#endif
+}
+
+using namespace std;
+
+
+
+int main(void)
+{
+ cout << "tnlVector on MIC test by Satan:" <<endl; //LOL
+
+ #ifdef HAVE_ICPC
+ cout << "ICPC in USE" <<endl; //LOL
+ #endif
+/*
+
+ tnlFile soubor;
+ soubor.open("/home/hanousek/pokus.tnl",tnlReadMode);
+ soubor.close();*/
+
+ //tnlVector<> vct;
+ tnlArray<double> arr;
+
+/*#pragma offload target(mic)
+ {
+ cout << "Hello " <<endl;
+ }
+*/
+
+ //tnlVector<tnlMIC> aa(10);
+
+ //cout << aa <<endl;
+
+
+ return 0;
+}