doplněny nov soubory

netbeans_build.sh

+#!/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

netbeans_run_debug_satanexperimental.sh

+#!/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

src/core/tnlDevice_Callable.h

+/*
+ * 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 */
+

src/core/tnlMIC.h

+/*
+ * 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 */
+

src/core/vectors/tnlVectorOperationsMIC_impl.h

+/*
+ * 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 */
+

tests/unit-tests/core/arrays/tnlSatanExperimentalTest.cpp

+/**
+ *  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;
+}

tests/unit-tests/core/vectors/tnlSatanMICVectorExperimentalTest.cpp

+/**
+ *  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;
+}