diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
index d260486c83a221fb34a161103ad9e5686f98d2d7..a26124c8ac383dd9e7616002678086309f8bbd46 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -21,7 +21,6 @@ stages:
WITH_OPENMP: "no"
WITH_CUDA: "no"
WITH_CUDA_ARCH: "auto"
- WITH_MIC: "no"
WITH_MPI: "no"
# configurations
WITH_TESTS: "no"
@@ -56,7 +55,6 @@ stages:
-DWITH_MPI=${WITH_MPI}
-DWITH_CUDA=${WITH_CUDA}
-DWITH_CUDA_ARCH=${WITH_CUDA_ARCH}
- -DWITH_MIC=${WITH_MIC}
-DWITH_TESTS=${WITH_TESTS}
-DWITH_DOC=${WITH_DOC}
-DWITH_COVERAGE=${WITH_COVERAGE}
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 9540fe0028c4647db0e16cb7c7864c119772c2fd..4d7e0cedf772a46526bac88be06afd6b18681830 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -17,7 +17,6 @@ set( tnlVersion "0.1" )
# declare all custom build options
option(OFFLINE_BUILD "Offline build (i.e. without downloading libraries such as pybind11)" OFF)
-option(WITH_MIC "Build with MIC support" OFF)
option(WITH_CUDA "Build with CUDA support" ON)
set(WITH_CUDA_ARCH "auto" CACHE STRING "Build for these CUDA architectures")
option(WITH_OPENMP "Build with OpenMP support" ON)
@@ -120,22 +119,6 @@ if( NOT DEFINED ENV{CI_JOB_NAME} )
endif()
endif()
-if( CMAKE_CXX_COMPILER_ID STREQUAL "Intel" )
- set( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHAVE_ICPC -wd2568 -wd2571 -wd2570")
- #####
- # Check for MIC
- #
- if( ${WITH_MIC} )
- message( "Enabled MIC support." )
- set( MIC_CXX_FLAGS "-DHAVE_MIC")
- # build all tests with MIC support
- set( CXX_TESTS_FLAGS ${CXX_TESTS_FLAGS} -DHAVE_MIC )
- set( WITH_CUDA OFF CACHE BOOL "Build with CUDA support" )
- else()
- set( MIC_CXX_FLAGS "")
- endif()
-endif()
-
# force colorized output in continuous integration
if( DEFINED ENV{CI_JOB_NAME} OR ${CMAKE_GENERATOR} STREQUAL "Ninja" )
message(STATUS "Continuous integration or Ninja detected -- forcing compilers to produce colorized output.")
@@ -355,7 +338,6 @@ INCLUDE( CPack )
# Print custom build options
message( "-- Build options:" )
message( " OFFLINE_BUILD = ${OFFLINE_BUILD}" )
-message( " WITH_MIC = ${WITH_MIC}" )
message( " WITH_CUDA = ${WITH_CUDA}" )
message( " WITH_CUDA_ARCH = ${WITH_CUDA_ARCH}" )
message( " WITH_OPENMP = ${WITH_OPENMP}" )
diff --git a/build b/build
index c1e0d3162a9585c41a1050d30cecc442fb8d2173..bcd590860ef1e3aaa877a5dcd1cc3de0bb9ad3fe 100755
--- a/build
+++ b/build
@@ -48,7 +48,6 @@ do
--offline-build ) OFFLINE_BUILD="yes" ;;
--with-clang=* ) WITH_CLANG="${option#*=}" ;;
--with-mpi=* ) WITH_MPI="${option#*=}" ;;
- --with-mic=* ) WITH_MIC="${option#*=}" ;;
--with-cuda=* ) WITH_CUDA="${option#*=}" ;;
--with-cuda-arch=* ) WITH_CUDA_ARCH="${option#*=}";;
--with-openmp=* ) WITH_OPENMP="${option#*=}" ;;
@@ -78,7 +77,6 @@ if [[ ${HELP} == "yes" ]]; then
echo " --install=yes/no Enables the installation of TNL files."
echo " --offline-build=yes/no Disables online updates during the build. 'no' by default."
echo " --with-mpi=yes/no Enables MPI. 'yes' by default (OpenMPI required)."
- echo " --with-mic=yes/no Enables MIC (Intel Xeon Phi). 'no' by default (Intel Compiler required)."
echo " --with-cuda=yes/no Enables CUDA. 'yes' by default (CUDA Toolkit is required)."
echo " --with-cuda-arch=all/auto/3.0/3.5/... Chooses CUDA architecture. 'auto' by default."
echo " --with-openmp=yes/no Enables OpenMP. 'yes' by default."
@@ -126,7 +124,6 @@ cmake_command=(
-DCMAKE_BUILD_TYPE=${BUILD}
-DCMAKE_INSTALL_PREFIX=${PREFIX}
-DOFFLINE_BUILD=${OFFLINE_BUILD}
- -DWITH_MIC=${WITH_MIC}
-DWITH_CUDA=${WITH_CUDA}
-DWITH_CUDA_ARCH=${WITH_CUDA_ARCH}
-DWITH_OPENMP=${WITH_OPENMP}
diff --git a/src/Benchmarks/HeatEquation/Tuning/GridTraverser_impl.h b/src/Benchmarks/HeatEquation/Tuning/GridTraverser_impl.h
index f3d9fbeec528dae97e4f3304f44b8440318d4529..816ee5e2c4724137118d221a7487f75651bbc8ac 100644
--- a/src/Benchmarks/HeatEquation/Tuning/GridTraverser_impl.h
+++ b/src/Benchmarks/HeatEquation/Tuning/GridTraverser_impl.h
@@ -8,8 +8,6 @@
/* See Copyright Notice in tnl/Copyright */
-#include <TNL/Devices/MIC.h>
-
#pragma once
#include "GridTraverser.h"
diff --git a/src/Benchmarks/ODESolvers/Euler.hpp b/src/Benchmarks/ODESolvers/Euler.hpp
index 1066e178c2f150c97514eb04dcd19a5a30932102..efb336aca076cb33d0ab5a79cec836d6bfd79e76 100644
--- a/src/Benchmarks/ODESolvers/Euler.hpp
+++ b/src/Benchmarks/ODESolvers/Euler.hpp
@@ -10,7 +10,6 @@
#pragma once
-#include <TNL/Devices/MIC.h>
#include <TNL/Communicators/MpiCommunicator.h>
#include <TNL/Communicators/NoDistrCommunicator.h>
#include "ComputeBlockResidue.h"
@@ -209,28 +208,7 @@ void Euler< Problem, SolverMonitor >::computeNewTimeLevel( DofVectorPointer& u,
}
#endif
}
-
- //MIC
- if( std::is_same< DeviceType, Devices::MIC >::value )
- {
-#ifdef HAVE_MIC
- Devices::MICHider<RealType> mu;
- mu.pointer=_u;
- Devices::MICHider<RealType> mk1;
- mk1.pointer=_k1;
- #pragma offload target(mic) in(mu,mk1,size) inout(localResidue)
- {
- #pragma omp parallel for reduction(+:localResidue) firstprivate( mu, mk1 )
- for( IndexType i = 0; i < size; i ++ )
- {
- const RealType add = tau * mk1.pointer[ i ];
- mu.pointer[ i ] += add;
- localResidue += std::fabs( add );
- }
- }
-#endif
- }
localResidue /= tau * ( RealType ) size;
Problem::CommunicatorType::Allreduce( &localResidue, ¤tResidue, 1, MPI_SUM, Problem::CommunicatorType::AllGroup );
//std::cerr << "Local residue = " << localResidue << " - globalResidue = " << currentResidue << std::endl;
diff --git a/src/Examples/heat-equation/CMakeLists.txt b/src/Examples/heat-equation/CMakeLists.txt
index c895199060481cd4e20e22c30a4a5b473faeb552..979c34076cd5588921cfaea29e10d4ef712f7a79 100644
--- a/src/Examples/heat-equation/CMakeLists.txt
+++ b/src/Examples/heat-equation/CMakeLists.txt
@@ -1,18 +1,17 @@
-set( tnl_heat_equation_SOURCES
+set( tnl_heat_equation_SOURCES
tnl-heat-equation.cpp
tnl-heat-equation-eoc.cpp
tnl-heat-equation.cu
tnl-heat-equation-eoc.cu )
-
+
IF( BUILD_CUDA )
CUDA_ADD_EXECUTABLE(tnl-heat-equation tnl-heat-equation.cu)
CUDA_ADD_EXECUTABLE(tnl-heat-equation-eoc-test tnl-heat-equation-eoc.cu)
target_link_libraries (tnl-heat-equation ${CUSPARSE_LIBRARY} )
target_link_libraries (tnl-heat-equation-eoc-test ${CUSPARSE_LIBRARY} )
-ELSE( BUILD_CUDA )
- ADD_EXECUTABLE(tnl-heat-equation tnl-heat-equation.cpp)
- ADD_EXECUTABLE(tnl-heat-equation-eoc-test tnl-heat-equation-eoc.cpp)
- TARGET_COMPILE_DEFINITIONS( tnl-heat-equation PUBLIC ${MIC_CXX_FLAGS} )
+ELSE( BUILD_CUDA )
+ ADD_EXECUTABLE(tnl-heat-equation tnl-heat-equation.cpp)
+ ADD_EXECUTABLE(tnl-heat-equation-eoc-test tnl-heat-equation-eoc.cpp)
ENDIF( BUILD_CUDA )
@@ -20,7 +19,7 @@ INSTALL( TARGETS tnl-heat-equation
tnl-heat-equation-eoc-test
RUNTIME DESTINATION bin
PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE )
-
+
INSTALL( FILES tnl-run-heat-equation-eoc-test
tnl-run-heat-equation
${tnl_heat_equation_SOURCES}
diff --git a/src/TNL/Allocators/Default.h b/src/TNL/Allocators/Default.h
index 6906a905c3a82d3e2400c4ba6a767848bf1be061..eed5c193b94891497ef9ccaa5b7e0d055f26c080 100644
--- a/src/TNL/Allocators/Default.h
+++ b/src/TNL/Allocators/Default.h
@@ -14,10 +14,8 @@
#include <TNL/Allocators/Host.h>
#include <TNL/Allocators/Cuda.h>
-#include <TNL/Allocators/MIC.h>
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
namespace TNL {
namespace Allocators {
@@ -45,13 +43,5 @@ struct Default< Devices::Cuda >
using Allocator = Allocators::Cuda< T >;
};
-//! Sets \ref Allocators::MIC as the default allocator for \ref Devices::MIC.
-template<>
-struct Default< Devices::MIC >
-{
- template< typename T >
- using Allocator = Allocators::MIC< T >;
-};
-
} // namespace Allocators
} // namespace TNL
diff --git a/src/TNL/Allocators/MIC.h b/src/TNL/Allocators/MIC.h
deleted file mode 100644
index c3599f449cd85f9f83c0ef0e5974bb015d04a6ef..0000000000000000000000000000000000000000
--- a/src/TNL/Allocators/MIC.h
+++ /dev/null
@@ -1,100 +0,0 @@
-/***************************************************************************
- MIC.h - description
- -------------------
- begin : Jul 2, 2019
- copyright : (C) 2019 by Tomas Oberhuber et al.
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Jakub Klinkovsky
-
-#pragma once
-
-#include <TNL/Devices/MIC.h>
-#include <TNL/Exceptions/MICSupportMissing.h>
-
-namespace TNL {
-namespace Allocators {
-
-/**
- * \brief Allocator for the MIC device memory space.
- */
-template< class T >
-struct MIC
-{
- using value_type = T;
- using size_type = std::size_t;
- using difference_type = std::ptrdiff_t;
-
- MIC() = default;
- MIC( const MIC& ) = default;
- MIC( MIC&& ) = default;
-
- MIC& operator=( const MIC& ) = default;
- MIC& operator=( MIC&& ) = default;
-
- template< class U >
- MIC( const MIC< U >& )
- {}
-
- template< class U >
- MIC( MIC< U >&& )
- {}
-
- template< class U >
- MIC& operator=( const MIC< U >& )
- {
- return *this;
- }
-
- template< class U >
- MIC& operator=( MIC< U >&& )
- {
- return *this;
- }
-
- value_type* allocate( size_type size )
- {
-#ifdef HAVE_MIC
- Devices::MICHider<void> hide_ptr;
- #pragma offload target(mic) out(hide_ptr) in(size)
- {
- hide_ptr.pointer = malloc(size * sizeof(value_type));
- }
- return hide_ptr.pointer;
-#else
- throw Exceptions::MICSupportMissing();
-#endif
- }
-
- void deallocate(value_type* ptr, size_type)
- {
-#ifdef HAVE_MIC
- Devices::MICHider<void> hide_ptr;
- hide_ptr.pointer=ptr;
- #pragma offload target(mic) in(hide_ptr)
- {
- free(hide_ptr.pointer);
- }
-#else
- throw Exceptions::MICSupportMissing();
-#endif
- }
-};
-
-template<class T1, class T2>
-bool operator==(const MIC<T1>&, const MIC<T2>&)
-{
- return true;
-}
-
-template<class T1, class T2>
-bool operator!=(const MIC<T1>& lhs, const MIC<T2>& rhs)
-{
- return !(lhs == rhs);
-}
-
-} // namespace Allocators
-} // namespace TNL
diff --git a/src/TNL/Assert.h b/src/TNL/Assert.h
index 27f3b11b28ed46a9741c3593c573a243f1e0a81d..3d91c8c763b8c9532713f90beee3a02fd2c64b6a 100644
--- a/src/TNL/Assert.h
+++ b/src/TNL/Assert.h
@@ -38,7 +38,7 @@
#define TNL_NVCC_HD_WARNING_DISABLE
#endif
-#if defined(NDEBUG) || defined(HAVE_MIC)
+#ifdef NDEBUG
// empty macros for optimized build
/**
diff --git a/src/TNL/Containers/Algorithms/ArrayOperations.h b/src/TNL/Containers/Algorithms/ArrayOperations.h
index ca62f5b7ea45254298cb02d0ac909ee2242e72f2..d4c35f5b1a14352a3c978217d61b2895e70405b4 100644
--- a/src/TNL/Containers/Algorithms/ArrayOperations.h
+++ b/src/TNL/Containers/Algorithms/ArrayOperations.h
@@ -12,7 +12,6 @@
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
namespace TNL {
namespace Containers {
@@ -194,95 +193,6 @@ struct ArrayOperations< Devices::Host, Devices::Cuda >
const Index size );
};
-
-template<>
-struct ArrayOperations< Devices::MIC >
-{
- template< typename Element >
- static void setElement( Element* data,
- const Element& value );
-
- template< typename Element >
- static Element getElement( const Element* data );
-
- template< typename Element, typename Index >
- static void set( Element* data,
- const Element& value,
- const Index size );
-
- template< typename DestinationElement,
- typename SourceElement,
- typename Index >
- static void copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size );
-
- template< typename DestinationElement,
- typename Index,
- typename SourceIterator >
- static void copyFromIterator( DestinationElement* destination,
- Index destinationSize,
- SourceIterator first,
- SourceIterator last );
-
- template< typename Element1,
- typename Element2,
- typename Index >
- static bool compare( const Element1* destination,
- const Element2* source,
- const Index size );
-
- template< typename Element,
- typename Index >
- static bool containsValue( const Element* data,
- const Index size,
- const Element& value );
-
- template< typename Element,
- typename Index >
- static bool containsOnlyValue( const Element* data,
- const Index size,
- const Element& value );
-};
-
-template<>
-struct ArrayOperations< Devices::MIC, Devices::Host >
-{
- public:
-
- template< typename DestinationElement,
- typename SourceElement,
- typename Index >
- static void copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size );
-
- template< typename DestinationElement,
- typename SourceElement,
- typename Index >
- static bool compare( const DestinationElement* destination,
- const SourceElement* source,
- const Index size );
-};
-
-template<>
-struct ArrayOperations< Devices::Host, Devices::MIC >
-{
- template< typename DestinationElement,
- typename SourceElement,
- typename Index >
- static void copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size );
-
- template< typename DestinationElement,
- typename SourceElement,
- typename Index >
- static bool compare( const DestinationElement* destination,
- const SourceElement* source,
- const Index size );
-};
-
} // namespace Algorithms
} // namespace Containers
} // namespace TNL
@@ -290,4 +200,3 @@ struct ArrayOperations< Devices::Host, Devices::MIC >
#include <TNL/Containers/Algorithms/ArrayOperationsStatic.hpp>
#include <TNL/Containers/Algorithms/ArrayOperationsHost.hpp>
#include <TNL/Containers/Algorithms/ArrayOperationsCuda.hpp>
-#include <TNL/Containers/Algorithms/ArrayOperationsMIC.hpp>
diff --git a/src/TNL/Containers/Algorithms/ArrayOperationsMIC.hpp b/src/TNL/Containers/Algorithms/ArrayOperationsMIC.hpp
deleted file mode 100644
index 4113bbcd90f0edce53d143cf65996a392c2a91b4..0000000000000000000000000000000000000000
--- a/src/TNL/Containers/Algorithms/ArrayOperationsMIC.hpp
+++ /dev/null
@@ -1,429 +0,0 @@
-/***************************************************************************
- ArrayOperationsMIC_impl.h - description
- -------------------
- begin : Mar 4, 2017
- copyright : (C) 2017 by Tomas Oberhuber
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Vit Hanousek
-
-#pragma once
-
-#include <iostream>
-
-#include <TNL/Math.h>
-#include <TNL/Exceptions/MICSupportMissing.h>
-#include <TNL/Containers/Algorithms/ArrayOperations.h>
-#include <TNL/Containers/Algorithms/Reduction.h>
-#include <TNL/Exceptions/NotImplementedError.h>
-
-namespace TNL {
-namespace Containers {
-namespace Algorithms {
-
-static constexpr std::size_t MIC_STACK_VAR_LIM = 5*1024*1024;
-
-template< typename Element >
-void
-ArrayOperations< Devices::MIC >::
-setElement( Element* data,
- const Element& value )
-{
- TNL_ASSERT( data, );
- ArrayOperations< Devices::MIC >::set( data, value, 1 );
-}
-
-template< typename Element >
-Element
-ArrayOperations< Devices::MIC >::
-getElement( const Element* data )
-{
- TNL_ASSERT( data, );
- Element result;
- ArrayOperations< Devices::Host, Devices::MIC >::copy< Element, Element, int >( &result, data, 1 );
- return result;
-}
-
-template< typename Element, typename Index >
-void
-ArrayOperations< Devices::MIC >::
-set( Element* data,
- const Element& value,
- const Index size )
-{
- TNL_ASSERT( data, );
-#ifdef HAVE_MIC
- Element tmp=value;
- Devices::MICHider<Element> hide_ptr;
- hide_ptr.pointer=data;
- #pragma offload target(mic) in(hide_ptr,tmp,size)
- {
- Element * dst= hide_ptr.pointer;
- for(int i=0;i<size;i++)
- dst[i]=tmp;
- }
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-template< typename DestinationElement,
- typename SourceElement,
- typename Index >
-void
-ArrayOperations< Devices::MIC >::
-copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
-{
- TNL_ASSERT( destination, );
- TNL_ASSERT( source, );
- #ifdef HAVE_MIC
- if( std::is_same< DestinationElement, SourceElement >::value )
- {
- Devices::MICHider<void> src_ptr;
- src_ptr.pointer=(void*)source;
- Devices::MICHider<void> dst_ptr;
- dst_ptr.pointer=(void*)destination;
- #pragma offload target(mic) in(src_ptr,dst_ptr,size)
- {
- memcpy(dst_ptr.pointer,src_ptr.pointer,size*sizeof(DestinationElement));
- }
- }
- else
- {
- Devices::MICHider<const SourceElement> src_ptr;
- src_ptr.pointer=source;
- Devices::MICHider<DestinationElement> dst_ptr;
- dst_ptr.pointer=destination;
- #pragma offload target(mic) in(src_ptr,dst_ptr,size)
- {
- for(int i=0;i<size;i++)
- dst_ptr.pointer[i]=src_ptr.pointer[i];
- }
- }
- #else
- throw Exceptions::MICSupportMissing();
- #endif
-}
-
-template< typename DestinationElement,
- typename Index,
- typename SourceIterator >
-void
-ArrayOperations< Devices::MIC >::
-copyFromIterator( DestinationElement* destination,
- Index destinationSize,
- SourceIterator first,
- SourceIterator last )
-{
- throw Exceptions::NotImplementedError();
-}
-
-template< typename Element1,
- typename Element2,
- typename Index >
-bool
-ArrayOperations< Devices::MIC >::
-compare( const Element1* destination,
- const Element2* source,
- const Index size )
-{
- TNL_ASSERT( destination, );
- TNL_ASSERT( source, );
-#ifdef HAVE_MIC
- if( std::is_same< Element1, Element2 >::value )
- {
- Devices::MICHider<void> src_ptr;
- src_ptr.pointer=(void*)source;
- Devices::MICHider<void> dst_ptr;
- dst_ptr.pointer=(void*)destination;
- int ret=0;
- #pragma offload target(mic) in(src_ptr,dst_ptr,size) out(ret)
- {
- ret=memcmp(dst_ptr.pointer,src_ptr.pointer,size*sizeof(Element1));
- }
- if(ret==0)
- return true;
- }
- else
- {
- Devices::MICHider<const Element2> src_ptr;
- src_ptr.pointer=source;
- Devices::MICHider<const Element1> dst_ptr;
- dst_ptr.pointer=destination;
- bool ret=false;
- #pragma offload target(mic) in(src_ptr,dst_ptr,size) out(ret)
- {
- int i=0;
- for(i=0;i<size;i++)
- if(dst_ptr.pointer[i]!=src_ptr.pointer[i])
- break;
- if(i==size)
- ret=true;
- else
- ret=false;
- }
- return ret;
- }
- return false;
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-template< typename Element,
- typename Index >
-bool
-ArrayOperations< Devices::MIC >::
-containsValue( const Element* data,
- const Index size,
- const Element& value )
-{
- TNL_ASSERT_TRUE( data, "Attempted to check data through a nullptr." );
- TNL_ASSERT_GE( size, 0, "" );
-#ifdef HAVE_MIC
- throw Exceptions::NotImplementedError();
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-template< typename Element,
- typename Index >
-bool
-ArrayOperations< Devices::MIC >::
-containsOnlyValue( const Element* data,
- const Index size,
- const Element& value )
-{
- TNL_ASSERT_TRUE( data, "Attempted to check data through a nullptr." );
- TNL_ASSERT_GE( size, 0, "" );
-#ifdef HAVE_MIC
- throw Exceptions::NotImplementedError();
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-
-
-/****
- * Operations MIC -> Host
- */
-
-template< typename DestinationElement,
- typename SourceElement,
- typename Index >
-void
-ArrayOperations< Devices::Host, Devices::MIC >::
-copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
-{
- TNL_ASSERT( destination, );
- TNL_ASSERT( source, );
-#ifdef HAVE_MIC
- if( std::is_same< DestinationElement, SourceElement >::value )
- {
- Devices::MICHider<void> src_ptr;
- src_ptr.pointer=(void*)source;
-
- //JAKA KONSTANTA se vejde do stacku 5MB?
- if(size<MIC_STACK_VAR_LIM)
- {
- uint8_t tmp[size*sizeof(SourceElement)];
-
- #pragma offload target(mic) in(src_ptr,size) out(tmp)
- {
- memcpy((void*)&tmp,src_ptr.pointer,size*sizeof(SourceElement));
- }
-
- memcpy((void*)destination,(void*)&tmp,size*sizeof(SourceElement));
- }
- else
- {
- //direct -- pomalejšĂ
- uint8_t* tmp=(uint8_t*)destination;
- #pragma offload target(mic) in(src_ptr,size) out(tmp:length(size))
- {
- memcpy((void*)tmp,src_ptr.pointer,size*sizeof(SourceElement));
- }
- }
- }
- else
- {
- Devices::MICHider<const SourceElement> src_ptr;
- src_ptr.pointer=source;
-
- if(size<MIC_STACK_VAR_LIM)
- {
- uint8_t tmp[size*sizeof(DestinationElement)];
-
- #pragma offload target(mic) in(src_ptr,size) out(tmp)
- {
- DestinationElement *dst=(DestinationElement*)&tmp;
- for(int i=0;i<size;i++)
- dst[i]=src_ptr.pointer[i];
- }
-
- memcpy((void*)destination,(void*)&tmp,size*sizeof(DestinationElement));
- }
- else
- {
- //direct pseudo heap-- pomalejšĂ
- uint8_t* tmp=(uint8_t*)destination;
- #pragma offload target(mic) in(src_ptr,size) out(tmp:length(size*sizeof(DestinationElement)))
- {
- DestinationElement *dst=(DestinationElement*)tmp;
- for(int i=0;i<size;i++)
- dst[i]=src_ptr.pointer[i];
- }
- }
- }
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-
-template< typename Element1,
- typename Element2,
- typename Index >
-bool
-ArrayOperations< Devices::Host, Devices::MIC >::
-compare( const Element1* destination,
- const Element2* source,
- const Index size )
-{
- /***
- * Here, destination is on host and source is on MIC device.
- */
- TNL_ASSERT( destination, );
- TNL_ASSERT( source, );
- TNL_ASSERT( size >= 0, std::cerr << "size = " << size );
-#ifdef HAVE_MIC
- Index compared( 0 );
- Index transfer( 0 );
- std::size_t max_transfer=MIC_STACK_VAR_LIM/sizeof(Element2);
- uint8_t host_buffer[max_transfer*sizeof(Element2)];
-
- Devices::MICHider<const Element2> src_ptr;
-
- while( compared < size )
- {
- transfer=min(size-compared,max_transfer);
- src_ptr.pointer=source+compared;
- #pragma offload target(mic) out(host_buffer) in(src_ptr,transfer)
- {
- memcpy((void*)&host_buffer,(void*)src_ptr.pointer,transfer*sizeof(Element2));
- }
- if( ! ArrayOperations< Devices::Host >::compare( &destination[ compared ], (Element2*)&host_buffer, transfer ) )
- {
- return false;
- }
- compared += transfer;
- }
- return true;
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-/****
- * Operations Host -> MIC
- */
-template< typename DestinationElement,
- typename SourceElement,
- typename Index >
-void
-ArrayOperations< Devices::MIC, Devices::Host >::
-copy( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
-{
- TNL_ASSERT( destination, );
- TNL_ASSERT( source, );
- TNL_ASSERT( size >= 0, std::cerr << "size = " << size );
-#ifdef HAVE_MIC
- if( std::is_same< DestinationElement, SourceElement >::value )
- {
- Devices::MICHider<void> dst_ptr;
- dst_ptr.pointer=(void*)destination;
-
- //JAKA KONSTANTA se vejde do stacku 5MB?
- if(size<MIC_STACK_VAR_LIM)
- {
- uint8_t tmp[size*sizeof(SourceElement)];
- memcpy((void*)&tmp,(void*)source,size*sizeof(SourceElement));
-
- #pragma offload target(mic) in(dst_ptr,tmp,size)
- {
- memcpy(dst_ptr.pointer,(void*)&tmp,size*sizeof(SourceElement));
- }
- }
- else
- {
- //direct pseudo heap-- pomalejšĂ
- uint8_t* tmp=(uint8_t*)source;
- #pragma offload target(mic) in(dst_ptr,size) in(tmp:length(size))
- {
- memcpy(dst_ptr.pointer,(void*)tmp,size*sizeof(SourceElement));
- }
- }
- }
- else
- {
- Devices::MICHider<DestinationElement> dst_ptr;
- dst_ptr.pointer=destination;
-
- if(size<MIC_STACK_VAR_LIM)
- {
- uint8_t tmp[size*sizeof(SourceElement)];
- memcpy((void*)&tmp,(void*)source,size*sizeof(SourceElement));
-
- #pragma offload target(mic) in(dst_ptr,size,tmp)
- {
- SourceElement *src=(SourceElement*)&tmp;
- for(int i=0;i<size;i++)
- dst_ptr.pointer[i]=src[i];
- }
- }
- else
- {
- //direct pseudo heap-- pomalejšĂ
- uint8_t* tmp=(uint8_t*)source;
- #pragma offload target(mic) in(dst_ptr,size) in(tmp:length(size*sizeof(SourceElement)))
- {
- SourceElement *src=(SourceElement*)tmp;
- for(int i=0;i<size;i++)
- dst_ptr.pointer[i]=src[i];
- }
- }
- }
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
-template< typename Element1,
- typename Element2,
- typename Index >
-bool
-ArrayOperations< Devices::MIC, Devices::Host >::
-compare( const Element1* hostData,
- const Element2* deviceData,
- const Index size )
-{
- TNL_ASSERT( hostData, );
- TNL_ASSERT( deviceData, );
- TNL_ASSERT( size >= 0, std::cerr << "size = " << size );
- return ArrayOperations< Devices::Host, Devices::MIC >::compare( deviceData, hostData, size );
-}
-
-} // namespace Algorithms
-} // namespace Containers
-} // namespace TNL
diff --git a/src/TNL/Devices/CudaCallable.h b/src/TNL/Devices/CudaCallable.h
index f9311443f12a0c85fb6fba9ebaf07ca47736b030..f63e4e430e01e76df7833001a651e4f220e9bab6 100644
--- a/src/TNL/Devices/CudaCallable.h
+++ b/src/TNL/Devices/CudaCallable.h
@@ -20,11 +20,9 @@
* This macro serves for definition of function which are supposed to be called
* even from device. If HAVE_CUDA is defined, the __cuda_callable__ function
* is compiled for both CPU and GPU. If HAVE_CUDA is not defined, this macro has
- * no effect. Support for Intel Xeon Phi is now in "hibernated" state.
+ * no effect.
*/
-#ifdef HAVE_MIC
- #define __cuda_callable__ __attribute__((target(mic)))
-#elif HAVE_CUDA
+#ifdef HAVE_CUDA
#define __cuda_callable__ __device__ __host__
#else
#define __cuda_callable__
diff --git a/src/TNL/Devices/MIC.h b/src/TNL/Devices/MIC.h
deleted file mode 100644
index f347a24d1f9e4fa6d5cceb7e2693807c7158065a..0000000000000000000000000000000000000000
--- a/src/TNL/Devices/MIC.h
+++ /dev/null
@@ -1,170 +0,0 @@
-/***************************************************************************
- MIC.h - description
- -------------------
- begin : Nov 7, 2016
- copyright : (C) 2016 by Tomas Oberhuber
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Vit Hanousek
-
-#pragma once
-
-#include <iostream>
-#include <cstring>
-#include <unistd.h>
-#include <TNL/String.h>
-#include <TNL/Assert.h>
-#include <TNL/Pointers/SmartPointersRegister.h>
-#include <TNL/Timer.h>
-
-#include <TNL/Devices/CudaCallable.h>
-
-
-namespace TNL {
-namespace Devices {
-namespace {
-
-//useful macros from Intel's tutorials -- but we do not use it, becaouse it is tricky (system of maping variables CPU-MIC)
-#define ALLOC alloc_if(1) //alloac variable at begining of offloaded block -- default
-#define FREE free_if(1) // delete variable at the end of offloaded block -- default
-#define RETAIN free_if(0) //do not delete variable at the end of offladed block
-#define REUSE alloc_if(0) //do not alloc variable at begin of offloaded block, reuse variable on MIC which was not deleted befeore
-
-//structure which hides pointer - bypass mapping of variables and addresses of arrays and allow get RAW addres of MIC memory to RAM
-template< typename Type >
-struct MICHider{
- Type *pointer;
-};
-
-//inflatable structure -- structures can be copied to MIC - classes not (viz paper published after CSJP 2016 in Krakow)
-//object can be copied in side this structure and then copied into MIC memory
-template <unsigned int VELIKOST>
-struct MICStruct{
- uint8_t data[VELIKOST];
-};
-
-//Macros which can make code better readeble --but they are tricky, creating variables with specific names...
-//version using inflatable structure
-#define TNLMICSTRUCT(bb,typ) Devices::MICStruct<sizeof(typ)> s ## bb; \
- memcpy((void*)& s ## bb,(void*)& bb,sizeof(typ));
-#define TNLMICSTRUCTOFF(bb,typ) s ## bb
-#define TNLMICSTRUCTUSE(bb,typ) typ * kernel ## bb = (typ*) &s ## bb;
-#define TNLMICSTRUCTALLOC(bb,typ) typ * kernel ## bb = (typ*) malloc (sizeof(typ)); \
- memcpy((void*)kernel ## bb,(void*) & s ## bb, sizeof(typ));
-
-//version which retypes pointer of object to pointer to array of uint8_t,
-//object can be copied using uint8_t pointer as array with same length as object size
-#define TNLMICHIDE(bb,typ) uint8_t * u ## bb=(uint8_t *)&bb; \
- MICHider<typ> kernel ## bb;
-#define TNLMICHIDEALLOCOFF(bb,typ) in(u ## bb:length(sizeof(typ))) out(kernel ## bb)
-#define TNLMICHIDEALLOC(bb,typ) kernel ## bb.pointer=(typ*)malloc(sizeof(typ)); \
- memcpy((void*)kernel ## bb.pointer,(void*)u ## bb,sizeof(typ));
-#define TNLMICHIDEFREEOFF(bb,typ) in(kernel ## bb)
-#define TNLMICHIDEFREE(bb,typ) free((void*)kernel ## bb.pointer
-
-class MIC
-{
- public:
-
- static String getDeviceType()
- {
- return String( "Devices::MIC" );
- };
-
- // TODO: Remove getDeviceType();
- static inline String getType() { return getDeviceType(); };
-
-#ifdef HAVE_MIC
-
- //useful debuging -- but produce warning
- __cuda_callable__ static inline void CheckMIC(void)
- {
- #ifdef __MIC__
- std::cout<<"ON MIC"<<std::endl;
- #else
- std::cout<<"ON CPU" <<std::endl;
- #endif
- };
-
-
- //old copying funciton -- deprecated
- template <typename TYP>
- static
- TYP * passToDevice(TYP &objektCPU)
- {
- uint8_t * uk=(uint8_t *)&objektCPU;
- MICHider<TYP> ret;
-
- #pragma offload target(mic) in(uk:length(sizeof(TYP))) out(ret)
- {
- ret.pointer=(TYP*)malloc(sizeof(TYP));
- std::memcpy((void*)ret.pointer,(void*)uk,sizeof(TYP));
- }
- return ret.pointer;
-
- std::cout << "Někdo mně volá :-D" <<std::endl;
- };
-
- //old cleaning function -- deprecated
- template <typename TYP>
- static
- void freeFromDevice(TYP *objektMIC)
- {
- MICHider<TYP> ptr;
- ptr.pointer=objektMIC;
- #pragma offload target(mic) in(ptr)
- {
- free((void*)ptr.pointer);
- }
- };
-
- static inline
- void CopyToMIC(void* mic_ptr,void* ptr,size_t size)
- {
- uint8_t image[size];
- std::memcpy((void*)&image,ptr,size);
- Devices::MICHider<void> hide_ptr;
- hide_ptr.pointer=mic_ptr;
- #pragma offload target(mic) in(hide_ptr) in(image) in(size)
- {
- std::memcpy((void*)hide_ptr.pointer,(void*)&image,size);
- }
- };
-
-#endif
-
- static void insertSmartPointer( Pointers::SmartPointer* pointer )
- {
- smartPointersRegister.insert( pointer, -1 );
- }
-
- static void removeSmartPointer( Pointers::SmartPointer* pointer )
- {
- smartPointersRegister.remove( pointer, -1 );
- }
-
- // Negative deviceId means that CudaDeviceInfo::getActiveDevice will be
- // called to get the device ID.
- static bool synchronizeDevice( int deviceId = -1 )
- {
- smartPointersSynchronizationTimer.start();
- bool b = smartPointersRegister.synchronizeDevice( deviceId );
- smartPointersSynchronizationTimer.stop();
- return b;
- }
-
- static Timer smartPointersSynchronizationTimer;
-
-protected:
- static Pointers::SmartPointersRegister smartPointersRegister;
-};
-
-Pointers::SmartPointersRegister MIC::smartPointersRegister;
-Timer MIC::smartPointersSynchronizationTimer;
-
-} // namespace <unnamed>
-} // namespace Devices
-} // namespace TNL
diff --git a/src/TNL/Exceptions/MICBadAlloc.h b/src/TNL/Exceptions/MICBadAlloc.h
deleted file mode 100644
index b8f3a9157c54d8155652a42a700ad71a221aa201..0000000000000000000000000000000000000000
--- a/src/TNL/Exceptions/MICBadAlloc.h
+++ /dev/null
@@ -1,31 +0,0 @@
-/***************************************************************************
- MICBadAlloc.h - description
- -------------------
- begin : Jul 31, 2017
- copyright : (C) 2017 by Tomas Oberhuber et al.
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Jakub Klinkovsky
-
-#pragma once
-
-#include <new>
-
-namespace TNL {
-namespace Exceptions {
-
-struct MICBadAlloc
- : public std::bad_alloc
-{
- const char* what() const throw()
- {
- return "Failed to allocate memory on the MIC device: "
- "most likely there is not enough space on the device memory.";
- }
-};
-
-} // namespace Exceptions
-} // namespace TNL
diff --git a/src/TNL/Exceptions/MICSupportMissing.h b/src/TNL/Exceptions/MICSupportMissing.h
deleted file mode 100644
index 6d4260e6addbbb9dd89a7c9d5a07833485c6a0c2..0000000000000000000000000000000000000000
--- a/src/TNL/Exceptions/MICSupportMissing.h
+++ /dev/null
@@ -1,30 +0,0 @@
-/***************************************************************************
- MICSupportMissing.h - description
- -------------------
- begin : Jul 31, 2017
- copyright : (C) 2017 by Tomas Oberhuber et al.
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Jakub Klinkovsky
-
-#pragma once
-
-#include <stdexcept>
-
-namespace TNL {
-namespace Exceptions {
-
-struct MICSupportMissing
- : public std::runtime_error
-{
- MICSupportMissing()
- : std::runtime_error( "MIC support is missing, but the program called a function which needs it. "
- "Please recompile the program with MIC support." )
- {}
-};
-
-} // namespace Exceptions
-} // namespace TNL
diff --git a/src/TNL/File.h b/src/TNL/File.h
index 1aa5615e5cbbf8f36d2c9ac3d98bdbc3ba4ada03..70eb013b770570f2f694151420e6bdf8c482bf27 100644
--- a/src/TNL/File.h
+++ b/src/TNL/File.h
@@ -16,7 +16,6 @@
#include <TNL/String.h>
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
namespace TNL {
@@ -154,14 +153,6 @@ class File
typename = void >
void load_impl( Type* buffer, std::streamsize elements );
- template< typename Type,
- typename SourceType,
- typename Device,
- typename = typename std::enable_if< std::is_same< Device, Devices::MIC >::value >::type,
- typename = void,
- typename = void >
- void load_impl( Type* buffer, std::streamsize elements );
-
template< typename Type,
typename TargetType,
typename Device,
@@ -175,14 +166,6 @@ class File
typename = void >
void save_impl( const Type* buffer, std::streamsize elements );
- template< typename Type,
- typename TargetType,
- typename Device,
- typename = typename std::enable_if< std::is_same< Device, Devices::MIC >::value >::type,
- typename = void,
- typename = void >
- void save_impl( const Type* buffer, std::streamsize elements );
-
std::fstream file;
String fileName;
diff --git a/src/TNL/File.hpp b/src/TNL/File.hpp
index f4edd2b9638e0331da973bdf16de06568d4b7c23..19a9eaa06a3a88e264de3b4528730baa0304a897 100644
--- a/src/TNL/File.hpp
+++ b/src/TNL/File.hpp
@@ -18,7 +18,6 @@
#include <TNL/File.h>
#include <TNL/Assert.h>
#include <TNL/Exceptions/CudaSupportMissing.h>
-#include <TNL/Exceptions/MICSupportMissing.h>
#include <TNL/Exceptions/FileSerializationError.h>
#include <TNL/Exceptions/FileDeserializationError.h>
#include <TNL/Exceptions/NotImplementedError.h>
@@ -168,48 +167,6 @@ void File::load_impl( Type* buffer, std::streamsize elements )
#endif
}
-// MIC
-template< typename Type,
- typename SourceType,
- typename Device,
- typename, typename, typename >
-void File::load_impl( Type* buffer, std::streamsize elements )
-{
-#ifdef HAVE_MIC
- const std::streamsize host_buffer_size = std::min( TransferBufferSize / (std::streamsize) sizeof(Type), elements );
- using BaseType = typename std::remove_cv< Type >::type;
- std::unique_ptr< BaseType[] > host_buffer{ new BaseType[ host_buffer_size ] };
-
- std::streamsize readElements = 0;
- if( std::is_same< Type, SourceType >::value )
- {
- while( readElements < elements )
- {
- const std::streamsize transfer = std::min( elements - readElements, host_buffer_size );
- file.read( reinterpret_cast<char*>(host_buffer.get()), sizeof(Type) * transfer );
-
- Devices::MICHider<Type> device_buff;
- device_buff.pointer=buffer;
- #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.get(), transfer*sizeof(Type) );
- }
-
- readElements += transfer;
- }
- free( host_buffer );
- }
- else
- throw Exceptions::NotImplementedError("Type conversion during loading is not implemented for MIC.");
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
template< typename Type,
typename TargetType,
typename Device >
@@ -303,48 +260,6 @@ void File::save_impl( const Type* buffer, std::streamsize elements )
#endif
}
-// MIC
-template< typename Type,
- typename TargetType,
- typename Device,
- typename, typename, typename >
-void File::save_impl( const Type* buffer, std::streamsize elements )
-{
-#ifdef HAVE_MIC
- const std::streamsize host_buffer_size = std::min( TransferBufferSize / (std::streamsize) sizeof(Type), elements );
- using BaseType = typename std::remove_cv< Type >::type;
- std::unique_ptr< BaseType[] > host_buffer{ new BaseType[ host_buffer_size ] };
-
- std::streamsize writtenElements = 0;
- if( std::is_same< Type, TargetType >::value )
- {
- while( this->writtenElements < elements )
- {
- const std::streamsize transfer = std::min( elements - writtenElements, host_buffer_size );
-
- Devices::MICHider<const Type> device_buff;
- device_buff.pointer=buffer;
- #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.get(), &(device_buff.pointer[writtenElements]), transfer*sizeof(Type) );
- }
-
- file.write( reinterpret_cast<const char*>(host_buffer.get()), sizeof(Type) * transfer );
- writtenElements += transfer;
- }
- }
- else
- throw Exceptions::NotImplementedError("Type conversion during saving is not implemented for MIC.");
-#else
- throw Exceptions::MICSupportMissing();
-#endif
-}
-
inline bool fileExists( const String& fileName )
{
std::fstream file;
diff --git a/src/TNL/Math.h b/src/TNL/Math.h
index cd73b020e4c35fbe08c969864c9c26e400bd76ef..b7591bf65942b8070ec90f6d8a440cecab4807b6 100644
--- a/src/TNL/Math.h
+++ b/src/TNL/Math.h
@@ -30,7 +30,7 @@ ResultType sum( const T1& a, const T2& b )
* \brief This function returns minimum of two numbers.
*
* GPU device code uses the functions defined in the CUDA's math_functions.h,
- * MIC uses trivial override and host uses the STL functions.
+ * host uses the STL functions.
*/
template< typename T1, typename T2, typename ResultType = typename std::common_type< T1, T2 >::type,
// enable_if is necessary to avoid ambiguity in vector expressions
@@ -44,8 +44,6 @@ ResultType min( const T1& a, const T2& b )
#else
#if defined(__CUDA_ARCH__)
return ::min( (ResultType) a, (ResultType) b );
- #elif defined(__MIC__)
- return a < b ? a : b;
#else
return std::min( (ResultType) a, (ResultType) b );
#endif
@@ -57,7 +55,7 @@ ResultType min( const T1& a, const T2& b )
* \brief This function returns maximum of two numbers.
*
* GPU device code uses the functions defined in the CUDA's math_functions.h,
- * MIC uses trivial override and host uses the STL functions.
+ * host uses the STL functions.
*/
template< typename T1, typename T2, typename ResultType = typename std::common_type< T1, T2 >::type,
// enable_if is necessary to avoid ambiguity in vector expressions
@@ -71,8 +69,6 @@ ResultType max( const T1& a, const T2& b )
#else
#if defined(__CUDA_ARCH__)
return ::max( (ResultType) a, (ResultType) b );
- #elif defined(__MIC__)
- return a > b ? a : b;
#else
return std::max( (ResultType) a, (ResultType) b );
#endif
@@ -92,10 +88,6 @@ T abs( const T& n )
return ::abs( n );
else
return ::fabs( n );
-#elif defined(__MIC__)
- if( n < ( T ) 0 )
- return -n;
- return n;
#else
return std::abs( n );
#endif
@@ -159,7 +151,7 @@ template< typename T1, typename T2, typename ResultType = typename std::common_t
__cuda_callable__ inline
ResultType pow( const T1& base, const T2& exp )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::pow( (ResultType) base, (ResultType) exp );
#else
return std::pow( (ResultType) base, (ResultType) exp );
@@ -173,7 +165,7 @@ template< typename T >
__cuda_callable__ inline
auto exp( const T& value ) -> decltype( std::exp(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::exp( value );
#else
return std::exp( value );
@@ -187,7 +179,7 @@ template< typename T >
__cuda_callable__ inline
auto sqrt( const T& value ) -> decltype( std::sqrt(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::sqrt( value );
#else
return std::sqrt( value );
@@ -201,7 +193,7 @@ template< typename T >
__cuda_callable__ inline
auto cbrt( const T& value ) -> decltype( std::cbrt(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::cbrt( value );
#else
return std::cbrt( value );
@@ -215,7 +207,7 @@ template< typename T >
__cuda_callable__ inline
auto log( const T& value ) -> decltype( std::log(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::log( value );
#else
return std::log( value );
@@ -229,7 +221,7 @@ template< typename T >
__cuda_callable__ inline
auto log10( const T& value ) -> decltype( std::log10(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::log10( value );
#else
return std::log10( value );
@@ -243,7 +235,7 @@ template< typename T >
__cuda_callable__ inline
auto log2( const T& value ) -> decltype( std::log2(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::log2( value );
#else
return std::log2( value );
@@ -257,7 +249,7 @@ template< typename T >
__cuda_callable__ inline
auto sin( const T& value ) -> decltype( std::sin(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::sin( value );
#else
return std::sin( value );
@@ -271,7 +263,7 @@ template< typename T >
__cuda_callable__ inline
auto cos( const T& value ) -> decltype( std::cos(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::cos( value );
#else
return std::cos( value );
@@ -285,7 +277,7 @@ template< typename T >
__cuda_callable__ inline
auto tan( const T& value ) -> decltype( std::tan(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::tan( value );
#else
return std::tan( value );
@@ -299,7 +291,7 @@ template< typename T >
__cuda_callable__ inline
auto asin( const T& value ) -> decltype( std::asin(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::asin( value );
#else
return std::asin( value );
@@ -313,7 +305,7 @@ template< typename T >
__cuda_callable__ inline
auto acos( const T& value ) -> decltype( std::acos(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::acos( value );
#else
return std::acos( value );
@@ -327,7 +319,7 @@ template< typename T >
__cuda_callable__ inline
auto atan( const T& value ) -> decltype( std::atan(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::atan( value );
#else
return std::atan( value );
@@ -341,7 +333,7 @@ template< typename T >
__cuda_callable__ inline
auto sinh( const T& value ) -> decltype( std::sinh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::sinh( value );
#else
return std::sinh( value );
@@ -355,7 +347,7 @@ template< typename T >
__cuda_callable__ inline
auto cosh( const T& value ) -> decltype( std::cosh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::cosh( value );
#else
return std::cosh( value );
@@ -369,7 +361,7 @@ template< typename T >
__cuda_callable__ inline
auto tanh( const T& value ) -> decltype( std::tanh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::tanh( value );
#else
return std::tanh( value );
@@ -383,7 +375,7 @@ template< typename T >
__cuda_callable__ inline
auto asinh( const T& value ) -> decltype( std::asinh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::asinh( value );
#else
return std::asinh( value );
@@ -397,7 +389,7 @@ template< typename T >
__cuda_callable__ inline
auto acosh( const T& value ) -> decltype( std::acosh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::acosh( value );
#else
return std::acosh( value );
@@ -411,7 +403,7 @@ template< typename T >
__cuda_callable__ inline
auto atanh( const T& value ) -> decltype( std::atanh(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::atanh( value );
#else
return std::atanh( value );
@@ -425,7 +417,7 @@ template< typename T >
__cuda_callable__ inline
auto floor( const T& value ) -> decltype( std::floor(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::floor( value );
#else
return std::floor( value );
@@ -439,7 +431,7 @@ template< typename T >
__cuda_callable__ inline
auto ceil( const T& value ) -> decltype( std::ceil(value) )
{
-#if defined(__CUDA_ARCH__) || defined(__MIC__)
+#if defined(__CUDA_ARCH__)
return ::ceil( value );
#else
return std::ceil( value );
diff --git a/src/TNL/Matrices/CSR_impl.h b/src/TNL/Matrices/CSR_impl.h
index 74ff682fdaaa94d0b32f8b02375f0b9678f21307..cddf6f9a75adcb585ad777381ebafa386f6a3c92 100644
--- a/src/TNL/Matrices/CSR_impl.h
+++ b/src/TNL/Matrices/CSR_impl.h
@@ -831,38 +831,6 @@ class CSRDeviceDependentCode< Devices::Host >
};
-#ifdef HAVE_MIC
-template<>
-class CSRDeviceDependentCode< Devices::MIC >
-{
- public:
-
- typedef Devices::MIC Device;
-
- template< typename Real,
- typename Index,
- typename InVector,
- typename OutVector >
- static void vectorProduct( const CSR< Real, Device, Index >& matrix,
- const InVector& inVector,
- OutVector& outVector )
- {
- throw Exceptions::NotImplementedError("CSRDeviceDependentCode is not implemented for MIC.");
- }
- /* const Index rows = matrix.getRows();
- const tnlCSRMatrix< Real, Device, Index >* matrixPtr = &matrix;
- const InVector* inVectorPtr = &inVector;
- OutVector* outVectorPtr = &outVector;
-#ifdef HAVE_OPENMP
-#pragma omp parallel for firstprivate( matrixPtr, inVectorPtr, outVectorPtr ), schedule(static ), if( Devices::Host::isOMPEnabled() )
-#endif
- for( Index row = 0; row < rows; row ++ )
- ( *outVectorPtr )[ row ] = matrixPtr->rowVectorProduct( row, *inVectorPtr );
- }*/
-
-};
-#endif
-
#ifdef HAVE_CUDA
template< typename Real,
typename Index,
diff --git a/src/TNL/Matrices/SlicedEllpack_impl.h b/src/TNL/Matrices/SlicedEllpack_impl.h
index 016edf6996e19fb5fa3cfedf65364ebc22fb53fd..4ce70d3ef134a05060e0a2e496c31617635692c8 100644
--- a/src/TNL/Matrices/SlicedEllpack_impl.h
+++ b/src/TNL/Matrices/SlicedEllpack_impl.h
@@ -638,9 +638,9 @@ template< typename Real,
SlicedEllpack< Real, Device, Index, SliceSize >&
SlicedEllpack< Real, Device, Index, SliceSize >::operator=( const SlicedEllpack< Real2, Device2, Index2, SliceSize >& matrix )
{
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::Cuda >::value || std::is_same< Device, Devices::MIC >::value,
+ static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::Cuda >::value,
"unknown device" );
- static_assert( std::is_same< Device2, Devices::Host >::value || std::is_same< Device2, Devices::Cuda >::value || std::is_same< Device2, Devices::MIC >::value,
+ static_assert( std::is_same< Device2, Devices::Host >::value || std::is_same< Device2, Devices::Cuda >::value,
"unknown device" );
this->setLike( matrix );
@@ -693,10 +693,6 @@ SlicedEllpack< Real, Device, Index, SliceSize >::operator=( const SlicedEllpack<
}
}
}
-
- if( std::is_same< Device, Devices::MIC >::value ) {
- throw Exceptions::NotImplementedError("Cross-device assignment for the SlicedEllpack format is not implemented for MIC.");
- }
return *this;
}
@@ -1064,61 +1060,6 @@ class SlicedEllpackDeviceDependentCode< Devices::Cuda >
cudaDeviceSynchronize();
#endif
}
-
-};
-
-template<>
-class SlicedEllpackDeviceDependentCode< Devices::MIC >
-{
- public:
-
- typedef Devices::MIC Device;
-
- template< typename Real,
- typename Index,
- int SliceSize >
- static void initRowTraverse( const SlicedEllpack< Real, Device, Index, SliceSize >& matrix,
- const Index row,
- Index& rowBegin,
- Index& rowEnd,
- Index& step )
- {
- throw Exceptions::NotImplementedError("Not Implemented yet SlicedEllpackDeviceDependentCode< Devices::MIC >::initRowTraverse");
- }
-
- template< typename Real,
- typename Index,
- int SliceSize >
- __cuda_callable__
- static void initRowTraverseFast( const SlicedEllpack< Real, Device, Index, SliceSize >& matrix,
- const Index row,
- Index& rowBegin,
- Index& rowEnd,
- Index& step )
- {
- throw Exceptions::NotImplementedError("Not Implemented yet SlicedEllpackDeviceDependentCode< Devices::MIC >::initRowTraverseFast");
- }
-
- template< typename Real,
- typename Index,
- int SliceSize >
- static bool computeMaximalRowLengthInSlices( SlicedEllpack< Real, Device, Index, SliceSize >& matrix,
- typename SlicedEllpack< Real, Device, Index >::ConstCompressedRowLengthsVectorView rowLengths )
- {
- throw Exceptions::NotImplementedError("Not Implemented yet SlicedEllpackDeviceDependentCode< Devices::MIC >::computeMaximalRowLengthInSlices");
- }
-
- template< typename Real,
- typename Index,
- typename InVector,
- typename OutVector,
- int SliceSize >
- static void vectorProduct( const SlicedEllpack< Real, Device, Index, SliceSize >& matrix,
- const InVector& inVector,
- OutVector& outVector )
- {
- throw Exceptions::NotImplementedError("Not Implemented yet SlicedEllpackDeviceDependentCode< Devices::MIC >::vectorProduct");
- }
};
} // namespace Matrices
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser.h b/src/TNL/Meshes/GridDetails/GridTraverser.h
index fb6b34da12fb750c0ad74cc3ba05b086727adf01..7ce106f5d9c9e05160490d59434b7d170298a993 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser.h
+++ b/src/TNL/Meshes/GridDetails/GridTraverser.h
@@ -89,38 +89,6 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::Cuda, Index > >
const int& stream = 0 );
};
-/****
- * 1D grid, Devices::MIC
- */
-template< typename Real,
- typename Index >
-class GridTraverser< Meshes::Grid< 1, Real, Devices::MIC, Index > >
-{
- public:
-
- typedef Meshes::Grid< 1, Real, Devices::MIC, Index > GridType;
- typedef Pointers::SharedPointer< GridType > GridPointer;
- typedef Real RealType;
- typedef Devices::MIC DeviceType;
- typedef Index IndexType;
- typedef typename GridType::CoordinatesType CoordinatesType;
-
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities >
- static void
- processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- GridTraverserMode mode = synchronousMode,
- const int& stream = 0 );
-};
-
-
/****
* 2D grid, Devices::Host
@@ -202,45 +170,6 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::Cuda, Index > >
const GridEntityParameters&... gridEntityParameters );
};
-/****
- * 2D grid, Devices::MIC
- */
-template< typename Real,
- typename Index >
-class GridTraverser< Meshes::Grid< 2, Real, Devices::MIC, Index > >
-{
- public:
-
- typedef Meshes::Grid< 2, Real, Devices::MIC, Index > GridType;
- typedef Pointers::SharedPointer< GridType > GridPointer;
- typedef Real RealType;
- typedef Devices::MIC DeviceType;
- typedef Index IndexType;
- typedef typename GridType::CoordinatesType CoordinatesType;
-
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities,
- int XOrthogonalBoundary = 1,
- int YOrthogonalBoundary = 1,
- typename... GridEntityParameters >
- static void
- processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- // FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
- //GridTraverserMode mode = synchronousMode,
- GridTraverserMode mode,
- // const int& stream = 0,
- const int& stream,
- // gridEntityParameters are passed to GridEntity's constructor
- // (i.e. orientation and basis for faces)
- const GridEntityParameters&... gridEntityParameters );
-};
/****
* 3D grid, Devices::Host
@@ -324,51 +253,9 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::Cuda, Index > >
const GridEntityParameters&... gridEntityParameters );
};
-/****
- * 3D grid, Devices::Cuda
- */
-template< typename Real,
- typename Index >
-class GridTraverser< Meshes::Grid< 3, Real, Devices::MIC, Index > >
-{
- public:
-
- typedef Meshes::Grid< 3, Real, Devices::MIC, Index > GridType;
- typedef Pointers::SharedPointer< GridType > GridPointer;
- typedef Real RealType;
- typedef Devices::MIC DeviceType;
- typedef Index IndexType;
- typedef typename GridType::CoordinatesType CoordinatesType;
-
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities,
- int XOrthogonalBoundary = 1,
- int YOrthogonalBoundary = 1,
- int ZOrthogonalBoundary = 1,
- typename... GridEntityParameters >
- static void
- processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- // FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
- //GridTraverserMode mode = synchronousMode,
- GridTraverserMode mode,
- // const int& stream = 0,
- const int& stream,
- // gridEntityParameters are passed to GridEntity's constructor
- // (i.e. orientation and basis for faces and edges)
- const GridEntityParameters&... gridEntityParameters );
-};
-
} // namespace Meshes
} // namespace TNL
#include <TNL/Meshes/GridDetails/GridTraverser_1D.hpp>
#include <TNL/Meshes/GridDetails/GridTraverser_2D.hpp>
#include <TNL/Meshes/GridDetails/GridTraverser_3D.hpp>
-
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
index 59989bb2a14a85443c2f9616c583ab945b727116..53370853824a8fde7f47bf28c0a68de76f84a2b0 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
@@ -14,7 +14,6 @@
#pragma once
-#include <TNL/Devices/MIC.h>
#include <TNL/Meshes/Grid.h>
#include <TNL/Pointers/SharedPointer.h>
#include <TNL/CudaStreamPool.h>
@@ -255,69 +254,5 @@ processEntities(
#endif
}
-/****
- * 1D traverser, MIC
- */
-
-template< typename Real,
- typename Index >
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities >
-void
-GridTraverser< Meshes::Grid< 1, Real, Devices::MIC, Index > >::
-processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- GridTraverserMode mode,
- const int& stream )
-{
- throw Exceptions::NotImplementedError("Not Implemented yet Grid Traverser <1, Real, Device::MIC>");
-/*
- auto& pool = CudaStreamPool::getInstance();
- const cudaStream_t& s = pool.getStream( stream );
-
- Devices::Cuda::synchronizeDevice();
- if( processOnlyBoundaryEntities )
- {
- dim3 cudaBlockSize( 2 );
- dim3 cudaBlocks( 1 );
- GridBoundaryTraverser1D< Real, Index, GridEntity, UserData, EntitiesProcessor >
- <<< cudaBlocks, cudaBlockSize, 0, s >>>
- ( &gridPointer.template getData< Devices::Cuda >(),
- userData,
- begin,
- end );
- }
- else
- {
- dim3 cudaBlockSize( 256 );
- dim3 cudaBlocks;
- cudaBlocks.x = Devices::Cuda::getNumberOfBlocks( end.x() - begin.x() + 1, cudaBlockSize.x );
- const IndexType cudaXGrids = Devices::Cuda::getNumberOfGrids( cudaBlocks.x );
-
- for( IndexType gridXIdx = 0; gridXIdx < cudaXGrids; gridXIdx ++ )
- GridTraverser1D< Real, Index, GridEntity, UserData, EntitiesProcessor >
- <<< cudaBlocks, cudaBlockSize, 0, s >>>
- ( &gridPointer.template getData< Devices::Cuda >(),
- userData,
- begin,
- end,
- gridXIdx );
- }
-
- // only launches into the stream 0 are synchronized
- if( stream == 0 )
- {
- cudaStreamSynchronize( s );
- TNL_CHECK_CUDA_DEVICE;
- }
-*/
-}
-
- } // namespace Meshes
+} // namespace Meshes
} // namespace TNL
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
index 50b30c0190bdda8c6c266385ecd785884f3282ac..3efdb478fd9d8874a730d48deff8849733f0b13e 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
@@ -10,7 +10,6 @@
#pragma once
-#include <TNL/Devices/MIC.h>
#include <TNL/Meshes/Grid.h>
#include <TNL/Pointers/SharedPointer.h>
#include <TNL/CudaStreamPool.h>
@@ -553,104 +552,5 @@ processEntities(
#endif
}
-
-/****
- * 2D traverser, MIC
- */
-template< typename Real,
- typename Index >
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities,
- int XOrthogonalBoundary,
- int YOrthogonalBoundary,
- typename... GridEntityParameters >
-void
-GridTraverser< Meshes::Grid< 2, Real, Devices::MIC, Index > >::
-processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- GridTraverserMode mode,
- const int& stream,
- const GridEntityParameters&... gridEntityParameters )
-{
-
-
-#ifdef HAVE_MIC
- Devices::MIC::synchronizeDevice();
-
- //TOHLE JE PRUSER -- nemim poslat vypustku --
- //GridEntity entity( gridPointer.template getData< Devices::MIC >(), begin, gridEntityParameters... );
-
-
- Devices::MICHider<const GridType> hMicGrid;
- hMicGrid.pointer=& gridPointer.template getData< Devices::MIC >();
- Devices::MICHider<UserData> hMicUserData;
- hMicUserData.pointer=& userDataPointer.template modifyData<Devices::MIC>();
- TNLMICSTRUCT(begin, const CoordinatesType);
- TNLMICSTRUCT(end, const CoordinatesType);
-
- #pragma offload target(mic) in(sbegin,send,hMicUserData,hMicGrid)
- {
-
- #pragma omp parallel firstprivate( sbegin, send )
- {
- TNLMICSTRUCTUSE(begin, const CoordinatesType);
- TNLMICSTRUCTUSE(end, const CoordinatesType);
- GridEntity entity( *(hMicGrid.pointer), *(kernelbegin) );
-
- if( processOnlyBoundaryEntities )
- {
- if( YOrthogonalBoundary )
- #pragma omp for
- for( auto k = kernelbegin->x();
- k <= kernelend->x();
- k ++ )
- {
- entity.getCoordinates().x() = k;
- entity.getCoordinates().y() = kernelbegin->y();
- entity.refresh();
- EntitiesProcessor::processEntity( entity.getMesh(), *(hMicUserData.pointer), entity );
- entity.getCoordinates().y() = kernelend->y();
- entity.refresh();
- EntitiesProcessor::processEntity( entity.getMesh(), *(hMicUserData.pointer), entity );
- }
- if( XOrthogonalBoundary )
- #pragma omp for
- for( auto k = kernelbegin->y();
- k <= kernelend->y();
- k ++ )
- {
- entity.getCoordinates().y() = k;
- entity.getCoordinates().x() = kernelbegin->x();
- entity.refresh();
- EntitiesProcessor::processEntity( entity.getMesh(), *(hMicUserData.pointer), entity );
- entity.getCoordinates().x() = kernelend->x();
- entity.refresh();
- EntitiesProcessor::processEntity( entity.getMesh(), *(hMicUserData.pointer), entity );
- }
- }
- else
- {
- #pragma omp for
- for( IndexType y = kernelbegin->y(); y <= kernelend->y(); y ++ )
- for( IndexType x = kernelbegin->x(); x <= kernelend->x(); x ++ )
- {
- // std::cerr << x << " " <<y << std::endl;
- entity.getCoordinates().x() = x;
- entity.getCoordinates().y() = y;
- entity.refresh();
- EntitiesProcessor::processEntity( entity.getMesh(), *(hMicUserData.pointer), entity );
- }
- }
- }
- }
-
-#endif
-}
- } // namespace Meshes
+} // namespace Meshes
} // namespace TNL
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
index 5a3cd568f93bcb20f40682e55959eadf50b5c67f..24200c15dd89661492142c5f74f3166feb5d7ed6 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
@@ -10,7 +10,6 @@
#pragma once
-#include <TNL/Devices/MIC.h>
#include <TNL/Meshes/Grid.h>
#include <TNL/Pointers/SharedPointer.h>
#include <TNL/CudaStreamPool.h>
@@ -488,68 +487,5 @@ processEntities(
#endif
}
-/****
- * 3D traverser, MIC
- */
-template< typename Real,
- typename Index >
- template<
- typename GridEntity,
- typename EntitiesProcessor,
- typename UserData,
- bool processOnlyBoundaryEntities,
- int XOrthogonalBoundary,
- int YOrthogonalBoundary,
- int ZOrthogonalBoundary,
- typename... GridEntityParameters >
-void
-GridTraverser< Meshes::Grid< 3, Real, Devices::MIC, Index > >::
-processEntities(
- const GridPointer& gridPointer,
- const CoordinatesType& begin,
- const CoordinatesType& end,
- UserData& userData,
- GridTraverserMode mode,
- const int& stream,
- const GridEntityParameters&... gridEntityParameters )
-{
- throw Exceptions::NotImplementedError("Not Implemented yet Grid Traverser <3, Real, Device::MIC>");
-
-/* HAVE_CUDA
- dim3 cudaBlockSize( 8, 8, 8 );
- dim3 cudaBlocks;
- cudaBlocks.x = Devices::Cuda::getNumberOfBlocks( end.x() - begin.x() + 1, cudaBlockSize.x );
- cudaBlocks.y = Devices::Cuda::getNumberOfBlocks( end.y() - begin.y() + 1, cudaBlockSize.y );
- cudaBlocks.z = Devices::Cuda::getNumberOfBlocks( end.z() - begin.z() + 1, cudaBlockSize.z );
- const IndexType cudaXGrids = Devices::Cuda::getNumberOfGrids( cudaBlocks.x );
- const IndexType cudaYGrids = Devices::Cuda::getNumberOfGrids( cudaBlocks.y );
- const IndexType cudaZGrids = Devices::Cuda::getNumberOfGrids( cudaBlocks.z );
-
- auto& pool = CudaStreamPool::getInstance();
- const cudaStream_t& s = pool.getStream( stream );
-
- Devices::Cuda::synchronizeDevice();
- for( IndexType gridZIdx = 0; gridZIdx < cudaZGrids; gridZIdx ++ )
- for( IndexType gridYIdx = 0; gridYIdx < cudaYGrids; gridYIdx ++ )
- for( IndexType gridXIdx = 0; gridXIdx < cudaXGrids; gridXIdx ++ )
- GridTraverser3D< Real, Index, GridEntity, UserData, EntitiesProcessor, processOnlyBoundaryEntities, GridEntityParameters... >
- <<< cudaBlocks, cudaBlockSize, 0, s >>>
- ( &gridPointer.template getData< Devices::Cuda >(),
- userData,
- begin,
- end,
- gridXIdx,
- gridYIdx,
- gridZIdx,
- gridEntityParameters... );
-
- // only launches into the stream 0 are synchronized
- if( stream == 0 )
- {
- cudaStreamSynchronize( s );
- TNL_CHECK_CUDA_DEVICE;
- }
- */
-}
- } // namespace Meshes
+} // namespace Meshes
} // namespace TNL
diff --git a/src/TNL/Object.h b/src/TNL/Object.h
index 24ced9a5c2ee9ac97f8d2a7a86e51e88085d13a4..356b91eda2f3b71381c14e4f9ff4933976bf33df 100644
--- a/src/TNL/Object.h
+++ b/src/TNL/Object.h
@@ -128,9 +128,7 @@ class Object
* Since it is not defined as \ref __cuda_callable__, objects inherited
* from Object should not be created in CUDA kernels.
*/
-#ifndef HAVE_MIC
virtual ~Object(){};
-#endif
};
/**
diff --git a/src/TNL/Pointers/DevicePointer.h b/src/TNL/Pointers/DevicePointer.h
index b0c0a934fa0dee01ebe4bb2d93abe3e6d0d36b68..7c0982dcaaddd6cd6deaddb6be2b28727ef3c40b 100644
--- a/src/TNL/Pointers/DevicePointer.h
+++ b/src/TNL/Pointers/DevicePointer.h
@@ -15,7 +15,6 @@
#include <TNL/Allocators/Default.h>
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
#include <TNL/Pointers/SmartPointer.h>
#include <cstring> // std::memcpy, std::memcmp
@@ -470,288 +469,9 @@ class DevicePointer< Object, Devices::Cuda > : public SmartPointer
Object* cuda_pointer;
};
-/****
- * Specialization for MIC
- */
-
-#ifdef HAVE_MIC
-template< typename Object >
-class DevicePointer< Object, Devices::MIC > : public SmartPointer
-{
- private:
- // Convenient template alias for controlling the selection of copy- and
- // move-constructors and assignment operators using SFINAE.
- // The type Object_ is "enabled" iff Object_ and Object are not the same,
- // but after removing const and volatile qualifiers they are the same.
- template< typename Object_ >
- using Enabler = std::enable_if< ! std::is_same< Object_, Object >::value &&
- std::is_same< typename std::remove_cv< Object >::type, Object_ >::value >;
-
- // friend class will be needed for templated assignment operators
- template< typename Object_, typename Device_ >
- friend class DevicePointer;
-
- public:
-
- typedef Object ObjectType;
- typedef Devices::MIC DeviceType;
-
- explicit DevicePointer( ObjectType& obj )
- : pointer( nullptr ),
- pd( nullptr ),
- mic_pointer( nullptr )
- {
- this->allocate( obj );
- }
-
- // this is needed only to avoid the default compiler-generated constructor
- DevicePointer( const DevicePointer& pointer )
- : pointer( pointer.pointer ),
- pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- this->pd->counter += 1;
- }
-
- // conditional constructor for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- DevicePointer( const DevicePointer< Object_, DeviceType >& pointer )
- : pointer( pointer.pointer ),
- pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- this->pd->counter += 1;
- }
-
- // this is needed only to avoid the default compiler-generated constructor
- DevicePointer( DevicePointer&& pointer )
- : pointer( pointer.pointer ),
- pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- pointer.pointer = nullptr;
- pointer.pd = nullptr;
- pointer.mic_pointer = nullptr;
- }
-
- // conditional constructor for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- DevicePointer( DevicePointer< Object_, DeviceType >&& pointer )
- : pointer( pointer.pointer ),
- pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- pointer.pointer = nullptr;
- pointer.pd = nullptr;
- pointer.mic_pointer = nullptr;
- }
-
- const Object* operator->() const
- {
- return this->pointer;
- }
-
- Object* operator->()
- {
- this->pd->maybe_modified = true;
- return this->pointer;
- }
-
- const Object& operator *() const
- {
- return *( this->pointer );
- }
-
- Object& operator *()
- {
- this->pd->maybe_modified = true;
- return *( this->pointer );
- }
-
- operator bool()
- {
- return this->pd;
- }
-
- template< typename Device = Devices::Host >
- __cuda_callable__
- const Object& getData() const
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT( this->pointer, );
- TNL_ASSERT( this->pd, );
- TNL_ASSERT( this->mic_pointer, );
- if( std::is_same< Device, Devices::Host >::value )
- return *( this->pointer );
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
- }
-
- template< typename Device = Devices::Host >
- __cuda_callable__
- Object& modifyData()
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT( this->pointer, );
- TNL_ASSERT( this->pd, );
- TNL_ASSERT( this->mic_pointer, );
- if( std::is_same< Device, Devices::Host >::value )
- {
- this->pd->maybe_modified = true;
- return *( this->pointer );
- }
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
- }
-
- // this is needed only to avoid the default compiler-generated operator
- const DevicePointer& operator=( const DevicePointer& ptr )
- {
- this->free();
- this->pointer = ptr.pointer;
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- this->pd->counter += 1;
- return *this;
- }
-
- // conditional operator for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- const DevicePointer& operator=( const DevicePointer< Object_, DeviceType >& ptr )
- {
- this->free();
- this->pointer = ptr.pointer;
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- this->pd->counter += 1;
- return *this;
- }
-
- // this is needed only to avoid the default compiler-generated operator
- const DevicePointer& operator=( DevicePointer&& ptr )
- {
- this->free();
- this->pointer = ptr.pointer;
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- ptr.pointer = nullptr;
- ptr.pd = nullptr;
- ptr.mic_pointer = nullptr;
- return *this;
- }
-
- // conditional operator for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- const DevicePointer& operator=( DevicePointer< Object_, DeviceType >&& ptr )
- {
- this->free();
- this->pointer = ptr.pointer;
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- ptr.pointer = nullptr;
- ptr.pd = nullptr;
- ptr.mic_pointer = nullptr;
- return *this;
- }
-
- bool synchronize()
- {
- if( ! this->pd )
- return true;
- if( this->modified() )
- {
- TNL_ASSERT( this->pointer, );
- TNL_ASSERT( this->mic_pointer, );
- Devices::MIC::CopyToMIC((void*) this->mic_pointer, (void*) this->pointer, sizeof( ObjectType ));
- this->set_last_sync_state();
- return true;
- }
- return true;
-
- }
-
- ~DevicePointer()
- {
- this->free();
- Devices::MIC::removeSmartPointer( this );
- }
-
- protected:
-
- struct PointerData
- {
- char data_image[ sizeof(Object) ];
- int counter = 1;
- bool maybe_modified = false;
- };
-
- bool allocate( ObjectType& obj )
- {
- this->pointer = &obj;
- this->pd = new PointerData();
- if( ! this->pd )
- return false;
- // pass to device
- this->mic_pointer = Allocators:::MIC< ObjectType >().allocate(1);
- if( ! this->mic_pointer )
- return false;
- Devices::MIC::CopyToMIC((void*)this->mic_pointer,(void*)this->pointer,sizeof(ObjectType));
-
- // set last-sync state
- this->set_last_sync_state();
- Devices::MIC::insertSmartPointer( this );
- return true;
- }
-
- void set_last_sync_state()
- {
- TNL_ASSERT( this->pointer, );
- TNL_ASSERT( this->pd, );
- std::memcpy( (void*) &this->pd->data_image, (void*) this->pointer, sizeof( Object ) );
- this->pd->maybe_modified = false;
- }
-
- bool modified()
- {
- TNL_ASSERT( this->pointer, );
- TNL_ASSERT( this->pd, );
- // optimization: skip bitwise comparison if we're sure that the data is the same
- if( ! this->pd->maybe_modified )
- return false;
- return std::memcmp( (void*) &this->pd->data_image, (void*) this->pointer, sizeof( Object ) ) != 0;
- }
-
- void free()
- {
- if( this->pd )
- {
- if( ! --this->pd->counter )
- {
- delete this->pd;
- this->pd = nullptr;
- if( this->mic_pointer )
- Allocators:::MIC< ObjectType >().deallocate(this->mic_pointer, 1);
- }
- }
- }
-
- Object* pointer;
-
- PointerData* pd;
-
- // mic_pointer can't be part of PointerData structure, since we would be
- // unable to dereference this-pd on the device
- Object* mic_pointer;
-};
-#endif
-
} // namespace Pointers
-#if (!defined(NDEBUG)) && (!defined(HAVE_MIC))
+#ifndef NDEBUG
namespace Assert {
template< typename Object, typename Device >
diff --git a/src/TNL/Pointers/SharedPointer.h b/src/TNL/Pointers/SharedPointer.h
index e6908e47953b330b612ea9ec8a2421d8c11bc8a9..51aff2a7829e60d84116eae7b602020a6fc0ae6b 100644
--- a/src/TNL/Pointers/SharedPointer.h
+++ b/src/TNL/Pointers/SharedPointer.h
@@ -49,7 +49,7 @@ class SharedPointer
} // namespace Pointers
-#if (!defined(NDEBUG)) && (!defined(HAVE_MIC))
+#ifndef NDEBUG
namespace Assert {
template< typename Object, typename Device >
@@ -72,4 +72,3 @@ struct Formatter< Pointers::SharedPointer< Object, Device > >
#include <TNL/Pointers/SharedPointerHost.h>
#include <TNL/Pointers/SharedPointerCuda.h>
-#include <TNL/Pointers/SharedPointerMic.h>
diff --git a/src/TNL/Pointers/SharedPointerMic.h b/src/TNL/Pointers/SharedPointerMic.h
deleted file mode 100644
index 0c2958b4ad7c6552f58363c98dca5104908f04cc..0000000000000000000000000000000000000000
--- a/src/TNL/Pointers/SharedPointerMic.h
+++ /dev/null
@@ -1,373 +0,0 @@
-/***************************************************************************
- SharedPointerMic.h - description
- -------------------
- begin : Aug 22, 2018
- copyright : (C) 2018 by Tomas Oberhuber et al.
- email : tomas.oberhuber@fjfi.cvut.cz
- ***************************************************************************/
-
-/* See Copyright Notice in tnl/Copyright */
-
-// Implemented by: Tomas Oberhuber, Jakub Klinkovsky
-
-#pragma once
-
-#include "SharedPointer.h"
-
-#include <TNL/Allocators/Default.h>
-#include <TNL/Devices/MIC.h>
-#include <TNL/Pointers/SmartPointer.h>
-
-#include <cstring> // std::memcpy, std::memcmp
-#include <cstddef> // std::nullptr_t
-#include <algorithm> // swap
-
-namespace TNL {
-namespace Pointers {
-
-#ifdef HAVE_MIC
-template< typename Object>
-class SharedPointer< Object, Devices::MIC > : public SmartPointer
-{
- private:
- // Convenient template alias for controlling the selection of copy- and
- // move-constructors and assignment operators using SFINAE.
- // The type Object_ is "enabled" iff Object_ and Object are not the same,
- // but after removing const and volatile qualifiers they are the same.
- template< typename Object_ >
- using Enabler = std::enable_if< ! std::is_same< Object_, Object >::value &&
- std::is_same< typename std::remove_cv< Object >::type, Object_ >::value >;
-
- // friend class will be needed for templated assignment operators
- template< typename Object_, typename Device_>
- friend class SharedPointer;
-
- public:
-
- using ObjectType = Object;
- using DeviceType = Devices::MIC;
-
- SharedPointer( std::nullptr_t )
- : pd( nullptr ),
- mic_pointer( nullptr )
- {}
-
- template< typename... Args >
- explicit SharedPointer( Args... args )
- : pd( nullptr ),
- mic_pointer( nullptr )
- {
- this->allocate( args... );
- }
-
- // this is needed only to avoid the default compiler-generated constructor
- SharedPointer( const SharedPointer& pointer )
- : pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- this->pd->counter += 1;
- }
-
- // conditional constructor for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- SharedPointer( const SharedPointer< Object_, DeviceType >& pointer )
- : pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- this->pd->counter += 1;
- }
-
- // this is needed only to avoid the default compiler-generated constructor
- SharedPointer( SharedPointer&& pointer )
- : pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- pointer.pd = nullptr;
- pointer.mic_pointer = nullptr;
- }
-
- // conditional constructor for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- SharedPointer( SharedPointer< Object_, DeviceType >&& pointer )
- : pd( (PointerData*) pointer.pd ),
- mic_pointer( pointer.mic_pointer )
- {
- pointer.pd = nullptr;
- pointer.mic_pointer = nullptr;
- }
-
- template< typename... Args >
- bool recreate( Args... args )
- {
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Recreating shared pointer to " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- if( ! this->pd )
- return this->allocate( args... );
-
- if( this->pd->counter == 1 )
- {
- /****
- * The object is not shared -> recreate it in-place, without reallocation
- */
- this->pd->data.~Object();
- new ( &this->pd->data ) Object( args... );
- Devices::MIC::CopyToMIC(this->mic_pointer,(void*) &this->pd->data,sizeof(Object));
- this->set_last_sync_state();
- return true;
- }
-
- // free will just decrement the counter
- this->free();
-
- return this->allocate( args... );
- }
-
- const Object* operator->() const
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- return &this->pd->data;
- }
-
- Object* operator->()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- this->pd->maybe_modified = true;
- return &this->pd->data;
- }
-
- const Object& operator *() const
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- return this->pd->data;
- }
-
- Object& operator *()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- this->pd->maybe_modified = true;
- return this->pd->data;
- }
-
- operator bool()
- {
- return this->pd;
- }
-
- template< typename Device = Devices::Host >
- __cuda_callable__
- const Object& getData() const
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- TNL_ASSERT_TRUE( this->mic_pointer, "Attempt to dereference a null pointer" );
- if( std::is_same< Device, Devices::Host >::value )
- return this->pd->data;
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
-
- }
-
- template< typename Device = Devices::Host >
- __cuda_callable__
- Object& modifyData()
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- TNL_ASSERT_TRUE( this->mic_pointer, "Attempt to dereference a null pointer" );
- if( std::is_same< Device, Devices::Host >::value )
- {
- this->pd->maybe_modified = true;
- return this->pd->data;
- }
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
-
- }
-
- // this is needed only to avoid the default compiler-generated operator
- const SharedPointer& operator=( const SharedPointer& ptr )
- {
- this->free();
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- if( this->pd != nullptr )
- this->pd->counter += 1;
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Copy-assigned shared pointer: counter = " << this->pd->counter << ", type: " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- return *this;
- }
-
- // conditional operator for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- const SharedPointer& operator=( const SharedPointer< Object_, DeviceType >& ptr )
- {
- this->free();
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- if( this->pd != nullptr )
- this->pd->counter += 1;
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Copy-assigned shared pointer: counter = " << this->pd->counter << ", type: " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- return *this;
- }
-
- // this is needed only to avoid the default compiler-generated operator
- const SharedPointer& operator=( SharedPointer&& ptr )
- {
- this->free();
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- ptr.pd = nullptr;
- ptr.mic_pointer = nullptr;
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Move-assigned shared pointer: counter = " << this->pd->counter << ", type: " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- return *this;
- }
-
- // conditional operator for non-const -> const data
- template< typename Object_,
- typename = typename Enabler< Object_ >::type >
- const SharedPointer& operator=( SharedPointer< Object_, DeviceType >&& ptr )
- {
- this->free();
- this->pd = (PointerData*) ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- ptr.pd = nullptr;
- ptr.mic_pointer = nullptr;
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Move-assigned shared pointer: counter = " << this->pd->counter << ", type: " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- return *this;
- }
-
- bool synchronize()
- {
- if( ! this->pd )
- return true;
-
- if( this->modified() )
- {
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Synchronizing shared pointer: counter = " << this->pd->counter << ", type: " << demangle(typeid(Object).name()) << std::endl;
- std::cerr << " ( " << sizeof( Object ) << " bytes, MIC adress " << this->mic_pointer << " )" << std::endl;
-#endif
- TNL_ASSERT( this->mic_pointer, );
-
- Devices::MIC::CopyToMIC((void*)this->mic_pointer,(void*) &this->pd->data,sizeof(Object));
- this->set_last_sync_state();
- return true;
- }
- return false; //??
- }
-
- void clear()
- {
- this->free();
- }
-
- void swap( SharedPointer& ptr2 )
- {
- std::swap( this->pd, ptr2.pd );
- std::swap( this->mic_pointer, ptr2.mic_pointer );
- }
-
- ~SharedPointer()
- {
- this->free();
- Devices::MIC::removeSmartPointer( this );
- }
-
- protected:
-
- struct PointerData
- {
- Object data;
- uint8_t data_image[ sizeof(Object) ];
- int counter;
- bool maybe_modified;
-
- template< typename... Args >
- explicit PointerData( Args... args )
- : data( args... ),
- counter( 1 ),
- maybe_modified( false )
- {}
- };
-
- template< typename... Args >
- bool allocate( Args... args )
- {
- this->pd = new PointerData( args... );
- if( ! this->pd )
- return false;
-
- mic_pointer = Allocators::MIC< Object >().allocate(1);
- Devices::MIC::CopyToMIC((void*)this->mic_pointer,(void*) &this->pd->data,sizeof(Object));
-
- if( ! this->mic_pointer )
- return false;
- // set last-sync state
- this->set_last_sync_state();
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Created shared pointer to " << demangle(typeid(ObjectType).name()) << " (mic_pointer = " << this->mic_pointer << ")" << std::endl;
-#endif
- Devices::MIC::insertSmartPointer( this );
- return true;
- }
-
- void set_last_sync_state()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- std::memcpy( (void*) &this->pd->data_image, (void*) &this->pd->data, sizeof( Object ) );
- this->pd->maybe_modified = false;
- }
-
- bool modified()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- // optimization: skip bitwise comparison if we're sure that the data is the same
- if( ! this->pd->maybe_modified )
- return false;
- return std::memcmp( (void*) &this->pd->data_image, (void*) &this->pd->data, sizeof( Object ) ) != 0;
- }
-
- void free()
- {
- if( this->pd )
- {
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "Freeing shared pointer: counter = " << this->pd->counter << ", mic_pointer = " << this->mic_pointer << ", type: " << demangle(typeid(ObjectType).name()) << std::endl;
-#endif
- if( ! --this->pd->counter )
- {
- delete this->pd;
- this->pd = nullptr;
- if( this->mic_pointer )
- {
- Allocators:::MIC< ObjectType >().deallocate(mic_pointer, 1);
- mic_pointer=nullptr;
- }
-#ifdef TNL_DEBUG_SHARED_POINTERS
- std::cerr << "...deleted data." << std::endl;
-#endif
- }
- }
- }
-
- PointerData* pd;
-
- // cuda_pointer can't be part of PointerData structure, since we would be
- // unable to dereference this-pd on the device -- NevĂm zda to platĂ pro MIC, asi jo
- Object* mic_pointer;
-};
-#endif
-
-} // namespace Pointers
-} // namespace TNL
diff --git a/src/TNL/Pointers/UniquePointer.h b/src/TNL/Pointers/UniquePointer.h
index cfb7b543fc3e94858ad5c34d4bf8e8c0faf85462..e85e18d1800539b7141629e6cce4a6c219908eab 100644
--- a/src/TNL/Pointers/UniquePointer.h
+++ b/src/TNL/Pointers/UniquePointer.h
@@ -15,7 +15,6 @@
#include <TNL/Allocators/Default.h>
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
#include <TNL/Pointers/SmartPointer.h>
#include <cstring> // std::memcpy, std::memcmp
@@ -311,187 +310,9 @@ class UniquePointer< Object, Devices::Cuda > : public SmartPointer
Object* cuda_pointer;
};
-#ifdef HAVE_MIC
-template< typename Object >
-class UniquePointer< Object, Devices::MIC > : public SmartPointer
-{
- public:
-
- typedef Object ObjectType;
- typedef Devices::MIC DeviceType;
-
- UniquePointer( std::nullptr_t )
- : pd( nullptr ),
- mic_pointer( nullptr )
- {}
-
- template< typename... Args >
- explicit UniquePointer( const Args... args )
- : pd( nullptr ),
- mic_pointer( nullptr )
- {
- this->allocate( args... );
- }
-
- const Object* operator->() const
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- return &this->pd->data;
- }
-
- Object* operator->()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- this->pd->maybe_modified = true;
- return &this->pd->data;
- }
-
- const Object& operator *() const
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- return this->pd->data;
- }
-
- Object& operator *()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- this->pd->maybe_modified = true;
- return this->pd->data;
- }
-
- operator bool()
- {
- return this->pd;
- }
-
- template< typename Device = Devices::Host >
- const Object& getData() const
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- TNL_ASSERT_TRUE( this->mic_pointer, "Attempt to dereference a null pointer" );
- if( std::is_same< Device, Devices::Host >::value )
- return this->pd->data;
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
- }
-
- template< typename Device = Devices::Host >
- Object& modifyData()
- {
- static_assert( std::is_same< Device, Devices::Host >::value || std::is_same< Device, Devices::MIC >::value, "Only Devices::Host or Devices::MIC devices are accepted here." );
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- TNL_ASSERT_TRUE( this->mic_pointer, "Attempt to dereference a null pointer" );
- if( std::is_same< Device, Devices::Host >::value )
- {
- this->pd->maybe_modified = true;
- return this->pd->data;
- }
- if( std::is_same< Device, Devices::MIC >::value )
- return *( this->mic_pointer );
- }
-
- const UniquePointer& operator=( UniquePointer& ptr )
- {
- this->free();
- this->pd = ptr.pd;
- this->mic_pointer = ptr.mic_pointer;
- ptr.pd = nullptr;
- ptr.mic_pointer = nullptr;
- return *this;
- }
-
- const UniquePointer& operator=( UniquePointer&& ptr )
- {
- return this->operator=( ptr );
- }
-
- bool synchronize()
- {
- if( ! this->pd )
- return true;
- if( this->modified() )
- {
- Devices::MIC::CopyToMIC(this->mic_pointer,(void*) &this->pd->data,sizeof(Object));
- this->set_last_sync_state();
- return true;
- }
- return true;//??
- }
-
- ~UniquePointer()
- {
- this->free();
- Devices::MIC::removeSmartPointer( this );
- }
-
- protected:
-
- struct PointerData
- {
- Object data;
- char data_image[ sizeof(Object) ];
- bool maybe_modified;
-
- template< typename... Args >
- explicit PointerData( Args... args )
- : data( args... ),
- maybe_modified( false )
- {}
- };
-
- template< typename... Args >
- bool allocate( Args... args )
- {
- this->pd = new PointerData( args... );
- if( ! this->pd )
- return false;
- // pass to device
- this->mic_pointer = Allocators::MIC< Object >().allocate(1);
- if( ! this->mic_pointer )
- return false;
- Devices::MIC::CopyToMIC((void*)mic_pointer,(void*)&this->pd->data,sizeof(Object));
- // set last-sync state
- this->set_last_sync_state();
- Devices::MIC::insertSmartPointer( this );
- return true;
- }
-
- void set_last_sync_state()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- std::memcpy( (void*) &this->pd->data_image, (void*) &this->pd->data, sizeof( ObjectType ) );
- this->pd->maybe_modified = false;
- }
-
- bool modified()
- {
- TNL_ASSERT_TRUE( this->pd, "Attempt to dereference a null pointer" );
- // optimization: skip bitwise comparison if we're sure that the data is the same
- if( ! this->pd->maybe_modified )
- return false;
- return std::memcmp( (void*) &this->pd->data_image, (void*) &this->pd->data, sizeof( ObjectType ) ) != 0;
- }
-
- void free()
- {
- if( this->pd )
- delete this->pd;
- if( this->mic_pointer )
- Allocators:::MIC< ObjectType >().deallocate(mic_pointer, 1);
- }
-
- PointerData* pd;
-
- // mic_pointer can't be part of PointerData structure, since we would be
- // unable to dereference this-pd on the device
- Object* mic_pointer;
-};
-#endif
-
} // namespace Pointers
-#if (!defined(NDEBUG)) && (!defined(HAVE_MIC))
+#ifndef NDEBUG
namespace Assert {
template< typename Object, typename Device >
diff --git a/src/TNL/Solvers/BuildConfigTags.h b/src/TNL/Solvers/BuildConfigTags.h
index 19bb42129563b48a7e1e1ba9baf8f107d25ee661..bcd4cdafcacff729b51b827348bcd7703f4bec21 100644
--- a/src/TNL/Solvers/BuildConfigTags.h
+++ b/src/TNL/Solvers/BuildConfigTags.h
@@ -27,10 +27,6 @@ template< typename ConfigTag, typename Device > struct ConfigTagDevice{ enum { e
template< typename ConfigTag > struct ConfigTagDevice< ConfigTag, Devices::Cuda >{ enum { enabled = false }; };
#endif
-#ifndef HAVE_MIC
-template< typename ConfigTag > struct ConfigTagDevice< ConfigTag, Devices::MIC >{ enum { enabled = false }; };
-#endif
-
/****
* All real types are enabled by default.
*/
diff --git a/src/TNL/Solvers/Linear/Preconditioners/ILU0.h b/src/TNL/Solvers/Linear/Preconditioners/ILU0.h
index 1fc2fa3fa69f964cb3486d6ee16dcf43fc8d3b9f..97bc854ce7e53cc1b90558860dbea6d6555c148b 100644
--- a/src/TNL/Solvers/Linear/Preconditioners/ILU0.h
+++ b/src/TNL/Solvers/Linear/Preconditioners/ILU0.h
@@ -199,29 +199,6 @@ public:
}
};
-template< typename Matrix, typename Real, typename Index >
-class ILU0_impl< Matrix, Real, Devices::MIC, Index >
-: public Preconditioner< Matrix >
-{
-public:
- using RealType = Real;
- using DeviceType = Devices::MIC;
- using IndexType = Index;
- using typename Preconditioner< Matrix >::VectorViewType;
- using typename Preconditioner< Matrix >::ConstVectorViewType;
- using typename Preconditioner< Matrix >::MatrixPointer;
-
- virtual void update( const MatrixPointer& matrixPointer ) override
- {
- throw Exceptions::NotImplementedError("Not Iplemented yet for MIC");
- }
-
- virtual void solve( ConstVectorViewType b, VectorViewType x ) const override
- {
- throw Exceptions::NotImplementedError("Not Iplemented yet for MIC");
- }
-};
-
} // namespace Preconditioners
} // namespace Linear
} // namespace Solvers
diff --git a/src/TNL/Solvers/Linear/Preconditioners/ILUT.h b/src/TNL/Solvers/Linear/Preconditioners/ILUT.h
index 8f4c27d7abd8d65566916b7ac79f34d269bc84e1..fa7c814fc48a888e84c91b748e974bff18509452 100644
--- a/src/TNL/Solvers/Linear/Preconditioners/ILUT.h
+++ b/src/TNL/Solvers/Linear/Preconditioners/ILUT.h
@@ -111,29 +111,6 @@ public:
}
};
-template< typename Matrix, typename Real, typename Index >
-class ILUT_impl< Matrix, Real, Devices::MIC, Index >
-: public Preconditioner< Matrix >
-{
-public:
- using RealType = Real;
- using DeviceType = Devices::MIC;
- using IndexType = Index;
- using typename Preconditioner< Matrix >::VectorViewType;
- using typename Preconditioner< Matrix >::ConstVectorViewType;
- using typename Preconditioner< Matrix >::MatrixPointer;
-
- virtual void update( const MatrixPointer& matrixPointer ) override
- {
- throw std::runtime_error("Not Iplemented yet for MIC");
- }
-
- virtual void solve( ConstVectorViewType b, VectorViewType x ) const override
- {
- throw std::runtime_error("Not Iplemented yet for MIC");
- }
-};
-
} // namespace Preconditioners
} // namespace Linear
} // namespace Solvers
diff --git a/src/TNL/Solvers/ODE/Euler.h b/src/TNL/Solvers/ODE/Euler.h
index 2ba128073ec65aba4f8e1bc5c7f6cad661f67303..508d77b6a97045d1f9250cc61bc8beffe7a6656f 100644
--- a/src/TNL/Solvers/ODE/Euler.h
+++ b/src/TNL/Solvers/ODE/Euler.h
@@ -10,12 +10,10 @@
#pragma once
-#include <math.h>
#include <TNL/Config/ConfigDescription.h>
#include <TNL/Solvers/ODE/ExplicitSolver.h>
#include <TNL/Solvers/DummyProblem.h>
#include <TNL/Config/ParameterContainer.h>
-#include <TNL/Timer.h>
namespace TNL {
namespace Solvers {
diff --git a/src/TNL/Solvers/ODE/Euler.hpp b/src/TNL/Solvers/ODE/Euler.hpp
index 12da6439bd15d4fdbe1e0a088910c940cfc90aa2..1cf5001ae4397daab76e978af97a8cbb86c24df9 100644
--- a/src/TNL/Solvers/ODE/Euler.hpp
+++ b/src/TNL/Solvers/ODE/Euler.hpp
@@ -10,9 +10,7 @@
#pragma once
-#include <TNL/Devices/MIC.h>
-#include <TNL/Communicators/MpiCommunicator.h>
-#include <TNL/Communicators/NoDistrCommunicator.h>
+#include <TNL/Solvers/ODE/Euler.h>
namespace TNL {
namespace Solvers {
@@ -77,7 +75,6 @@ bool Euler< Problem, SolverMonitor > :: solve( DofVectorPointer& _u )
/****
* First setup the supporting meshes k1...k5 and k_tmp.
*/
- //timer.start();
_k1->setLike( *_u );
auto k1 = _k1->getView();
auto u = _u->getView();
@@ -104,9 +101,7 @@ bool Euler< Problem, SolverMonitor > :: solve( DofVectorPointer& _u )
/****
* Compute the RHS
*/
- //timer.stop();
this->problem->getExplicitUpdate( time, currentTau, _u, _k1 );
- //timer.start();
RealType lastResidue = this->getResidue();
RealType maxResidue( 0.0 );
diff --git a/src/TNL/Solvers/SolverConfig_impl.h b/src/TNL/Solvers/SolverConfig_impl.h
index 701c5eb730b99e2487e56bd5e56a9ffdec0b916d..e5673d5c1ed45ea9a28f8615cd4f099284bb8875 100644
--- a/src/TNL/Solvers/SolverConfig_impl.h
+++ b/src/TNL/Solvers/SolverConfig_impl.h
@@ -67,12 +67,6 @@ bool SolverConfig< ConfigTag, ProblemConfig >::configSetup( Config::ConfigDescri
if( ConfigTagDevice< ConfigTag, Devices::Cuda >::enabled )
config.addEntryEnum( "cuda" );
#endif
-
-#ifdef HAVE_MIC
- if( ConfigTagDevice< ConfigTag, Devices::MIC >::enabled )
- config.addEntryEnum( "mic" );
-#endif
-
/****
* Setup index type.
diff --git a/src/TNL/Solvers/SolverInitiator_impl.h b/src/TNL/Solvers/SolverInitiator_impl.h
index c6bc5ca7f494abd8922f1a0fcb45b4814277094f..e54a8fe308c4478a7242a24f8032473be8431d1e 100644
--- a/src/TNL/Solvers/SolverInitiator_impl.h
+++ b/src/TNL/Solvers/SolverInitiator_impl.h
@@ -12,7 +12,6 @@
#include <TNL/Devices/Host.h>
#include <TNL/Devices/Cuda.h>
-#include <TNL/Devices/MIC.h>
#include <TNL/Config/ParameterContainer.h>
#include <TNL/Meshes/TypeResolver/TypeResolver.h>
#include <TNL/Solvers/BuildConfigTags.h>
@@ -92,8 +91,6 @@ class SolverInitiatorRealResolver< ProblemSetter, Real, ConfigTag, true >
return SolverInitiatorDeviceResolver< ProblemSetter, Real, Devices::Host, ConfigTag >::run( parameters );
if( device == "cuda" )
return SolverInitiatorDeviceResolver< ProblemSetter, Real, Devices::Cuda, ConfigTag >::run( parameters );
- if(device == "mic")
- return SolverInitiatorDeviceResolver< ProblemSetter, Real, Devices::MIC, ConfigTag >::run( parameters );
std::cerr << "The device '" << device << "' is not defined. " << std::endl;
return false;
}
diff --git a/src/UnitTests/Containers/ArrayTest.h b/src/UnitTests/Containers/ArrayTest.h
index a18471a4a8c24447b48fe673627471aa28e639f6..7151ed441477f93a885b288b1845ca3f38527b93 100644
--- a/src/UnitTests/Containers/ArrayTest.h
+++ b/src/UnitTests/Containers/ArrayTest.h
@@ -101,27 +101,6 @@ using ArrayTypes = ::testing::Types<
Array< float, Devices::Cuda, long >,
Array< double, Devices::Cuda, long >,
Array< MyData, Devices::Cuda, long >
-#endif
-#ifdef HAVE_MIC
- ,
- Array< int, Devices::MIC, short >,
- Array< long, Devices::MIC, short >,
- Array< float, Devices::MIC, short >,
- Array< double, Devices::MIC, short >,
- // TODO: MyData does not work on MIC
-// Array< MyData, Devices::MIC, short >,
- Array< int, Devices::MIC, int >,
- Array< long, Devices::MIC, int >,
- Array< float, Devices::MIC, int >,
- Array< double, Devices::MIC, int >,
- // TODO: MyData does not work on MIC
-// Array< MyData, Devices::MIC, int >,
- Array< int, Devices::MIC, long >,
- Array< long, Devices::MIC, long >,
- Array< float, Devices::MIC, long >,
- Array< double, Devices::MIC, long >
- // TODO: MyData does not work on MIC
-// Array< MyData, Devices::MIC, long >
#endif
// all array tests should also work with Vector
@@ -136,11 +115,6 @@ using ArrayTypes = ::testing::Types<
Vector< float, Devices::Cuda, long >,
Vector< double, Devices::Cuda, long >
#endif
-#ifdef HAVE_MIC
- ,
- Vector< float, Devices::MIC, long >,
- Vector< double, Devices::MIC, long >
-#endif
>;
TYPED_TEST_SUITE( ArrayTest, ArrayTypes );
@@ -353,14 +327,6 @@ void testArrayElementwiseAccess( Array< Value, Devices::Cuda, Index >&& u )
#endif
}
-template< typename Value, typename Index >
-void testArrayElementwiseAccess( Array< Value, Devices::MIC, Index >&& u )
-{
-#ifdef HAVE_MIC
- // TODO
-#endif
-}
-
TYPED_TEST( ArrayTest, elementwiseAccess )
{
using ArrayType = typename TestFixture::ArrayType;
diff --git a/src/UnitTests/Containers/ArrayViewTest.h b/src/UnitTests/Containers/ArrayViewTest.h
index 35344eecf1041725b33c84664b8bca2769f3d525..7f1fb6941205e56fd9a4d6c021d4b46ab57b44de 100644
--- a/src/UnitTests/Containers/ArrayViewTest.h
+++ b/src/UnitTests/Containers/ArrayViewTest.h
@@ -98,27 +98,6 @@ using ViewTypes = ::testing::Types<
ArrayView< float, Devices::Cuda, long >,
ArrayView< double, Devices::Cuda, long >,
ArrayView< MyData, Devices::Cuda, long >
-#endif
-#ifdef HAVE_MIC
- ,
- ArrayView< int, Devices::MIC, short >,
- ArrayView< long, Devices::MIC, short >,
- ArrayView< float, Devices::MIC, short >,
- ArrayView< double, Devices::MIC, short >,
- // TODO: MyData does not work on MIC
-// ArrayView< MyData, Devices::MIC, short >,
- ArrayView< int, Devices::MIC, int >,
- ArrayView< long, Devices::MIC, int >,
- ArrayView< float, Devices::MIC, int >,
- ArrayView< double, Devices::MIC, int >,
- // TODO: MyData does not work on MIC
-// ArrayView< MyData, Devices::MIC, int >,
- ArrayView< int, Devices::MIC, long >,
- ArrayView< long, Devices::MIC, long >,
- ArrayView< float, Devices::MIC, long >,
- ArrayView< double, Devices::MIC, long >,
- // TODO: MyData does not work on MIC
-// ArrayView< MyData, Devices::MIC, long >,
#endif
// all ArrayView tests should also work with VectorView
@@ -133,11 +112,6 @@ using ViewTypes = ::testing::Types<
VectorView< float, Devices::Cuda, long >,
VectorView< double, Devices::Cuda, long >
#endif
-#ifdef HAVE_MIC
- ,
- VectorView< float, Devices::MIC, long >,
- VectorView< double, Devices::MIC, long >
-#endif
>;
TYPED_TEST_SUITE( ArrayViewTest, ViewTypes );
@@ -289,14 +263,6 @@ void testArrayViewElementwiseAccess( Array< Value, Devices::Cuda, Index >&& u )
#endif
}
-template< typename Value, typename Index >
-void testArrayViewElementwiseAccess( Array< Value, Devices::MIC, Index >&& u )
-{
-#ifdef HAVE_MIC
- // TODO
-#endif
-}
-
TYPED_TEST( ArrayViewTest, elementwiseAccess )
{
using ArrayType = typename TestFixture::ArrayType;
diff --git a/src/UnitTests/Containers/VectorTestSetup.h b/src/UnitTests/Containers/VectorTestSetup.h
index 5c342dced87f713824344cb43ee1c9922dbf0ef6..c8ec42bea482a1691fca97efecb8342985f8207d 100644
--- a/src/UnitTests/Containers/VectorTestSetup.h
+++ b/src/UnitTests/Containers/VectorTestSetup.h
@@ -76,21 +76,6 @@ using VectorTypes = ::testing::Types<
//Vector< Quad< float >, Devices::Cuda, long >,
//Vector< Quad< double >, Devices::Cuda, long >
#endif
-#ifdef HAVE_MIC
- ,
- Vector< int, Devices::MIC, short >,
- Vector< long, Devices::MIC, short >,
- Vector< float, Devices::MIC, short >,
- Vector< double, Devices::MIC, short >,
- Vector< int, Devices::MIC, int >,
- Vector< long, Devices::MIC, int >,
- Vector< float, Devices::MIC, int >,
- Vector< double, Devices::MIC, int >,
- Vector< int, Devices::MIC, long >,
- Vector< long, Devices::MIC, long >,
- Vector< float, Devices::MIC, long >,
- Vector< double, Devices::MIC, long >
-#endif
>;
TYPED_TEST_SUITE( VectorTest, VectorTypes );