Commit 2b407a76 authored by Vít Hanousek's avatar Vít Hanousek
Browse files

Implementováno tnlArray na Intel MIC. Vše je zavřeno v #ifdef HAVE_MIC pro... 

Implementováno tnlArray na Intel MIC. Vše je zavřeno v #ifdef HAVE_MIC pro oddělení kódu bez podpory MIC a s podporou MIC.
parent 896e81a8

netbeans_build.sh

+1 −0
Original line number Diff line number Diff line 
#!/bin/bash



module load gcc-5.3.0 cmake-3.4.3 intel_parallel_studio_ex-2016.1

#make -j 6 tnlSatanExperimentalTest-dbg

make -j 6 tnlSatanMICVectorExperimentalTest-dbg

#make -j 6 tnl-image-converter-dbg
 
 No newline at end of file

src/core/arrays/tnlArray.h

+4 −2
Original line number Diff line number Diff line
@@ -187,6 +187,8 @@ template< typename Element, typename Device, typename Index > 
ostream& operator << ( ostream& str, const tnlArray< Element, Device, Index >& v );



#include <core/arrays/tnlArray_impl.h>

//#include <core/arrays/tnlArrayMIC_impl.h>



#ifdef HAVE_MIC

    //MIC specializaton of Araray

    #include <core/arrays/tnlArrayMIC_impl.h>

#endif

#endif /* TNLARRAY_H_ */

src/core/arrays/tnlArrayMIC_impl.h

+1 −1
Original line number Diff line number Diff line
@@ -22,7 +22,7 @@ 
#ifndef TNLARRAYMIC_IMPL_H

#define TNLARRAYMIC_IMPL_H



#ifdef ENABLE_MIC_SATAN

#ifdef HAVE_MIC



#include <core/tnlMIC.h>

src/core/tnlDevice_Callable.h

+1 −1
Original line number Diff line number Diff line
@@ -22,7 +22,7 @@ 
#endif



//NEW, better  __device_callable__ --used only with MIC touch code

#ifdef HAVE_ICPC 

#ifdef HAVE_MIC 

    #define __device_callable__ __attribute__((target(mic)))

#elif HAVE_CUDA

    #define __device_callable__ __device__ __host__

src/core/tnlFile_impl.h

+12 −10
Original line number Diff line number Diff line
@@ -130,10 +130,9 @@ bool tnlFile :: read( Type* buffer, 
#endif

   }

   //MIC

   /*

   if( Device :: getDeviceType() == "tnlMIC" )

   {

       

#ifdef HAVE_MIC           

        Type * host_buffer = (Type *)malloc( sizeof( Type ) * host_buffer_size );

        readElements = 0;

        if( ! host_buffer )
@@ -156,12 +155,12 @@ bool tnlFile :: read( Type* buffer, 
            }

           satanHider<Type> device_buff;

           device_buff.pointer=buffer;

          // #pragma offload target(mic) in(device_buff,readElements) in(host_buffer:length(transfer))

           #pragma offload target(mic) in(device_buff,readElements) in(host_buffer:length(transfer))

           {

               /*

               for(int i=0;i<transfer;i++)

                    device_buff.pointer[readElements+i]=host_buffer[i];

                

                */                

               memcpy(&(device_buff.pointer[readElements]),host_buffer, transfer*sizeof(Type) );

           }
@@ -169,7 +168,8 @@ bool tnlFile :: read( Type* buffer, 
      }

      free( host_buffer );

      return true;

   }*/

#endif

   }

   

   return true;

};
@@ -266,8 +266,9 @@ bool tnlFile :: write( const Type* buffer, 
#endif

   }

   //MIC

   /*if( Device :: getDeviceType() == "tnlMIC" )

   if( Device :: getDeviceType() == "tnlMIC" )

   {

#ifdef HAVE_MIC

         Type * host_buffer = (Type *)malloc( sizeof( Type ) * host_buffer_size );

         if( ! host_buffer )

         {
@@ -282,12 +283,12 @@ bool tnlFile :: write( const Type* buffer, 
            

           satanHider<const Type> device_buff;

           device_buff.pointer=buffer;

           //#pragma offload target(mic) in(device_buff,writtenElements) out(host_buffer:length(transfer))

           #pragma offload target(mic) in(device_buff,writtenElements) out(host_buffer:length(transfer))

           {

               //THIS SHOULD WORK... BUT NOT WHY?

               /*for(int i=0;i<transfer;i++)

                    host_buffer[i]=device_buff.pointer[writtenElements+i];

               

                */              

               

               memcpy(host_buffer,&(device_buff.pointer[writtenElements]), transfer*sizeof(Type) );

            }
@@ -305,7 +306,8 @@ bool tnlFile :: write( const Type* buffer, 
         }

         free( host_buffer );

         return true;

   } */

#endif

   } 

   return true;

};