diff --git a/src/TNL/CMakeLists.txt b/src/TNL/CMakeLists.txt
index 8d713d4ad7bce5337dc803cee14659a922e29c2e..8cdf8c581ed95272a81e40942bfe0e142d8534b5 100755
--- a/src/TNL/CMakeLists.txt
+++ b/src/TNL/CMakeLists.txt
@@ -26,8 +26,6 @@ set( headers
File_impl.h
FileName.h
Object.h
- List.h
- List_impl.h
Logger.h
Logger_impl.h
Math.h
diff --git a/src/TNL/Config/ConfigDescription.cpp b/src/TNL/Config/ConfigDescription.cpp
index 6822aec9ce7d0551a2e598ef93f394d95e8bd57b..6145103ef64c155cd8b4387d381e135ce47cba70 100644
--- a/src/TNL/Config/ConfigDescription.cpp
+++ b/src/TNL/Config/ConfigDescription.cpp
@@ -129,7 +129,7 @@ bool Config::ConfigDescription :: checkMissingEntries( Config::ParameterContaine
{
int i;
const int size = entries. getSize();
- List< String > missingParameters;
+ Containers::List< String > missingParameters;
for( i = 0; i < size; i ++ )
{
const char* entry_name = entries[ i ] -> name. getString();
diff --git a/src/TNL/Config/ConfigDescription.h b/src/TNL/Config/ConfigDescription.h
index 5c210fa6df578267f1134697092ba6cefd454dfb..4c93e2482bc4bc9ff7d460bcc23aa98bf18fc2d4 100644
--- a/src/TNL/Config/ConfigDescription.h
+++ b/src/TNL/Config/ConfigDescription.h
@@ -11,7 +11,7 @@
#pragma once
#include <TNL/String.h>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/param-types.h>
#include <TNL/Config/ConfigEntryType.h>
#include <TNL/Config/ConfigEntry.h>
@@ -166,7 +166,7 @@ class ConfigDescription
protected:
- List< ConfigEntryBase* > entries;
+ Containers::List< ConfigEntryBase* > entries;
ConfigEntryBase* currentEntry;
diff --git a/src/TNL/Config/ConfigEntry.h b/src/TNL/Config/ConfigEntry.h
index fa452fb4ea0ba23a6bc320434f8de170d9620c24..759a77ef3d52d991c4a1f879a7e1e5390023159d 100644
--- a/src/TNL/Config/ConfigEntry.h
+++ b/src/TNL/Config/ConfigEntry.h
@@ -11,6 +11,7 @@
#pragma once
#include <TNL/Config/ConfigEntryBase.h>
+#include <TNL/Containers/List.h>
namespace TNL {
namespace Config {
@@ -20,7 +21,7 @@ struct ConfigEntry : public ConfigEntryBase
{
EntryType defaultValue;
- List< EntryType > enumValues;
+ Containers::List< EntryType > enumValues;
public:
@@ -61,7 +62,7 @@ struct ConfigEntry : public ConfigEntryBase
return convertToString( defaultValue );
};
- List< EntryType >& getEnumValues()
+ Containers::List< EntryType >& getEnumValues()
{
return this->enumValues;
}
diff --git a/src/TNL/Config/ConfigEntryList.h b/src/TNL/Config/ConfigEntryList.h
index 01db24d6d6543bda7b3a12d24fa0772b3313c5ff..e8fabf5c58338177e38fa8a1ef652400c2ac79c7 100644
--- a/src/TNL/Config/ConfigEntryList.h
+++ b/src/TNL/Config/ConfigEntryList.h
@@ -11,6 +11,7 @@
#pragma once
#include <TNL/Config/ConfigEntryBase.h>
+#include <TNL/Containers/List.h>
namespace TNL {
namespace Config {
@@ -20,7 +21,7 @@ struct ConfigEntryList : public ConfigEntryBase
{
EntryType defaultValue;
- List< EntryType > enumValues;
+ Containers::List< EntryType > enumValues;
public:
@@ -49,12 +50,12 @@ struct ConfigEntryList : public ConfigEntryBase
String getEntryType() const
{
- return TNL::getType< List< EntryType > >();
+ return TNL::getType< Containers::List< EntryType > >();
}
String getUIEntryType() const
{
- return TNL::Config::getUIEntryType< List< EntryType > >();
+ return TNL::Config::getUIEntryType< Containers::List< EntryType > >();
}
String printDefaultValue() const
@@ -62,7 +63,7 @@ struct ConfigEntryList : public ConfigEntryBase
return convertToString( defaultValue );
};
- List< EntryType >& getEnumValues()
+ Containers::List< EntryType >& getEnumValues()
{
return this->enumValues;
}
@@ -84,7 +85,7 @@ struct ConfigEntryList : public ConfigEntryBase
std::cout << " ";
}
- bool checkValue( const List< EntryType >& values ) const
+ bool checkValue( const Containers::List< EntryType >& values ) const
{
if( this->enumValues.getSize() != 0 )
{
diff --git a/src/TNL/Config/ConfigEntryType.h b/src/TNL/Config/ConfigEntryType.h
index 0d104cb63add85408a70ecdbad400426e3ee1ee5..4ba052d022dd160ad633983d255b5dbfbb318edb 100644
--- a/src/TNL/Config/ConfigEntryType.h
+++ b/src/TNL/Config/ConfigEntryType.h
@@ -10,6 +10,8 @@
#pragma once
+#include <TNL/Containers/List.h>
+
namespace TNL {
namespace Config {
@@ -21,10 +23,10 @@ template<> inline String getUIEntryType< bool >() { return "bool"; };
template<> inline String getUIEntryType< int >() { return "integer"; };
template<> inline String getUIEntryType< double >() { return "real"; };
-template<> inline String getUIEntryType< List< String > >() { return "list of string"; };
-template<> inline String getUIEntryType< List< bool > >() { return "list of bool"; };
-template<> inline String getUIEntryType< List< int > >() { return "list of integer"; };
-template<> inline String getUIEntryType< List< double > >() { return "list of real"; };
+template<> inline String getUIEntryType< Containers::List< String > >() { return "list of string"; };
+template<> inline String getUIEntryType< Containers::List< bool > >() { return "list of bool"; };
+template<> inline String getUIEntryType< Containers::List< int > >() { return "list of integer"; };
+template<> inline String getUIEntryType< Containers::List< double > >() { return "list of real"; };
struct ConfigEntryType
{
diff --git a/src/TNL/Config/ParameterContainer.cpp b/src/TNL/Config/ParameterContainer.cpp
index ef2b335dae862cef51cff879eb8cf26cfa8f4a3d..b2e77dd5e21d4747cb89dfddf6cb53f4ef423a16 100644
--- a/src/TNL/Config/ParameterContainer.cpp
+++ b/src/TNL/Config/ParameterContainer.cpp
@@ -200,7 +200,7 @@ parseCommandLine( int argc, char* argv[],
std::cerr << "Missing value for the parameter " << option << "." << std::endl;
return false;
}
- List< String > parsedEntryType;
+ Containers::List< String > parsedEntryType;
if( ! parseObjectType( entryType, parsedEntryType ) )
{
std::cerr << "Internal error: Uknown config entry type " << entryType << "." << std::endl;
@@ -208,26 +208,26 @@ parseCommandLine( int argc, char* argv[],
}
if( parsedEntryType[ 0 ] == "List" )
{
- List< String >* string_list( 0 );
- List< bool >* bool_list( 0 );
- List< int >* integer_list( 0 );
- List< double >* real_list( 0 );
+ Containers::List< String >* string_list( 0 );
+ Containers::List< bool >* bool_list( 0 );
+ Containers::List< int >* integer_list( 0 );
+ Containers::List< double >* real_list( 0 );
if( parsedEntryType[ 1 ] == "String" )
- string_list = new List< String >;
+ string_list = new Containers::List< String >;
if( parsedEntryType[ 1 ] == "bool" )
- bool_list = new List< bool >;
+ bool_list = new Containers::List< bool >;
if( parsedEntryType[ 1 ] == "int" )
- integer_list = new List< int >;
+ integer_list = new Containers::List< int >;
if( parsedEntryType[ 1 ] == "double" )
- real_list = new List< double >;
+ real_list = new Containers::List< double >;
while( i < argc && ( ( argv[ i ] )[ 0 ] != '-' || ( atof( argv[ i ] ) < 0.0 && ( integer_list || real_list ) ) ) )
{
const char* value = argv[ i ++ ];
if( string_list )
{
- /*if( ! ( ( ConfigEntry< List< String > >* ) entry )->checkValue( String( value ) ) )
+ /*if( ! ( ( ConfigEntry< Containers::List< String > >* ) entry )->checkValue( String( value ) ) )
{
delete string_list;
return false;
@@ -246,7 +246,7 @@ parseCommandLine( int argc, char* argv[],
}
if( integer_list )
{
- /*if( ! ( ConfigEntry< List< int > >* ) entry->checkValue( atoi( value ) ) )
+ /*if( ! ( ConfigEntry< Containers::List< int > >* ) entry->checkValue( atoi( value ) ) )
{
delete integer_list;
return false;
@@ -255,7 +255,7 @@ parseCommandLine( int argc, char* argv[],
}
if( real_list )
{
- /*if( ! ( ConfigEntry< List< double > >* ) entry->checkValue( atof( value ) ) )
+ /*if( ! ( ConfigEntry< Containers::List< double > >* ) entry->checkValue( atof( value ) ) )
{
delete real_list;
return false;
@@ -265,22 +265,22 @@ parseCommandLine( int argc, char* argv[],
}
if( string_list )
{
- parameters. addParameter< List< String > >( option, *string_list );
+ parameters. addParameter< Containers::List< String > >( option, *string_list );
delete string_list;
}
if( bool_list )
{
- parameters. addParameter< List< bool > >( option, *bool_list );
+ parameters. addParameter< Containers::List< bool > >( option, *bool_list );
delete bool_list;
}
if( integer_list )
{
- parameters. addParameter< List< int > >( option, *integer_list );
+ parameters. addParameter< Containers::List< int > >( option, *integer_list );
delete integer_list;
}
if( real_list )
{
- parameters. addParameter< List< double > >( option, *real_list );
+ parameters. addParameter< Containers::List< double > >( option, *real_list );
delete real_list;
}
if( i < argc ) i --;
diff --git a/src/TNL/Config/ParameterContainer.h b/src/TNL/Config/ParameterContainer.h
index 8230490a38de13b80613c201b046995b4ae1ae19..5f6df343000fe3a28eb1afe5edc1f54156ba451d 100644
--- a/src/TNL/Config/ParameterContainer.h
+++ b/src/TNL/Config/ParameterContainer.h
@@ -10,7 +10,7 @@
#pragma once
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/Config/ConfigDescription.h>
#include <TNL/mpi-supp.h>
#include <TNL/param-types.h>
@@ -96,7 +96,7 @@ class ParameterContainer
protected:
- List< tnlParameterBase* > parameters;
+ Containers::List< tnlParameterBase* > parameters;
};
diff --git a/src/TNL/Containers/ArrayOperations.h b/src/TNL/Containers/Algorithms/ArrayOperations.h
similarity index 96%
rename from src/TNL/Containers/ArrayOperations.h
rename to src/TNL/Containers/Algorithms/ArrayOperations.h
index 042270d5b7455aeb06500a98ddc4b161613b477d..e32c7fd288c098ae2e648c1ad2fdbc7a6ccfc324 100644
--- a/src/TNL/Containers/ArrayOperations.h
+++ b/src/TNL/Containers/Algorithms/ArrayOperations.h
@@ -15,6 +15,7 @@
namespace TNL {
namespace Containers {
+namespace Algorithms {
template< typename DestinationDevice,
typename SourceDevice = DestinationDevice >
@@ -64,7 +65,6 @@ class ArrayOperations< Devices::Host >
static bool compareMemory( const Element1* destination,
const Element2* source,
const Index size );
-
};
template<>
@@ -152,8 +152,9 @@ class ArrayOperations< Devices::Host, Devices::Cuda >
const Index size );
};
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
-#include <TNL/Containers/ArrayOperationsHost_impl.h>
-#include <TNL/Containers/ArrayOperationsCuda_impl.h>
+#include <TNL/Containers/Algorithms/ArrayOperationsHost_impl.h>
+#include <TNL/Containers/Algorithms/ArrayOperationsCuda_impl.h>
diff --git a/src/TNL/Containers/ArrayOperationsCuda_impl.h b/src/TNL/Containers/Algorithms/ArrayOperationsCuda_impl.h
similarity index 92%
rename from src/TNL/Containers/ArrayOperationsCuda_impl.h
rename to src/TNL/Containers/Algorithms/ArrayOperationsCuda_impl.h
index 68b58e4621afe3e3c7e7b45394d2e25c752bf6a2..0793201d5a8d274638dc0044d6368bf9ba85ae26 100644
--- a/src/TNL/Containers/ArrayOperationsCuda_impl.h
+++ b/src/TNL/Containers/Algorithms/ArrayOperationsCuda_impl.h
@@ -11,17 +11,22 @@
#pragma once
#include <iostream>
+
#include <TNL/tnlConfig.h>
#include <TNL/Math.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Containers/Algorithms/Reduction.h>
#include <TNL/Containers/Algorithms/reduction-operations.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
template< typename Element, typename Index >
-bool ArrayOperations< Devices::Cuda >::allocateMemory( Element*& data,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda >::
+allocateMemory( Element*& data,
+ const Index size )
{
#ifdef HAVE_CUDA
checkCudaDevice;
@@ -36,7 +41,9 @@ bool ArrayOperations< Devices::Cuda >::allocateMemory( Element*& data,
}
template< typename Element >
-bool ArrayOperations< Devices::Cuda >::freeMemory( Element* data )
+bool
+ArrayOperations< Devices::Cuda >::
+freeMemory( Element* data )
{
Assert( data, );
#ifdef HAVE_CUDA
@@ -50,15 +57,19 @@ bool ArrayOperations< Devices::Cuda >::freeMemory( Element* data )
}
template< typename Element >
-void ArrayOperations< Devices::Cuda >::setMemoryElement( Element* data,
- const Element& value )
+void
+ArrayOperations< Devices::Cuda >::
+setMemoryElement( Element* data,
+ const Element& value )
{
Assert( data, );
ArrayOperations< Devices::Cuda >::setMemory( data, value, 1 );
}
template< typename Element >
-Element ArrayOperations< Devices::Cuda >::getMemoryElement( const Element* data )
+Element
+ArrayOperations< Devices::Cuda >::
+getMemoryElement( const Element* data )
{
Assert( data, );
Element result;
@@ -67,14 +78,18 @@ Element ArrayOperations< Devices::Cuda >::getMemoryElement( const Element* data
}
template< typename Element, typename Index >
-Element& ArrayOperations< Devices::Cuda >::getArrayElementReference( Element* data, const Index i )
+Element&
+ArrayOperations< Devices::Cuda >::
+getArrayElementReference( Element* data, const Index i )
{
Assert( data, );
return data[ i ];
}
template< typename Element, typename Index >
-const Element& ArrayOperations< Devices::Cuda >::getArrayElementReference( const Element* data, const Index i )
+const
+Element& ArrayOperations< Devices::Cuda >::
+getArrayElementReference( const Element* data, const Index i )
{
Assert( data, );
return data[ i ];
@@ -83,9 +98,10 @@ const Element& ArrayOperations< Devices::Cuda >::getArrayElementReference( const
#ifdef HAVE_CUDA
template< typename Element, typename Index >
-__global__ void setArrayValueCudaKernel( Element* data,
- const Index size,
- const Element value )
+__global__ void
+setArrayValueCudaKernel( Element* data,
+ const Index size,
+ const Element value )
{
Index elementIdx = blockDim. x * blockIdx. x + threadIdx. x;
const Index maxGridSize = blockDim. x * gridDim. x;
@@ -98,9 +114,11 @@ __global__ void setArrayValueCudaKernel( Element* data,
#endif
template< typename Element, typename Index >
-bool ArrayOperations< Devices::Cuda >::setMemory( Element* data,
- const Element& value,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda >::
+setMemory( Element* data,
+ const Element& value,
+ const Index size )
{
Assert( data, );
#ifdef HAVE_CUDA
@@ -120,9 +138,10 @@ bool ArrayOperations< Devices::Cuda >::setMemory( Element* data,
template< typename DestinationElement,
typename SourceElement,
typename Index >
-__global__ void copyMemoryCudaToCudaKernel( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+__global__ void
+copyMemoryCudaToCudaKernel( DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
Index elementIdx = blockDim. x * blockIdx. x + threadIdx. x;
const Index maxGridSize = blockDim. x * gridDim. x;
@@ -137,9 +156,11 @@ __global__ void copyMemoryCudaToCudaKernel( DestinationElement* destination,
template< typename DestinationElement,
typename SourceElement,
typename Index >
-bool ArrayOperations< Devices::Cuda >::copyMemory( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda >::
+copyMemory( DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
Assert( destination, );
Assert( source, );
@@ -170,9 +191,11 @@ bool ArrayOperations< Devices::Cuda >::copyMemory( DestinationElement* destinati
template< typename Element1,
typename Element2,
typename Index >
-bool ArrayOperations< Devices::Cuda >::compareMemory( const Element1* destination,
- const Element2* source,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda >::
+compareMemory( const Element1* destination,
+ const Element2* source,
+ const Index size )
{
Assert( destination, );
Assert( source, );
@@ -190,9 +213,11 @@ bool ArrayOperations< Devices::Cuda >::compareMemory( const Element1* destinatio
template< typename DestinationElement,
typename SourceElement,
typename Index >
-bool ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+bool
+ArrayOperations< Devices::Host, Devices::Cuda >::
+copyMemory( DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
Assert( destination, );
Assert( source, );
@@ -251,9 +276,11 @@ bool ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory( DestinationEle
template< typename Element1,
typename Element2,
typename Index >
-bool ArrayOperations< Devices::Host, Devices::Cuda >::compareMemory( const Element1* destination,
- const Element2* source,
- const Index size )
+bool
+ArrayOperations< Devices::Host, Devices::Cuda >::
+compareMemory( const Element1* destination,
+ const Element2* source,
+ const Index size )
{
/***
* Here, destination is on host and source is on CUDA device.
@@ -303,9 +330,11 @@ bool ArrayOperations< Devices::Host, Devices::Cuda >::compareMemory( const Eleme
template< typename DestinationElement,
typename SourceElement,
typename Index >
-bool ArrayOperations< Devices::Cuda, Devices::Host >::copyMemory( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda, Devices::Host >::
+copyMemory( DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
Assert( destination, );
Assert( source, );
@@ -364,9 +393,11 @@ bool ArrayOperations< Devices::Cuda, Devices::Host >::copyMemory( DestinationEle
template< typename Element1,
typename Element2,
typename Index >
-bool ArrayOperations< Devices::Cuda, Devices::Host >::compareMemory( const Element1* hostData,
- const Element2* deviceData,
- const Index size )
+bool
+ArrayOperations< Devices::Cuda, Devices::Host >::
+compareMemory( const Element1* hostData,
+ const Element2* deviceData,
+ const Index size )
{
Assert( hostData, );
Assert( deviceData, );
@@ -651,5 +682,6 @@ extern template bool ArrayOperations< Devices::Cuda >::setMemory< long double, l
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
diff --git a/src/TNL/Containers/ArrayOperationsHost_impl.h b/src/TNL/Containers/Algorithms/ArrayOperationsHost_impl.h
similarity index 91%
rename from src/TNL/Containers/ArrayOperationsHost_impl.h
rename to src/TNL/Containers/Algorithms/ArrayOperationsHost_impl.h
index 6bebb34a4e5ff52c5f6cb4ba6804a3073e7ad15f..8d6df347765c636f67bd6122801fa26971670501 100644
--- a/src/TNL/Containers/ArrayOperationsHost_impl.h
+++ b/src/TNL/Containers/Algorithms/ArrayOperationsHost_impl.h
@@ -11,15 +11,20 @@
#pragma once
#include <type_traits>
-#include <TNL/tnlConfig.h>
#include <string.h>
+#include <TNL/tnlConfig.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
+
namespace TNL {
namespace Containers {
+namespace Algorithms {
template< typename Element, typename Index >
-bool ArrayOperations< Devices::Host >::allocateMemory( Element*& data,
- const Index size )
+bool
+ArrayOperations< Devices::Host >::
+allocateMemory( Element*& data,
+ const Index size )
{
if( ! ( data = new Element[ size ] ) )
return false;
@@ -27,42 +32,54 @@ bool ArrayOperations< Devices::Host >::allocateMemory( Element*& data,
}
template< typename Element >
-bool ArrayOperations< Devices::Host >::freeMemory( Element* data )
+bool
+ArrayOperations< Devices::Host >::
+freeMemory( Element* data )
{
delete[] data;
return true;
}
template< typename Element >
-void ArrayOperations< Devices::Host >::setMemoryElement( Element* data,
- const Element& value )
+void
+ArrayOperations< Devices::Host >::
+setMemoryElement( Element* data,
+ const Element& value )
{
*data = value;
};
template< typename Element >
-Element ArrayOperations< Devices::Host >::getMemoryElement( Element* data )
+Element
+ArrayOperations< Devices::Host >::
+getMemoryElement( Element* data )
{
return *data;
};
template< typename Element, typename Index >
-Element& ArrayOperations< Devices::Host >::getArrayElementReference( Element* data,
- const Index i )
+Element&
+ArrayOperations< Devices::Host >::
+getArrayElementReference( Element* data,
+ const Index i )
{
return data[ i ];
};
template< typename Element, typename Index >
-const Element& ArrayOperations< Devices::Host >::getArrayElementReference( const Element* data,
- const Index i )
+const Element&
+ArrayOperations< Devices::Host >::
+getArrayElementReference( const Element* data,
+ const Index i )
{
return data[ i ];
};
template< typename Element, typename Index >
-bool ArrayOperations< Devices::Host >::setMemory( Element* data,
- const Element& value,
- const Index size )
+bool
+ArrayOperations< Devices::Host >::
+setMemory( Element* data,
+ const Element& value,
+ const Index size )
{
for( Index i = 0; i < size; i ++ )
data[ i ] = value;
@@ -72,9 +89,11 @@ bool ArrayOperations< Devices::Host >::setMemory( Element* data,
template< typename DestinationElement,
typename SourceElement,
typename Index >
-bool ArrayOperations< Devices::Host >::copyMemory( DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+bool
+ArrayOperations< Devices::Host >::
+copyMemory( DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
if( std::is_same< DestinationElement, SourceElement >::value )
memcpy( destination, source, size * sizeof( DestinationElement ) );
@@ -87,9 +106,11 @@ bool ArrayOperations< Devices::Host >::copyMemory( DestinationElement* destinati
template< typename DestinationElement,
typename SourceElement,
typename Index >
-bool ArrayOperations< Devices::Host >::compareMemory( const DestinationElement* destination,
- const SourceElement* source,
- const Index size )
+bool
+ArrayOperations< Devices::Host >::
+compareMemory( const DestinationElement* destination,
+ const SourceElement* source,
+ const Index size )
{
if( std::is_same< DestinationElement, SourceElement >::value )
{
@@ -284,5 +305,6 @@ extern template bool ArrayOperations< Devices::Host >::setMemory< long double, l
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
diff --git a/src/TNL/Containers/Algorithms/CMakeLists.txt b/src/TNL/Containers/Algorithms/CMakeLists.txt
index b98458026dfd7a372f456ab0fd2a4296f35dfba3..4b2744aced7e565692ee3a2469580fff1f7186a8 100755
--- a/src/TNL/Containers/Algorithms/CMakeLists.txt
+++ b/src/TNL/Containers/Algorithms/CMakeLists.txt
@@ -1,6 +1,9 @@
ADD_SUBDIRECTORY( TemplateExplicitInstantiation )
-set( headers cuda-prefix-sum.h
+set( headers ArrayOperations.h
+ ArrayOperationsHost_impl.h
+ ArrayOperationsCuda_impl.h
+ cuda-prefix-sum.h
cuda-prefix-sum_impl.h
reduction-operations.h
CublasWrapper.h
@@ -11,6 +14,9 @@ set( headers cuda-prefix-sum.h
Multireduction_impl.h
Reduction.h
Reduction_impl.h
+ VectorOperations.h
+ VectorOperationsHost_impl.h
+ VectorOperationsCuda_impl.h
)
INSTALL( FILES ${headers} DESTINATION include/tnl-${tnlVersion}/TNL/Containers/Algorithms )
diff --git a/src/TNL/Containers/Algorithms/CudaMultireductionKernel.h b/src/TNL/Containers/Algorithms/CudaMultireductionKernel.h
index e1dd4c36ef83d70597e43bb15066b13a224cfdec..ae0ee698ca923123cbb19389cfc563c7eb615648 100644
--- a/src/TNL/Containers/Algorithms/CudaMultireductionKernel.h
+++ b/src/TNL/Containers/Algorithms/CudaMultireductionKernel.h
@@ -50,9 +50,7 @@ CudaMultireductionKernel( Operation operation,
typedef typename Operation::IndexType IndexType;
typedef typename Operation::ResultType ResultType;
- extern __shared__ __align__ ( 8 ) char __sdata[];
-
- ResultType* sdata = reinterpret_cast< ResultType* >( __sdata );
+ ResultType* sdata = Devices::Cuda::getSharedMemory< ResultType >();
/***
* Get thread id (tid) and global element id (gid).
diff --git a/src/TNL/Containers/Algorithms/CudaReductionKernel.h b/src/TNL/Containers/Algorithms/CudaReductionKernel.h
index 607779eff4cded4a8fe1ad91c902733bda159c34..c6d7d49eca08b780181a095a5303e7f98d55e9df 100644
--- a/src/TNL/Containers/Algorithms/CudaReductionKernel.h
+++ b/src/TNL/Containers/Algorithms/CudaReductionKernel.h
@@ -48,9 +48,7 @@ CudaReductionKernel( Operation operation,
typedef typename Operation::IndexType IndexType;
typedef typename Operation::ResultType ResultType;
- extern __shared__ __align__ ( 8 ) char __sdata[];
-
- ResultType* sdata = reinterpret_cast< ResultType* >( __sdata );
+ ResultType* sdata = Devices::Cuda::getSharedMemory< ResultType >();
/***
* Get thread id (tid) and global thread id (gid).
diff --git a/src/TNL/Containers/Algorithms/Multireduction_impl.h b/src/TNL/Containers/Algorithms/Multireduction_impl.h
index 1eff176f014355a78afc7f0207d086a0f218e8c4..2de63dd37849750f284bd2f9d51f30aaaf6cb511 100644
--- a/src/TNL/Containers/Algorithms/Multireduction_impl.h
+++ b/src/TNL/Containers/Algorithms/Multireduction_impl.h
@@ -16,7 +16,7 @@
#include <TNL/Assert.h>
#include <TNL/Containers/Algorithms/reduction-operations.h>
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Containers/Algorithms/CudaMultireductionKernel.h>
#ifdef CUDA_REDUCTION_PROFILING
diff --git a/src/TNL/Containers/Algorithms/Reduction_impl.h b/src/TNL/Containers/Algorithms/Reduction_impl.h
index bbc312da0c574c90cfaa7ab637a5c8efaf035dc6..cd4d636c824481c56c56e30fb6f285045df0fe7c 100644
--- a/src/TNL/Containers/Algorithms/Reduction_impl.h
+++ b/src/TNL/Containers/Algorithms/Reduction_impl.h
@@ -16,7 +16,7 @@
#include <TNL/Assert.h>
#include <TNL/Containers/Algorithms/reduction-operations.h>
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Containers/Algorithms/CudaReductionKernel.h>
#ifdef CUDA_REDUCTION_PROFILING
@@ -59,10 +59,10 @@ reductionOnCudaDevice( Operation& operation,
{
RealType hostArray1[ minGPUReductionDataSize ];
RealType hostArray2[ minGPUReductionDataSize ];
- if( ! Containers::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( hostArray1, deviceInput1, size ) )
+ if( ! ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( hostArray1, deviceInput1, size ) )
return false;
if( deviceInput2 && !
- Containers::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( hostArray2, deviceInput2, size ) )
+ ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( hostArray2, deviceInput2, size ) )
return false;
result = operation.initialValue();
for( IndexType i = 0; i < size; i ++ )
@@ -96,7 +96,7 @@ reductionOnCudaDevice( Operation& operation,
* Transfer the reduced data from device to host.
*/
ResultType resultArray[ reducedSize ];
- if( ! Containers::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< ResultType, ResultType, IndexType >( resultArray, deviceAux1, reducedSize ) )
+ if( ! ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< ResultType, ResultType, IndexType >( resultArray, deviceAux1, reducedSize ) )
return false;
#ifdef CUDA_REDUCTION_PROFILING
diff --git a/src/TNL/Containers/ArrayOperationsCuda_impl.cpp b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cpp
similarity index 99%
rename from src/TNL/Containers/ArrayOperationsCuda_impl.cpp
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cpp
index 019b4f9ea48237e58237d9b964ce8403f50fee03..33a2aa1e10b4bf43383a3964f35c4bdc60466f32 100644
--- a/src/TNL/Containers/ArrayOperationsCuda_impl.cpp
+++ b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cpp
@@ -8,10 +8,11 @@
/* See Copyright Notice in tnl/Copyright */
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
@@ -290,8 +291,6 @@ template bool ArrayOperations< Devices::Cuda >::setMemory< long double, long int
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
-
-
-
diff --git a/src/TNL/Containers/ArrayOperationsCuda_impl.cu b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cu
similarity index 99%
rename from src/TNL/Containers/ArrayOperationsCuda_impl.cu
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cu
index 24e2de2063ed17cadba0834cb76de32cadfa7adb..8690a9256c613c07110c6c25df0186683c9254d0 100644
--- a/src/TNL/Containers/ArrayOperationsCuda_impl.cu
+++ b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsCuda_impl.cu
@@ -8,10 +8,11 @@
/* See Copyright Notice in tnl/Copyright */
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
@@ -290,5 +291,6 @@ template bool ArrayOperations< Devices::Cuda >::setMemory< long double, long int
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
diff --git a/src/TNL/Containers/ArrayOperationsHost_impl.cpp b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cpp
similarity index 99%
rename from src/TNL/Containers/ArrayOperationsHost_impl.cpp
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cpp
index a684b3608e20d627a92939822b4d9e0db3228cb8..28b2d6d494f968d00f2856c285ec73b8fba45c26 100644
--- a/src/TNL/Containers/ArrayOperationsHost_impl.cpp
+++ b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cpp
@@ -8,10 +8,11 @@
/* See Copyright Notice in tnl/Copyright */
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
@@ -194,5 +195,6 @@ template bool ArrayOperations< Devices::Host >::setMemory< long double, long int
#endif
+} // namespace Algorithms
} // namespace Containers
-} // namespace TNL
\ No newline at end of file
+} // namespace TNL
diff --git a/src/TNL/Containers/ArrayOperationsHost_impl.cu b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cu
similarity index 99%
rename from src/TNL/Containers/ArrayOperationsHost_impl.cu
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cu
index 2be45e40951538866bbc6998c9a451988eba916a..5c4f0b8789145a3328088d1f596d65b0ce2caee1 100644
--- a/src/TNL/Containers/ArrayOperationsHost_impl.cu
+++ b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/ArrayOperationsHost_impl.cu
@@ -8,10 +8,11 @@
/* See Copyright Notice in tnl/Copyright */
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
@@ -194,5 +195,6 @@ template bool ArrayOperations< Devices::Host >::setMemory< long double, long int
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
diff --git a/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/CMakeLists.txt b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/CMakeLists.txt
index 266952e328fc11a0235a8b737f612cb2ffdfa7dc..8770de6f389c2aa8ed7418c620067f8da053f681 100755
--- a/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/CMakeLists.txt
+++ b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/CMakeLists.txt
@@ -1,7 +1,12 @@
-SET( CURRENT_DIR ${CMAKE_SOURCE_DIR}/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation )
+SET( CURRENT_DIR ${CMAKE_SOURCE_DIR}/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation )
+set( common_SOURCES
+ ${CURRENT_DIR}/VectorOperationsHost_impl.cpp
+)
IF( BUILD_CUDA )
set( tnl_core_cuda_CUDA__SOURCES
${common_SOURCES}
+ ${CURRENT_DIR}/ArrayOperationsHost_impl.cu
+ ${CURRENT_DIR}/ArrayOperationsCuda_impl.cu
${CURRENT_DIR}/cuda-reduction-sum_impl.cu
${CURRENT_DIR}/cuda-reduction-min_impl.cu
${CURRENT_DIR}/cuda-reduction-max_impl.cu
@@ -21,13 +26,20 @@ IF( BUILD_CUDA )
${CURRENT_DIR}/cuda-reduction-diff-abs-sum_impl.cu
${CURRENT_DIR}/cuda-reduction-diff-abs-min_impl.cu
${CURRENT_DIR}/cuda-reduction-diff-abs-max_impl.cu
- ${CURRENT_DIR}/cuda-reduction-diff-l2-norm_impl.cu
- ${CURRENT_DIR}/cuda-reduction-diff-lp-norm_impl.cu
+ ${CURRENT_DIR}/cuda-reduction-diff-l2-norm_impl.cu
+ ${CURRENT_DIR}/cuda-reduction-diff-lp-norm_impl.cu
${CURRENT_DIR}/cuda-prefix-sum_impl.cu
- PARENT_SCOPE )
-endif()
+ ${CURRENT_DIR}/VectorOperationsCuda_impl.cu
+ PARENT_SCOPE )
+ELSE()
+ set( common_SOURCES
+ ${common_SOURCES}
+ ${CURRENT_DIR}/ArrayOperationsHost_impl.cpp
+ ${CURRENT_DIR}/ArrayOperationsCuda_impl.cpp
+ )
+ENDIF()
set( tnl_core_cuda_SOURCES
${common_SOURCES}
${CURRENT_DIR}/cuda-reduction_impl.cpp
- PARENT_SCOPE )
+ PARENT_SCOPE )
diff --git a/src/TNL/Containers/VectorOperationsCuda_impl.cpp b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsCuda_impl.cpp
similarity index 100%
rename from src/TNL/Containers/VectorOperationsCuda_impl.cpp
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsCuda_impl.cpp
diff --git a/src/TNL/Containers/VectorOperationsCuda_impl.cu b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsCuda_impl.cu
similarity index 100%
rename from src/TNL/Containers/VectorOperationsCuda_impl.cu
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsCuda_impl.cu
diff --git a/src/TNL/Containers/VectorOperationsHost_impl.cpp b/src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsHost_impl.cpp
similarity index 100%
rename from src/TNL/Containers/VectorOperationsHost_impl.cpp
rename to src/TNL/Containers/Algorithms/TemplateExplicitInstantiation/VectorOperationsHost_impl.cpp
diff --git a/src/TNL/Containers/VectorOperations.h b/src/TNL/Containers/Algorithms/VectorOperations.h
similarity index 87%
rename from src/TNL/Containers/VectorOperations.h
rename to src/TNL/Containers/Algorithms/VectorOperations.h
index 19df2add38bc06620046fa23e7e79a69a8458f64..fa490959f475fbe8a801fa33ef18767eb13d498d 100644
--- a/src/TNL/Containers/VectorOperations.h
+++ b/src/TNL/Containers/Algorithms/VectorOperations.h
@@ -17,6 +17,7 @@
namespace TNL {
namespace Containers {
+namespace Algorithms {
template< typename Device >
class VectorOperations{};
@@ -64,36 +65,36 @@ class VectorOperations< Devices::Host >
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceMax( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceMin( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceAbsMax( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceAbsMin( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceL1Norm( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceL2Norm( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceLpNorm( const Vector1& v1,
- const Vector2& v2,
- const typename Vector1::RealType& p );
+ const Vector2& v2,
+ const typename Vector1::RealType& p );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceSum( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector >
@@ -102,7 +103,7 @@ class VectorOperations< Devices::Host >
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getScalarProduct( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static void addVector( Vector1& y,
@@ -127,7 +128,6 @@ class VectorOperations< Devices::Host >
static void computeExclusivePrefixSum( Vector& v,
const typename Vector::IndexType begin,
const typename Vector::IndexType end );
-
};
template<>
@@ -177,11 +177,11 @@ class VectorOperations< Devices::Cuda >
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceMin( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceAbsMax( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceAbsMin( const Vector1& v1,
@@ -197,12 +197,12 @@ class VectorOperations< Devices::Cuda >
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceLpNorm( const Vector1& v1,
- const Vector2& v2,
- const typename Vector1::RealType& p );
+ const Vector2& v2,
+ const typename Vector1::RealType& p );
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getVectorDifferenceSum( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector >
static void vectorScalarMultiplication( Vector& v,
@@ -210,7 +210,7 @@ class VectorOperations< Devices::Cuda >
template< typename Vector1, typename Vector2 >
static typename Vector1::RealType getScalarProduct( const Vector1& v1,
- const Vector2& v2 );
+ const Vector2& v2 );
template< typename Vector1, typename Vector2 >
static void addVector( Vector1& y,
@@ -238,9 +238,9 @@ class VectorOperations< Devices::Cuda >
const typename Vector::IndexType end );
};
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
-#include <TNL/Containers/VectorOperationsHost_impl.h>
-#include <TNL/Containers/VectorOperationsCuda_impl.h>
-
+#include <TNL/Containers/Algorithms/VectorOperationsHost_impl.h>
+#include <TNL/Containers/Algorithms/VectorOperationsCuda_impl.h>
diff --git a/src/TNL/Containers/VectorOperationsCuda_impl.h b/src/TNL/Containers/Algorithms/VectorOperationsCuda_impl.h
similarity index 73%
rename from src/TNL/Containers/VectorOperationsCuda_impl.h
rename to src/TNL/Containers/Algorithms/VectorOperationsCuda_impl.h
index 3c6774613ef6448be062dcc9c899fcffeb9ad71e..a71c28018d5d97a14c7df2a13091998961b23fc1 100644
--- a/src/TNL/Containers/VectorOperationsCuda_impl.h
+++ b/src/TNL/Containers/Algorithms/VectorOperationsCuda_impl.h
@@ -11,96 +11,110 @@
#pragma once
#include <TNL/tnlConfig.h>
+#include <TNL/Containers/Algorithms/VectorOperations.h>
#include <TNL/Containers/Algorithms/cuda-prefix-sum.h>
#include <TNL/Containers/Algorithms/CublasWrapper.h>
namespace TNL {
namespace Containers {
+namespace Algorithms {
template< typename Vector >
-void VectorOperations< Devices::Cuda >::addElement( Vector& v,
- const typename Vector::IndexType i,
- const typename Vector::RealType& value )
+void
+VectorOperations< Devices::Cuda >::
+addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value )
{
v[ i ] += value;
}
template< typename Vector >
-void VectorOperations< Devices::Cuda >::addElement( Vector& v,
- const typename Vector::IndexType i,
- const typename Vector::RealType& value,
- const typename Vector::RealType& thisElementMultiplicator )
+void
+VectorOperations< Devices::Cuda >::
+addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value,
+ const typename Vector::RealType& thisElementMultiplicator )
{
v[ i ] = thisElementMultiplicator * v[ i ] + value;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Cuda > :: getVectorMax( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Cuda >::
+getVectorMax( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionMax< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Cuda > :: getVectorMin( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Cuda >::
+getVectorMin( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionMin< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Cuda > :: getVectorAbsMax( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Cuda >::
+getVectorAbsMax( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionAbsMax< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Cuda > :: getVectorAbsMin( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Cuda >::
+getVectorAbsMin( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionAbsMin< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
@@ -111,16 +125,16 @@ typename Vector::RealType
VectorOperations< Devices::Cuda >::
getVectorL1Norm( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionAbsSum< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
@@ -131,16 +145,16 @@ typename Vector::RealType
VectorOperations< Devices::Cuda >::
getVectorL2Norm( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionL2Norm< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return std::sqrt( result );
@@ -153,10 +167,10 @@ VectorOperations< Devices::Cuda >::
getVectorLpNorm( const Vector& v,
const typename Vector::RealType& p )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Assert( p > 0.0,
std::cerr << " p = " << p );
@@ -166,110 +180,120 @@ getVectorLpNorm( const Vector& v,
return getVectorL2Norm( v );
Real result( 0 );
Algorithms::tnlParallelReductionLpNorm< Real, Index > operation;
- operation. setPower( p );
+ operation.setPower( p );
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return std::pow( result, 1.0 / p );
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Cuda > :: getVectorSum( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Cuda >::
+getVectorSum( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
Real result( 0 );
Algorithms::tnlParallelReductionSum< Real, Index > operation;
reductionOnCudaDevice( operation,
- v. getSize(),
- v. getData(),
+ v.getSize(),
+ v.getData(),
( Real* ) 0,
result );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getVectorDifferenceMax( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getVectorDifferenceMax( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffMax< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getVectorDifferenceMin( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getVectorDifferenceMin( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffMin< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getVectorDifferenceAbsMax( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getVectorDifferenceAbsMax( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffAbsMax< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getVectorDifferenceAbsMin( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getVectorDifferenceAbsMin( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffAbsMin< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
@@ -280,18 +304,18 @@ VectorOperations< Devices::Cuda >::
getVectorDifferenceL1Norm( const Vector1& v1,
const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffAbsSum< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
@@ -302,18 +326,18 @@ VectorOperations< Devices::Cuda >::
getVectorDifferenceL2Norm( const Vector1& v1,
const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffL2Norm< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return ::sqrt( result );
}
@@ -324,55 +348,58 @@ typename Vector1::RealType
VectorOperations< Devices::Cuda >::
getVectorDifferenceLpNorm( const Vector1& v1,
const Vector2& v2,
- const typename Vector1 :: RealType& p )
+ const typename Vector1::RealType& p )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
Assert( p > 0.0,
std::cerr << " p = " << p );
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffLpNorm< Real, Index > operation;
operation.setPower( p );
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return ::pow( result, 1.0 / p );
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getVectorDifferenceSum( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getVectorDifferenceSum( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
Algorithms::tnlParallelReductionDiffSum< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
#ifdef HAVE_CUDA
template< typename Real, typename Index >
-__global__ void vectorScalarMultiplicationCudaKernel( Real* data,
- Index size,
- Real alpha )
+__global__ void
+vectorScalarMultiplicationCudaKernel( Real* data,
+ Index size,
+ Real alpha )
{
- Index elementIdx = blockDim. x * blockIdx. x + threadIdx. x;
- const Index maxGridSize = blockDim. x * gridDim. x;
+ Index elementIdx = blockDim.x * blockIdx.x + threadIdx.x;
+ const Index maxGridSize = blockDim.x * gridDim.x;
while( elementIdx < size )
{
data[ elementIdx ] *= alpha;
@@ -382,20 +409,22 @@ __global__ void vectorScalarMultiplicationCudaKernel( Real* data,
#endif
template< typename Vector >
-void VectorOperations< Devices::Cuda > :: vectorScalarMultiplication( Vector& v,
- const typename Vector::RealType& alpha )
+void
+VectorOperations< Devices::Cuda >::
+vectorScalarMultiplication( Vector& v,
+ const typename Vector::RealType& alpha )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
#ifdef HAVE_CUDA
dim3 blockSize( 0 ), gridSize( 0 );
const Index& size = v.getSize();
- blockSize. x = 256;
- Index blocksNumber = ceil( ( double ) size / ( double ) blockSize. x );
- gridSize. x = min( blocksNumber, Devices::Cuda::getMaxGridSize() );
+ blockSize.x = 256;
+ Index blocksNumber = ceil( ( double ) size / ( double ) blockSize.x );
+ gridSize.x = min( blocksNumber, Devices::Cuda::getMaxGridSize() );
vectorScalarMultiplicationCudaKernel<<< gridSize, blockSize >>>( v.getData(),
size,
alpha );
@@ -407,14 +436,16 @@ void VectorOperations< Devices::Cuda > :: vectorScalarMultiplication( Vector& v,
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getScalarProduct( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Cuda >::
+getScalarProduct( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0 );
/*#if defined HAVE_CUBLAS && defined HAVE_CUDA
@@ -425,9 +456,9 @@ typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getScalarProdu
#endif*/
Algorithms::tnlParallelReductionScalarProduct< Real, Index > operation;
reductionOnCudaDevice( operation,
- v1. getSize(),
- v1. getData(),
- v2. getData(),
+ v1.getSize(),
+ v1.getData(),
+ v2.getData(),
result );
return result;
}
@@ -435,14 +466,15 @@ typename Vector1 :: RealType VectorOperations< Devices::Cuda > :: getScalarProdu
#ifdef HAVE_CUDA
template< typename Real,
typename Index >
-__global__ void vectorAddVectorCudaKernel( Real* y,
- const Real* x,
- const Index size,
- const Real alpha,
- const Real thisMultiplicator )
+__global__ void
+vectorAddVectorCudaKernel( Real* y,
+ const Real* x,
+ const Index size,
+ const Real alpha,
+ const Real thisMultiplicator )
{
- Index elementIdx = blockDim. x * blockIdx. x + threadIdx. x;
- const Index maxGridSize = blockDim. x * gridDim. x;
+ Index elementIdx = blockDim.x * blockIdx.x + threadIdx.x;
+ const Index maxGridSize = blockDim.x * gridDim.x;
if( thisMultiplicator == 1.0 )
while( elementIdx < size )
{
@@ -455,25 +487,25 @@ __global__ void vectorAddVectorCudaKernel( Real* y,
y[ elementIdx ] = thisMultiplicator * y[ elementIdx ] + alpha * x[ elementIdx ];
elementIdx += maxGridSize;
}
-
}
#endif
template< typename Vector1, typename Vector2 >
-void VectorOperations< Devices::Cuda > :: addVector( Vector1& y,
- const Vector2& x,
- const typename Vector2::RealType& alpha,
- const typename Vector1::RealType& thisMultiplicator )
+void
+VectorOperations< Devices::Cuda >::
+addVector( Vector1& y,
+ const Vector2& x,
+ const typename Vector2::RealType& alpha,
+ const typename Vector1::RealType& thisMultiplicator )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( y. getSize() > 0, );
- Assert( y. getSize() == x. getSize(), );
+ Assert( y.getSize() > 0, );
+ Assert( y.getSize() == x.getSize(), );
Assert( y.getData() != 0, );
Assert( x.getData() != 0, );
-
#ifdef HAVE_CUDA
dim3 blockSize( 0 ), gridSize( 0 );
@@ -496,16 +528,17 @@ void VectorOperations< Devices::Cuda > :: addVector( Vector1& y,
#ifdef HAVE_CUDA
template< typename Real,
typename Index >
-__global__ void vectorAddVectorsCudaKernel( Real* v,
- const Real* v1,
- const Real* v2,
- const Index size,
- const Real multiplicator1,
- const Real multiplicator2,
- const Real thisMultiplicator )
+__global__ void
+vectorAddVectorsCudaKernel( Real* v,
+ const Real* v1,
+ const Real* v2,
+ const Index size,
+ const Real multiplicator1,
+ const Real multiplicator2,
+ const Real thisMultiplicator )
{
- Index elementIdx = blockDim. x * blockIdx. x + threadIdx. x;
- const Index maxGridSize = blockDim. x * gridDim. x;
+ Index elementIdx = blockDim.x * blockIdx.x + threadIdx.x;
+ const Index maxGridSize = blockDim.x * gridDim.x;
if( thisMultiplicator == 1.0 )
while( elementIdx < size )
{
@@ -524,7 +557,6 @@ __global__ void vectorAddVectorsCudaKernel( Real* v,
}
#endif
-
template< typename Vector1,
typename Vector2,
typename Vector3 >
@@ -537,8 +569,8 @@ addVectors( Vector1& v,
const typename Vector3::RealType& multiplicator2,
const typename Vector1::RealType& thisMultiplicator )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
Assert( v.getSize() > 0, );
Assert( v.getSize() == v1.getSize(), );
@@ -571,50 +603,54 @@ addVectors( Vector1& v,
}
template< typename Vector >
-void VectorOperations< Devices::Cuda >::computePrefixSum( Vector& v,
- typename Vector::IndexType begin,
- typename Vector::IndexType end )
+void
+VectorOperations< Devices::Cuda >::
+computePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
{
#ifdef HAVE_CUDA
- typedef Algorithms::tnlParallelReductionSum< typename Vector::RealType,
- typename Vector::IndexType > OperationType;
+ typedef Algorithms::tnlParallelReductionSum< typename Vector::RealType, typename Vector::IndexType > OperationType;
OperationType operation;
Algorithms::cudaPrefixSum< typename Vector::RealType,
- OperationType,
- typename Vector::IndexType >( end - begin,
- 256,
- &v.getData()[ begin ],
- &v.getData()[ begin ],
- operation,
- Algorithms::inclusivePrefixSum );
+ OperationType,
+ typename Vector::IndexType >
+ ( end - begin,
+ 256,
+ &v.getData()[ begin ],
+ &v.getData()[ begin ],
+ operation,
+ Algorithms::inclusivePrefixSum );
#else
CudaSupportMissingMessage;;
#endif
}
template< typename Vector >
-void VectorOperations< Devices::Cuda >::computeExclusivePrefixSum( Vector& v,
- typename Vector::IndexType begin,
- typename Vector::IndexType end )
+void
+VectorOperations< Devices::Cuda >::
+computeExclusivePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
{
#ifdef HAVE_CUDA
- typedef Algorithms::tnlParallelReductionSum< typename Vector::RealType,
- typename Vector::IndexType > OperationType;
+ typedef Algorithms::tnlParallelReductionSum< typename Vector::RealType, typename Vector::IndexType > OperationType;
OperationType operation;
-
Algorithms::cudaPrefixSum< typename Vector::RealType,
- OperationType,
- typename Vector::IndexType >( end - begin,
- 256,
- &v.getData()[ begin ],
- &v.getData()[ begin ],
- operation,
- Algorithms::exclusivePrefixSum );
+ OperationType,
+ typename Vector::IndexType >
+ ( end - begin,
+ 256,
+ &v.getData()[ begin ],
+ &v.getData()[ begin ],
+ operation,
+ Algorithms::exclusivePrefixSum );
#endif
}
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
@@ -624,6 +660,7 @@ void VectorOperations< Devices::Cuda >::computeExclusivePrefixSum( Vector& v,
namespace TNL {
namespace Containers {
+namespace Algorithms {
/****
* Max
@@ -875,6 +912,7 @@ extern template long double VectorOperations< Devices::Cuda >::getVectorDifferen
#endif
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
diff --git a/src/TNL/Containers/VectorOperationsHost_impl.h b/src/TNL/Containers/Algorithms/VectorOperationsHost_impl.h
similarity index 78%
rename from src/TNL/Containers/VectorOperationsHost_impl.h
rename to src/TNL/Containers/Algorithms/VectorOperationsHost_impl.h
index 102cc083d9db6959297e1c6dc6a3fc72a28c06e8..ff2519003f2ad1be3b587fe8a9e516615a78cd66 100644
--- a/src/TNL/Containers/VectorOperationsHost_impl.h
+++ b/src/TNL/Containers/Algorithms/VectorOperationsHost_impl.h
@@ -10,79 +10,102 @@
#pragma once
+#include <TNL/Containers/Algorithms/VectorOperations.h>
+
namespace TNL {
namespace Containers {
+namespace Algorithms {
static const int OpenMPVectorOperationsThreshold = 65536; // TODO: check this threshold
static const int PrefetchDistance = 128;
template< typename Vector >
-void VectorOperations< Devices::Host >::addElement( Vector& v,
- const typename Vector::IndexType i,
- const typename Vector::RealType& value )
+void
+VectorOperations< Devices::Host >::
+addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value )
{
v[ i ] += value;
}
template< typename Vector >
-void VectorOperations< Devices::Host >::addElement( Vector& v,
- const typename Vector::IndexType i,
- const typename Vector::RealType& value,
- const typename Vector::RealType& thisElementMultiplicator )
+void
+VectorOperations< Devices::Host >::
+addElement( Vector& v,
+ const typename Vector::IndexType i,
+ const typename Vector::RealType& value,
+ const typename Vector::RealType& thisElementMultiplicator )
{
v[ i ] = thisElementMultiplicator * v[ i ] + value;
}
template< typename Vector >
-typename Vector::RealType VectorOperations< Devices::Host >::getVectorMax( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Host >::
+getVectorMax( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
- Real result = v. getElement( 0 );
- const Index n = v. getSize();
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
+
+ Real result = v.getElement( 0 );
+ const Index n = v.getSize();
for( Index i = 1; i < n; i ++ )
- result = max( result, v. getElement( i ) );
+ result = max( result, v.getElement( i ) );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Host > :: getVectorMin( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Host >::
+getVectorMin( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
- Real result = v. getElement( 0 );
- const Index n = v. getSize();
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
+
+ Real result = v.getElement( 0 );
+ const Index n = v.getSize();
for( Index i = 1; i < n; i ++ )
- result = min( result, v. getElement( i ) );
+ result = min( result, v.getElement( i ) );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Host > :: getVectorAbsMax( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Host >::
+getVectorAbsMax( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
- Real result = std::fabs( v. getElement( 0 ) );
- const Index n = v. getSize();
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
+
+ Real result = std::fabs( v.getElement( 0 ) );
+ const Index n = v.getSize();
for( Index i = 1; i < n; i ++ )
- result = max( result, ( Real ) std::fabs( v. getElement( i ) ) );
+ result = max( result, ( Real ) std::fabs( v.getElement( i ) ) );
return result;
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Host > :: getVectorAbsMin( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Host >::
+getVectorAbsMin( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
- Real result = std::fabs( v. getElement( 0 ) );
- const Index n = v. getSize();
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
+
+ Real result = std::fabs( v.getElement( 0 ) );
+ const Index n = v.getSize();
for( Index i = 1; i < n; i ++ )
- result = min( result, ( Real ) std::fabs( v. getElement( i ) ) );
+ result = min( result, ( Real ) std::fabs( v.getElement( i ) ) );
return result;
}
@@ -91,12 +114,13 @@ typename Vector::RealType
VectorOperations< Devices::Host >::
getVectorL1Norm( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
Real result( 0.0 );
- const Index n = v. getSize();
+ const Index n = v.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() && n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -110,11 +134,12 @@ typename Vector::RealType
VectorOperations< Devices::Host >::
getVectorL2Norm( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
- Assert( v. getSize() > 0, );
- const Index n = v. getSize();
+ const Index n = v.getSize();
#ifdef OPTIMIZED_VECTOR_HOST_OPERATIONS
#ifdef __GNUC__
@@ -170,20 +195,22 @@ template< typename Vector >
typename Vector::RealType
VectorOperations< Devices::Host >::
getVectorLpNorm( const Vector& v,
- const typename Vector :: RealType& p )
+ const typename Vector::RealType& p )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
Assert( p > 0.0,
std::cerr << " p = " << p );
+
if( p == 1.0 )
return getVectorL1Norm( v );
if( p == 2.0 )
return getVectorL2Norm( v );
Real result( 0.0 );
- const Index n = v. getSize();
+ const Index n = v.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -193,14 +220,17 @@ getVectorLpNorm( const Vector& v,
}
template< typename Vector >
-typename Vector :: RealType VectorOperations< Devices::Host > :: getVectorSum( const Vector& v )
+typename Vector::RealType
+VectorOperations< Devices::Host >::
+getVectorSum( const Vector& v )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
- Assert( v. getSize() > 0, );
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
+
+ Assert( v.getSize() > 0, );
Real result( 0.0 );
- const Index n = v. getSize();
+ const Index n = v.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -210,66 +240,76 @@ typename Vector :: RealType VectorOperations< Devices::Host > :: getVectorSum( c
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Host > :: getVectorDifferenceMax( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getVectorDifferenceMax( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
- Real result = v1. getElement( 0 ) - v2. getElement( 0 );
- const Index n = v1. getSize();
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
+
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
+
+ Real result = v1.getElement( 0 ) - v2.getElement( 0 );
+ const Index n = v1.getSize();
for( Index i = 1; i < n; i ++ )
- result = max( result, v1. getElement( i ) - v2. getElement( i ) );
+ result = max( result, v1.getElement( i ) - v2.getElement( i ) );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Host > :: getVectorDifferenceMin( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getVectorDifferenceMin( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
- Real result = v1. getElement( 0 ) - v2. getElement( 0 );
- const Index n = v1. getSize();
+ Real result = v1.getElement( 0 ) - v2.getElement( 0 );
+ const Index n = v1.getSize();
for( Index i = 1; i < n; i ++ )
- result = min( result, v1. getElement( i ) - v2. getElement( i ) );
+ result = min( result, v1.getElement( i ) - v2.getElement( i ) );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Host > :: getVectorDifferenceAbsMax( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getVectorDifferenceAbsMax( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
- Real result = std::fabs( v1. getElement( 0 ) - v2. getElement( 0 ) );
- const Index n = v1. getSize();
+ Real result = std::fabs( v1.getElement( 0 ) - v2.getElement( 0 ) );
+ const Index n = v1.getSize();
for( Index i = 1; i < n; i ++ )
- result = max( result, ( Real ) std::fabs( v1. getElement( i ) - v2. getElement( i ) ) );
+ result = max( result, ( Real ) std::fabs( v1.getElement( i ) - v2.getElement( i ) ) );
return result;
}
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Host > :: getVectorDifferenceAbsMin( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getVectorDifferenceAbsMin( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result = std::fabs( v1[ 0 ] - v2[ 0 ] );
- const Index n = v1. getSize();
+ const Index n = v1.getSize();
for( Index i = 1; i < n; i ++ )
result = min( result, ( Real ) std::fabs( v1[ i ] - v2[ i ] ) );
return result;
@@ -281,14 +321,14 @@ VectorOperations< Devices::Host >::
getVectorDifferenceL1Norm( const Vector1& v1,
const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0.0 );
- const Index n = v1. getSize();
+ const Index n = v1.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -303,14 +343,14 @@ VectorOperations< Devices::Host >::
getVectorDifferenceL2Norm( const Vector1& v1,
const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0.0 );
- const Index n = v1. getSize();
+ const Index n = v1.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -330,13 +370,13 @@ getVectorDifferenceLpNorm( const Vector1& v1,
const Vector2& v2,
const typename Vector1::RealType& p )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
Assert( p > 0.0,
std::cerr << " p = " << p );
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
if( p == 1.0 )
return getVectorDifferenceL1Norm( v1, v2 );
@@ -344,46 +384,50 @@ getVectorDifferenceLpNorm( const Vector1& v1,
return getVectorDifferenceL2Norm( v1, v2 );
Real result( 0.0 );
- const Index n = v1. getSize();
+ const Index n = v1.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
for( Index i = 0; i < n; i ++ )
- result += std::pow( std::fabs( v1. getElement( i ) - v2. getElement( i ) ), p );
+ result += std::pow( std::fabs( v1.getElement( i ) - v2.getElement( i ) ), p );
return std::pow( result, 1.0 / p );
}
template< typename Vector1, typename Vector2 >
-typename Vector1::RealType VectorOperations< Devices::Host > :: getVectorDifferenceSum( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getVectorDifferenceSum( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
Real result( 0.0 );
- const Index n = v1. getSize();
+ const Index n = v1.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for reduction(+:result) if( TNL::Devices::Host::isOMPEnabled() &&n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
for( Index i = 0; i < n; i ++ )
- result += v1. getElement( i ) - v2. getElement( i );
+ result += v1.getElement( i ) - v2.getElement( i );
return result;
}
template< typename Vector >
-void VectorOperations< Devices::Host > :: vectorScalarMultiplication( Vector& v,
- const typename Vector :: RealType& alpha )
+void
+VectorOperations< Devices::Host >::
+vectorScalarMultiplication( Vector& v,
+ const typename Vector::RealType& alpha )
{
- typedef typename Vector :: RealType Real;
- typedef typename Vector :: IndexType Index;
+ typedef typename Vector::RealType Real;
+ typedef typename Vector::IndexType Index;
- Assert( v. getSize() > 0, );
+ Assert( v.getSize() > 0, );
- const Index n = v. getSize();
+ const Index n = v.getSize();
#ifdef HAVE_OPENMP
#pragma omp parallel for if( n > OpenMPVectorOperationsThreshold ) // TODO: check this threshold
#endif
@@ -393,15 +437,17 @@ void VectorOperations< Devices::Host > :: vectorScalarMultiplication( Vector& v,
template< typename Vector1, typename Vector2 >
-typename Vector1 :: RealType VectorOperations< Devices::Host > :: getScalarProduct( const Vector1& v1,
- const Vector2& v2 )
+typename Vector1::RealType
+VectorOperations< Devices::Host >::
+getScalarProduct( const Vector1& v1,
+ const Vector2& v2 )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( v1. getSize() > 0, );
- Assert( v1. getSize() == v2. getSize(), );
- const Index n = v1. getSize();
+ Assert( v1.getSize() > 0, );
+ Assert( v1.getSize() == v2.getSize(), );
+ const Index n = v1.getSize();
#ifdef OPTIMIZED_VECTOR_HOST_OPERATIONS
#ifdef __GNUC__
@@ -453,17 +499,20 @@ typename Vector1 :: RealType VectorOperations< Devices::Host > :: getScalarProdu
}
template< typename Vector1, typename Vector2 >
-void VectorOperations< Devices::Host > :: addVector( Vector1& y,
- const Vector2& x,
- const typename Vector2::RealType& alpha,
- const typename Vector1::RealType& thisMultiplicator )
+void
+VectorOperations< Devices::Host >::
+addVector( Vector1& y,
+ const Vector2& x,
+ const typename Vector2::RealType& alpha,
+ const typename Vector1::RealType& thisMultiplicator )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
- Assert( x. getSize() > 0, );
- Assert( x. getSize() == y. getSize(), );
- const Index n = y. getSize();
+ Assert( x.getSize() > 0, );
+ Assert( x.getSize() == y.getSize(), );
+
+ const Index n = y.getSize();
#ifdef OPTIMIZED_VECTOR_HOST_OPERATIONS
#ifdef __GNUC__
@@ -527,8 +576,8 @@ addVectors( Vector1& v,
const typename Vector3::RealType& multiplicator2,
const typename Vector1::RealType& thisMultiplicator )
{
- typedef typename Vector1 :: RealType Real;
- typedef typename Vector1 :: IndexType Index;
+ typedef typename Vector1::RealType Real;
+ typedef typename Vector1::IndexType Index;
Assert( v.getSize() > 0, );
Assert( v.getSize() == v1.getSize(), );
@@ -550,22 +599,28 @@ addVectors( Vector1& v,
}
template< typename Vector >
-void VectorOperations< Devices::Host >::computePrefixSum( Vector& v,
- typename Vector::IndexType begin,
- typename Vector::IndexType end )
+void
+VectorOperations< Devices::Host >::
+computePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
{
typedef typename Vector::IndexType Index;
+
for( Index i = begin + 1; i < end; i++ )
v[ i ] += v[ i - 1 ];
}
template< typename Vector >
-void VectorOperations< Devices::Host >::computeExclusivePrefixSum( Vector& v,
- typename Vector::IndexType begin,
- typename Vector::IndexType end )
+void
+VectorOperations< Devices::Host >::
+computeExclusivePrefixSum( Vector& v,
+ typename Vector::IndexType begin,
+ typename Vector::IndexType end )
{
typedef typename Vector::IndexType Index;
typedef typename Vector::RealType Real;
+
Real aux( v[ begin ] );
v[ begin ] = 0.0;
for( Index i = begin + 1; i < end; i++ )
@@ -576,8 +631,9 @@ void VectorOperations< Devices::Host >::computeExclusivePrefixSum( Vector& v,
}
}
+} // namespace Algorithms
} // namespace Containers
-} //namespace TNL
+} // namespace TNL
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
@@ -585,6 +641,7 @@ void VectorOperations< Devices::Host >::computeExclusivePrefixSum( Vector& v,
namespace TNL {
namespace Containers {
+namespace Algorithms {
/****
* Max
@@ -837,8 +894,8 @@ extern template long double VectorOperations< Devices::Host >::getVectorDifferen
#endif
#endif
+} // namespace Algorithms
} // namespace Containers
} // namespace TNL
-#endif
-
+#endif
diff --git a/src/TNL/Containers/Algorithms/cuda-prefix-sum_impl.h b/src/TNL/Containers/Algorithms/cuda-prefix-sum_impl.h
index d0591050d556aaca779ee2e861c6c6166b9d7346..971544a024454ef488db269cbddd3b4c755d055f 100644
--- a/src/TNL/Containers/Algorithms/cuda-prefix-sum_impl.h
+++ b/src/TNL/Containers/Algorithms/cuda-prefix-sum_impl.h
@@ -31,7 +31,7 @@ __global__ void cudaFirstPhaseBlockPrefixSum( const enumPrefixSumType prefixSumT
DataType* output,
DataType* auxArray )
{
- DataType* sharedData = TNL::Devices::getSharedMemory< DataType >();
+ DataType* sharedData = TNL::Devices::Cuda::getSharedMemory< DataType >();
DataType* auxData = &sharedData[ elementsInBlock + elementsInBlock / Devices::Cuda::getNumberOfSharedMemoryBanks() + 2 ];
DataType* warpSums = &auxData[ blockDim. x ];
diff --git a/src/TNL/Containers/Array.h b/src/TNL/Containers/Array.h
index 8eb769622002c12ac6d5d55d68cc22d4ffeea175..92a3ccd6fd519b50435751e3753218b7060778a4 100644
--- a/src/TNL/Containers/Array.h
+++ b/src/TNL/Containers/Array.h
@@ -11,22 +11,13 @@
#pragma once
#include <TNL/Object.h>
-#include <TNL/Containers/SharedArray.h>
+#include <TNL/File.h>
+#include <TNL/Devices/Host.h>
-// Forward declarations
namespace TNL {
-class File;
-
-namespace Devices {
- class Host;
-}
-
-
namespace Containers {
-
-template< typename Element, typename Device, typename Index >
-class SharedArray;
+template< int, typename > class StaticArray;
/****
* Array handles memory allocation and sharing of the same data between more Arrays.
diff --git a/src/TNL/Containers/Array_impl.h b/src/TNL/Containers/Array_impl.h
index 60ff37e2cf3759e6fa4b255d85337a0220e6ee9c..992d27b92bade5d6cc205578ccba47d77c198aac 100644
--- a/src/TNL/Containers/Array_impl.h
+++ b/src/TNL/Containers/Array_impl.h
@@ -15,7 +15,7 @@
#include <TNL/File.h>
#include <TNL/Math.h>
#include <TNL/param-types.h>
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Containers/ArrayIO.h>
#include <TNL/Containers/Array.h>
@@ -23,8 +23,8 @@ namespace TNL {
namespace Containers {
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index >::
Array()
: size( 0 ),
@@ -35,8 +35,8 @@ Array()
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index >::
Array( const IndexType& size )
: size( 0 ),
@@ -48,11 +48,11 @@ Array( const IndexType& size )
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index >::
Array( Element* data,
- const IndexType& size )
+ const IndexType& size )
: size( size ),
data( data ),
allocationPointer( 0 ),
@@ -61,12 +61,12 @@ Array( Element* data,
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index >::
Array( Array< Element, Device, Index >& array,
- const IndexType& begin,
- const IndexType& size )
+ const IndexType& begin,
+ const IndexType& size )
: size( size ),
data( &array.getData()[ begin ] ),
allocationPointer( array.allocationPointer ),
@@ -106,8 +106,8 @@ getType()
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
String
Array< Element, Device, Index >::
getTypeVirtual() const
@@ -116,8 +116,8 @@ getTypeVirtual() const
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
String
Array< Element, Device, Index >::
getSerializationType()
@@ -126,8 +126,8 @@ getSerializationType()
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
String
Array< Element, Device, Index >::
getSerializationTypeVirtual() const
@@ -136,8 +136,8 @@ getSerializationTypeVirtual() const
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
void
Array< Element, Device, Index >::
releaseData() const
@@ -146,14 +146,14 @@ releaseData() const
{
if( --*this->referenceCounter == 0 )
{
- ArrayOperations< Device >::freeMemory( this->allocationPointer );
+ Algorithms::ArrayOperations< Device >::freeMemory( this->allocationPointer );
delete this->referenceCounter;
//std::cerr << "Deallocating reference counter " << this->referenceCounter << std::endl;
}
}
else
if( allocationPointer )
- ArrayOperations< Device >::freeMemory( this->allocationPointer );
+ Algorithms::ArrayOperations< Device >::freeMemory( this->allocationPointer );
this->allocationPointer = 0;
this->data = 0;
this->size = 0;
@@ -172,7 +172,7 @@ setSize( const Index size )
<< "New size: " << size << std::endl );
if( this->size == size && allocationPointer && ! referenceCounter ) return true;
this->releaseData();
- ArrayOperations< Device >::allocateMemory( this->allocationPointer, size );
+ Algorithms::ArrayOperations< Device >::allocateMemory( this->allocationPointer, size );
this->data = this->allocationPointer;
this->size = size;
if( ! this->allocationPointer )
@@ -186,8 +186,8 @@ setSize( const Index size )
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
template< typename ArrayT >
bool
Array< Element, Device, Index >::
@@ -200,8 +200,8 @@ setLike( const ArrayT& array )
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
void
Array< Element, Device, Index >::
bind( Element* data,
@@ -213,8 +213,8 @@ bind( Element* data,
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
template< typename ArrayT >
void
Array< Element, Device, Index >::
@@ -255,8 +255,8 @@ bind( const ArrayT& array,
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
template< int Size >
void
Array< Element, Device, Index >::
@@ -303,8 +303,8 @@ getSize() const
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
void
Array< Element, Device, Index >::
setElement( const Index& i, const Element& x )
@@ -313,12 +313,12 @@ setElement( const Index& i, const Element& x )
std::cerr << "Wrong index for setElement method in Array "
<< " index is " << i
<< " and array size is " << this->getSize() );
- return ArrayOperations< Device > :: setMemoryElement( &( this->data[ i ] ), x );
+ return Algorithms::ArrayOperations< Device > :: setMemoryElement( &( this->data[ i ] ), x );
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Element
Array< Element, Device, Index >::
getElement( const Index& i ) const
@@ -327,7 +327,7 @@ getElement( const Index& i ) const
std::cerr << "Wrong index for getElement method in Array "
<< " index is " << i
<< " and array size is " << this->getSize() );
- return ArrayOperations< Device >::getMemoryElement( & ( this->data[ i ] ) );
+ return Algorithms::ArrayOperations< Device >::getMemoryElement( & ( this->data[ i ] ) );
};
template< typename Element,
@@ -346,8 +346,8 @@ operator[] ( const Index& i )
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
__cuda_callable__
inline const Element&
Array< Element, Device, Index >::
@@ -361,8 +361,8 @@ operator[] ( const Index& i ) const
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index >&
Array< Element, Device, Index >::
operator = ( const Array< Element, Device, Index >& array )
@@ -370,19 +370,19 @@ operator = ( const Array< Element, Device, Index >& array )
Assert( array. getSize() == this->getSize(),
std::cerr << "Source size: " << array. getSize() << std::endl
<< "Target size: " << this->getSize() << std::endl );
- ArrayOperations< Device > ::
- template copyMemory< Element,
- Element,
- Index >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ Algorithms::ArrayOperations< Device >::
+ template copyMemory< Element,
+ Element,
+ Index >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
return ( *this );
};
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
template< typename ArrayT >
Array< Element, Device, Index >&
Array< Element, Device, Index >::
@@ -391,13 +391,13 @@ operator = ( const ArrayT& array )
Assert( array. getSize() == this->getSize(),
std::cerr << "Source size: " << array. getSize() << std::endl
<< "Target size: " << this->getSize() << std::endl );
- ArrayOperations< Device, typename ArrayT::DeviceType > ::
- template copyMemory< Element,
- typename ArrayT::ElementType,
- typename ArrayT::IndexType >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ Algorithms::ArrayOperations< Device, typename ArrayT::DeviceType >::
+ template copyMemory< Element,
+ typename ArrayT::ElementType,
+ typename ArrayT::IndexType >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
return ( *this );
};
@@ -411,13 +411,13 @@ operator == ( const ArrayT& array ) const
{
if( array. getSize() != this->getSize() )
return false;
- return ArrayOperations< Device, typename ArrayT::DeviceType > ::
- template compareMemory< typename ArrayT::ElementType,
- Element,
- typename ArrayT::IndexType >
- ( this->getData(),
- array.getData(),
- array.getSize() );
+ return Algorithms::ArrayOperations< Device, typename ArrayT::DeviceType >::
+ template compareMemory< typename ArrayT::ElementType,
+ Element,
+ typename ArrayT::IndexType >
+ ( this->getData(),
+ array.getData(),
+ array.getSize() );
}
template< typename Element,
@@ -436,12 +436,12 @@ template< typename Element,
void Array< Element, Device, Index > :: setValue( const Element& e )
{
Assert( this->getData(),);
- ArrayOperations< Device > :: setMemory( this->getData(), e, this->getSize() );
+ Algorithms::ArrayOperations< Device >::setMemory( this->getData(), e, this->getSize() );
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
__cuda_callable__
const Element* Array< Element, Device, Index > :: getData() const
{
@@ -449,8 +449,8 @@ const Element* Array< Element, Device, Index > :: getData() const
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
__cuda_callable__
Element* Array< Element, Device, Index > :: getData()
{
@@ -458,8 +458,8 @@ Element* Array< Element, Device, Index > :: getData()
}
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Array< Element, Device, Index > :: operator bool() const
{
return data != 0;
@@ -467,8 +467,8 @@ Array< Element, Device, Index > :: operator bool() const
template< typename Element,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
template< typename IndexType2 >
void Array< Element, Device, Index > :: touch( IndexType2 touches ) const
{
diff --git a/src/TNL/Containers/CMakeLists.txt b/src/TNL/Containers/CMakeLists.txt
index d67b8f8b673bc08eeccdfb5a827d1ca016da957a..b35fa285d9623457ddd5c45a887428a55161fde1 100755
--- a/src/TNL/Containers/CMakeLists.txt
+++ b/src/TNL/Containers/CMakeLists.txt
@@ -3,12 +3,11 @@ ADD_SUBDIRECTORY( Algorithms )
set( headers Array.h
Array_impl.h
ArrayIO.h
- ArrayOperations.h
- ArrayOperationsHost_impl.h
- ArrayOperationsCuda_impl.h
DynamicTypeTag.h
IndexedSet.h
IndexedSet_impl.h
+ List.h
+ List_impl.h
MultiArray.h
MultiArray1D_impl.h
MultiArray2D_impl.h
@@ -29,9 +28,6 @@ set( headers Array.h
MultiVector2D_impl.h
MultiVector3D_impl.h
MultiVector4D_impl.h
- VectorOperations.h
- VectorOperationsHost_impl.h
- VectorOperationsCuda_impl.h
SharedVector.h
SharedVector_impl.h
StaticVector.h
@@ -46,7 +42,6 @@ set( common_SOURCES
${CURRENT_DIR}/MultiArray_impl.cpp
${CURRENT_DIR}/Array_impl.cpp
${CURRENT_DIR}/StaticArray_impl.cpp
- ${CURRENT_DIR}/VectorOperationsHost_impl.cpp
${CURRENT_DIR}/MultiVector_impl.cpp
${CURRENT_DIR}/SharedVector_impl.cpp
${CURRENT_DIR}/Vector_impl.cpp
@@ -56,26 +51,17 @@ set( common_SOURCES
IF( BUILD_CUDA )
set( tnl_containers_CUDA__SOURCES
${common_SOURCES}
- ${CURRENT_DIR}/ArrayOperationsHost_impl.cu
- ${CURRENT_DIR}/ArrayOperationsCuda_impl.cu
${CURRENT_DIR}/Array_impl.cu
${CURRENT_DIR}/SharedArray_impl.cu
${CURRENT_DIR}/MultiArray_impl.cu
${CURRENT_DIR}/StaticArray_impl.cu
- ${CURRENT_DIR}/VectorOperationsCuda_impl.cu
${CURRENT_DIR}/Vector_impl.cu
${CURRENT_DIR}/StaticVector_impl.cu
PARENT_SCOPE )
-ELSE()
- set( common_SOURCES
- ${common_SOURCES}
- ${CURRENT_DIR}/ArrayOperationsHost_impl.cpp
- ${CURRENT_DIR}/ArrayOperationsCuda_impl.cpp
- )
ENDIF()
set( tnl_containers_SOURCES
${common_SOURCES}
PARENT_SCOPE )
-INSTALL( FILES ${headers} DESTINATION include/tnl-${tnlVersion}/TNL/Containers )
\ No newline at end of file
+INSTALL( FILES ${headers} DESTINATION include/tnl-${tnlVersion}/TNL/Containers )
diff --git a/src/TNL/Containers/ConstSharedArray_impl.h b/src/TNL/Containers/ConstSharedArray_impl.h
index d6a036459116503fd74f0ca505b2b9ca0ea42e44..8892c0a346461a38d1c49a01177cae5f46dd434a 100644
--- a/src/TNL/Containers/ConstSharedArray_impl.h
+++ b/src/TNL/Containers/ConstSharedArray_impl.h
@@ -13,7 +13,7 @@
#include <iostream>
#include <TNL/File.h>
#include <TNL/Containers/Array.h>
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Math.h>
#include <TNL/param-types.h>
@@ -134,7 +134,7 @@ Element tnlConstSharedArray< Element, Device, Index > :: getElement( Index i ) c
std::cerr << "Wrong index for getElement method in tnlConstSharedArray with name "
<< " index is " << i
<< " and array size is " << this->getSize() );
- return ArrayOperations< Device >::getMemoryElement( &( this->data[ i ] ) );
+ return Algorithms::ArrayOperations< Device >::getMemoryElement( &( this->data[ i ] ) );
};
template< typename Element,
@@ -148,7 +148,7 @@ const Element& tnlConstSharedArray< Element, Device, Index > :: operator[] ( Ind
<< " index is " << i
<< " and array size is " << this->getSize() );
// TODO: add static assert - this does not make sense for Devices::CudaDevice
- return ArrayOperations< Device >::getArrayElementReference( this->data, i );
+ return Algorithms::ArrayOperations< Device >::getArrayElementReference( this->data, i );
};
template< typename Element,
@@ -179,14 +179,13 @@ bool tnlConstSharedArray< Element, Device, Index > :: operator == ( const Array&
{
if( array. getSize() != this->getSize() )
return false;
- return ArrayOperations< Device,
- typename Array :: DeviceType > ::
- template compareMemory< typename Array :: ElementType,
- Element,
- typename Array :: IndexType >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ return Algorithms::ArrayOperations< Device, typename Array :: DeviceType >::
+ template compareMemory< typename Array :: ElementType,
+ Element,
+ typename Array :: IndexType >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
}
template< typename Element,
diff --git a/src/TNL/List.h b/src/TNL/Containers/List.h
similarity index 84%
rename from src/TNL/List.h
rename to src/TNL/Containers/List.h
index 9d0f2cdb1a12a58087863576dfcfd1cbd793d117..3692a9a2ca56404151110eadb96acb211345f34a 100644
--- a/src/TNL/List.h
+++ b/src/TNL/Containers/List.h
@@ -10,16 +10,18 @@
#pragma once
-#include <TNL/Assert.h>
#include <stdlib.h>
#include <iostream>
+
+#include <TNL/Assert.h>
+#include <TNL/File.h>
#include <TNL/String.h>
#include <TNL/param-types.h>
namespace TNL {
+namespace Containers {
-class File;
-template< class T > class DataElement;
+template< class T > class ListDataElement;
//! Template for double linked lists
/*! To acces elements in the list one can use method getSize() and
@@ -114,18 +116,18 @@ template< class T > class List
protected:
//! Pointer to the first element
- DataElement< T >* first;
+ ListDataElement< T >* first;
//! Pointer to the last element
/*! We use pointer to last element while adding new element to keep order of elements
*/
- DataElement< T >* last;
+ ListDataElement< T >* last;
//! List size
int size;
//! Iterator
- mutable DataElement< T >* iterator;
+ mutable ListDataElement< T >* iterator;
//! Iterator index
mutable int index;
@@ -136,33 +138,33 @@ template< class T > class List
template< typename T > std::ostream& operator << ( std::ostream& str, const List< T >& list );
//! Data element for List and mStack
-template< class T > class DataElement
+template< class T > class ListDataElement
{
//! Main data
T data;
//! Pointer to the next element
- DataElement< T >* next;
+ ListDataElement< T >* next;
//! Pointer to the previous element
- DataElement< T >* previous;
+ ListDataElement< T >* previous;
public:
//! Basic constructor
- DataElement()
+ ListDataElement()
: next( 0 ),
previous( 0 ){};
//! Constructor with given data and possibly pointer to next element
- DataElement( const T& dt,
- DataElement< T >* prv = 0,
- DataElement< T >* nxt = 0 )
+ ListDataElement( const T& dt,
+ ListDataElement< T >* prv = 0,
+ ListDataElement< T >* nxt = 0 )
: data( dt ),
next( nxt ),
previous( prv ){};
//! Destructor
- ~DataElement(){};
+ ~ListDataElement(){};
//! Return data for non-const instances
T& Data() { return data; };
@@ -171,19 +173,19 @@ template< class T > class DataElement
const T& Data() const { return data; };
//! Return pointer to the next element for non-const instances
- DataElement< T >*& Next() { return next; };
+ ListDataElement< T >*& Next() { return next; };
//! Return pointer to the next element for const instances
- const DataElement< T >* Next() const { return next; };
+ const ListDataElement< T >* Next() const { return next; };
//! Return pointer to the previous element for non-const instances
- DataElement< T >*& Previous() { return previous; };
+ ListDataElement< T >*& Previous() { return previous; };
//! Return pointer to the previous element for const instances
- const DataElement< T >* Previous() const { return previous; };
-
+ const ListDataElement< T >* Previous() const { return previous; };
};
+} // namespace Containers
} // namespace TNL
-#include <TNL/List_impl.h>
+#include <TNL/Containers/List_impl.h>
diff --git a/src/TNL/List_impl.h b/src/TNL/Containers/List_impl.h
similarity index 93%
rename from src/TNL/List_impl.h
rename to src/TNL/Containers/List_impl.h
index 389fb1732b63cea1cd95336a7a9ac5bf9f80dc5f..30634a8104c95b54df5add8250ea8cc1fb97cc68 100644
--- a/src/TNL/List_impl.h
+++ b/src/TNL/Containers/List_impl.h
@@ -10,9 +10,10 @@
#pragma once
-#include <TNL/File.h>
+#include <TNL/Containers/List.h>
namespace TNL {
+namespace Containers {
template< typename T >
List< T >::List()
@@ -111,12 +112,12 @@ bool List< T >::Append( const T& data )
if( ! first )
{
Assert( ! last, );
- first = last = new DataElement< T >( data );
+ first = last = new ListDataElement< T >( data );
if( ! first ) return false;
}
else
{
- DataElement< T >* new_element = new DataElement< T >( data, last, 0 );
+ ListDataElement< T >* new_element = new ListDataElement< T >( data, last, 0 );
if( ! new_element ) return false;
Assert( last, );
last = last -> Next() = new_element;
@@ -131,12 +132,12 @@ bool List< T >::Prepend( const T& data )
if( ! first )
{
Assert( ! last, );
- first = last = new DataElement< T >( data );
+ first = last = new ListDataElement< T >( data );
if( ! first ) return false;
}
else
{
- DataElement< T >* new_element = new DataElement< T >( data, 0, first );
+ ListDataElement< T >* new_element = new ListDataElement< T >( data, 0, first );
if( ! new_element ) return false;
first = first -> Previous() = new_element;
}
@@ -152,8 +153,8 @@ bool List< T >::Insert( const T& data, const int& ind )
if( ind == 0 ) return Prepend( data );
if( ind == size ) return Append( data );
operator[]( ind );
- DataElement< T >* new_el =
- new DataElement< T >( data,
+ ListDataElement< T >* new_el =
+ new ListDataElement< T >( data,
iterator -> Previous(),
iterator );
if( ! new_el ) return false;
@@ -200,7 +201,7 @@ template< typename T >
void List< T >::Erase( const int& ind )
{
operator[]( ind );
- DataElement< T >* tmp_it = iterator;
+ ListDataElement< T >* tmp_it = iterator;
if( iterator -> Next() )
iterator -> Next() -> Previous() = iterator -> Previous();
if( iterator -> Previous() )
@@ -229,7 +230,7 @@ template< typename T >
void List< T >::reset()
{
iterator = first;
- DataElement< T >* tmp_it;
+ ListDataElement< T >* tmp_it;
while( iterator )
{
Assert( iterator, );
@@ -245,7 +246,7 @@ template< typename T >
void List< T >::DeepEraseAll()
{
iterator = first;
- DataElement< T >* tmp_it;
+ ListDataElement< T >* tmp_it;
int i( 0 );
while( iterator )
{
@@ -381,6 +382,5 @@ std::ostream& operator << ( std::ostream& str, const List< T >& list )
return str;
};
+} // namespace Containers
} // namespace TNL
-
-
diff --git a/src/TNL/Containers/SharedArray_impl.h b/src/TNL/Containers/SharedArray_impl.h
index 60eeec9359382d5035f62cdb227bf3ab2a3fd273..e60929af5b3f0917a17987472c64e73df213f62f 100644
--- a/src/TNL/Containers/SharedArray_impl.h
+++ b/src/TNL/Containers/SharedArray_impl.h
@@ -14,7 +14,7 @@
#include <TNL/File.h>
#include <TNL/Containers/Array.h>
#include <TNL/Containers/StaticArray.h>
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Math.h>
#include <TNL/param-types.h>
@@ -191,7 +191,7 @@ void SharedArray< Element, Device, Index > :: setElement( const Index& i, const
std::cerr << "Wrong index for setElement method in SharedArray "
<< " index is " << i
<< " and array size is " << this->getSize() );
- return ArrayOperations< Device >::setMemoryElement( & ( this->data[ i ] ), x );
+ return Algorithms::ArrayOperations< Device >::setMemoryElement( & ( this->data[ i ] ), x );
};
template< typename Element,
@@ -203,7 +203,7 @@ Element SharedArray< Element, Device, Index > :: getElement( const Index& i ) co
std::cerr << "Wrong index for getElement method in SharedArray "
<< " index is " << i
<< " and array size is " << this->getSize() );
- return ArrayOperations< Device >::getMemoryElement( &( this->data[ i ] ) );
+ return Algorithms::ArrayOperations< Device >::getMemoryElement( &( this->data[ i ] ) );
};
template< typename Element,
@@ -241,13 +241,13 @@ SharedArray< Element, Device, Index >&
Assert( array. getSize() == this->getSize(),
std::cerr << "Source size: " << array. getSize() << std::endl
<< "Target size: " << this->getSize() << std::endl );
- ArrayOperations< Device > ::
- template copyMemory< Element,
- Element,
- Index >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ Algorithms::ArrayOperations< Device > ::
+ template copyMemory< Element,
+ Element,
+ Index >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
return ( *this );
};
@@ -260,14 +260,13 @@ SharedArray< Element, Device, Index >& SharedArray< Element, Device, Index > ::
Assert( array. getSize() == this->getSize(),
std::cerr << "Source size: " << array. getSize() << std::endl
<< "Target size: " << this->getSize() << std::endl );
- ArrayOperations< typename Array :: DeviceType,
- Device > ::
- template copyMemory< Element,
- typename Array :: ElementType,
- typename Array :: IndexType >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ Algorithms::ArrayOperations< typename Array::DeviceType, Device >::
+ template copyMemory< Element,
+ typename Array :: ElementType,
+ typename Array :: IndexType >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
return ( *this );
};
@@ -279,14 +278,13 @@ bool SharedArray< Element, Device, Index > :: operator == ( const Array& array )
{
if( array. getSize() != this->getSize() )
return false;
- return ArrayOperations< Device,
- typename Array :: DeviceType > ::
- template compareMemory< typename Array :: ElementType,
- Element,
- typename Array :: IndexType >
- ( this->getData(),
- array. getData(),
- array. getSize() );
+ return Algorithms::ArrayOperations< Device, typename Array::DeviceType >::
+ template compareMemory< typename Array :: ElementType,
+ Element,
+ typename Array :: IndexType >
+ ( this->getData(),
+ array. getData(),
+ array. getSize() );
}
template< typename Element,
@@ -304,7 +302,7 @@ template< typename Element,
void SharedArray< Element, Device, Index > :: setValue( const Element& e )
{
Assert( this->size != 0, );
- ArrayOperations< Device >::template setMemory< Element, Index >
+ Algorithms::ArrayOperations< Device >::template setMemory< Element, Index >
( this->getData(), e, this->getSize() );
}
diff --git a/src/TNL/Containers/SharedVector_impl.h b/src/TNL/Containers/SharedVector_impl.h
index d086bfc6426a2d6e6ceffda40c73beeb67ad7613..869f0aeae455b9b3891d9147f8d3ea4285503a02 100644
--- a/src/TNL/Containers/SharedVector_impl.h
+++ b/src/TNL/Containers/SharedVector_impl.h
@@ -10,7 +10,8 @@
#pragma once
-#include <TNL/Containers/VectorOperations.h>
+#include <TNL/Containers/SharedVector.h>
+#include <TNL/Containers/Algorithms/VectorOperations.h>
namespace TNL {
namespace Containers {
@@ -92,7 +93,7 @@ template< typename Real,
void SharedVector< Real, Device, Index >::addElement( const IndexType i,
const RealType& value )
{
- VectorOperations< Device >::addElement( *this, i, value );
+ Algorithms::VectorOperations< Device >::addElement( *this, i, value );
}
template< typename Real,
@@ -102,7 +103,7 @@ void SharedVector< Real, Device, Index >::addElement( const IndexType i,
const RealType& value,
const RealType& thisElementMultiplicator )
{
- VectorOperations< Device >::addElement( *this, i, value, thisElementMultiplicator );
+ Algorithms::VectorOperations< Device >::addElement( *this, i, value, thisElementMultiplicator );
}
template< typename Real,
@@ -169,7 +170,7 @@ template< typename Real,
typename Index >
SharedVector< Real, Device, Index >& SharedVector< Real, Device, Index > :: operator *= ( const RealType& c )
{
- VectorOperations< Device >::vectorScalarMultiplication( *this, c );
+ Algorithms::VectorOperations< Device >::vectorScalarMultiplication( *this, c );
return *this;
}
@@ -178,7 +179,7 @@ template< typename Real,
typename Index >
SharedVector< Real, Device, Index >& SharedVector< Real, Device, Index > :: operator /= ( const RealType& c )
{
- VectorOperations< Device >::vectorScalarMultiplication( *this, 1.0/ c );
+ Algorithms::VectorOperations< Device >::vectorScalarMultiplication( *this, 1.0/ c );
return *this;
}
@@ -187,7 +188,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: max() const
{
- return VectorOperations< Device > :: getVectorMax( *this );
+ return Algorithms::VectorOperations< Device > :: getVectorMax( *this );
}
template< typename Real,
@@ -195,7 +196,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: min() const
{
- return VectorOperations< Device > :: getVectorMin( *this );
+ return Algorithms::VectorOperations< Device > :: getVectorMin( *this );
}
@@ -204,7 +205,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: absMax() const
{
- return VectorOperations< Device > :: getVectorAbsMax( *this );
+ return Algorithms::VectorOperations< Device > :: getVectorAbsMax( *this );
}
template< typename Real,
@@ -212,7 +213,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: absMin() const
{
- return VectorOperations< Device > :: getVectorAbsMin( *this );
+ return Algorithms::VectorOperations< Device > :: getVectorAbsMin( *this );
}
template< typename Real,
@@ -220,7 +221,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: lpNorm( const Real& p ) const
{
- return VectorOperations< Device > :: getVectorLpNorm( *this, p );
+ return Algorithms::VectorOperations< Device > :: getVectorLpNorm( *this, p );
}
@@ -229,7 +230,7 @@ template< typename Real,
typename Index >
Real SharedVector< Real, Device, Index > :: sum() const
{
- return VectorOperations< Device > :: getVectorSum( *this );
+ return Algorithms::VectorOperations< Device > :: getVectorSum( *this );
}
@@ -239,7 +240,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceMax( const Vector& v ) const
{
- return VectorOperations< Device > :: getVectorDifferenceMax( *this, v );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceMax( *this, v );
}
@@ -249,7 +250,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceMin( const Vector& v ) const
{
- return VectorOperations< Device > :: getVectorDifferenceMin( *this, v );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceMin( *this, v );
}
@@ -259,7 +260,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceAbsMax( const Vector& v ) const
{
- return VectorOperations< Device > :: getVectorDifferenceAbsMax( *this, v );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceAbsMax( *this, v );
}
template< typename Real,
@@ -268,7 +269,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceAbsMin( const Vector& v ) const
{
- return VectorOperations< Device > :: getVectorDifferenceAbsMin( *this, v );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceAbsMin( *this, v );
}
template< typename Real,
@@ -277,7 +278,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceLpNorm( const Vector& v, const Real& p ) const
{
- return VectorOperations< Device > :: getVectorDifferenceLpNorm( *this, v, p );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceLpNorm( *this, v, p );
}
@@ -287,7 +288,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: differenceSum( const Vector& v ) const
{
- return VectorOperations< Device > :: getVectorDifferenceSum( *this, v );
+ return Algorithms::VectorOperations< Device > :: getVectorDifferenceSum( *this, v );
}
@@ -296,7 +297,7 @@ template< typename Real,
typename Index >
void SharedVector< Real, Device, Index > :: scalarMultiplication( const Real& alpha )
{
- VectorOperations< Device > :: vectorScalarMultiplication( *this, alpha );
+ Algorithms::VectorOperations< Device > :: vectorScalarMultiplication( *this, alpha );
}
@@ -306,7 +307,7 @@ template< typename Real,
template< typename Vector >
Real SharedVector< Real, Device, Index > :: scalarProduct( const Vector& v )
{
- return VectorOperations< Device > :: getScalarProduct( *this, v );
+ return Algorithms::VectorOperations< Device > :: getScalarProduct( *this, v );
}
template< typename Real,
@@ -317,7 +318,7 @@ void SharedVector< Real, Device, Index > :: addVector( const Vector& x,
const Real& alpha,
const Real& thisMultiplicator )
{
- VectorOperations< Device > :: addVector( *this, x, alpha, thisMultiplicator );
+ Algorithms::VectorOperations< Device > :: addVector( *this, x, alpha, thisMultiplicator );
}
template< typename Real,
@@ -332,7 +333,7 @@ addVectors( const Vector& v1,
const Real& multiplicator2,
const Real& thisMultiplicator )
{
- VectorOperations< Device >::addVectors( *this, v1, multiplicator1, v2, multiplicator2, thisMultiplicator );
+ Algorithms::VectorOperations< Device >::addVectors( *this, v1, multiplicator1, v2, multiplicator2, thisMultiplicator );
}
template< typename Real,
@@ -340,7 +341,7 @@ template< typename Real,
typename Index >
void SharedVector< Real, Device, Index > :: computePrefixSum()
{
- VectorOperations< Device >::computePrefixSum( *this, 0, this->getSize() );
+ Algorithms::VectorOperations< Device >::computePrefixSum( *this, 0, this->getSize() );
}
template< typename Real,
@@ -349,7 +350,7 @@ template< typename Real,
void SharedVector< Real, Device, Index > :: computePrefixSum( const IndexType begin,
const IndexType end )
{
- VectorOperations< Device >::computePrefixSum( *this, begin, end );
+ Algorithms::VectorOperations< Device >::computePrefixSum( *this, begin, end );
}
template< typename Real,
@@ -357,7 +358,7 @@ template< typename Real,
typename Index >
void SharedVector< Real, Device, Index > :: computeExclusivePrefixSum()
{
- VectorOperations< Device >::computeExclusivePrefixSum( *this, 0, this->getSize() );
+ Algorithms::VectorOperations< Device >::computeExclusivePrefixSum( *this, 0, this->getSize() );
}
template< typename Real,
@@ -366,7 +367,7 @@ template< typename Real,
void SharedVector< Real, Device, Index > :: computeExclusivePrefixSum( const IndexType begin,
const IndexType end )
{
- VectorOperations< Device >::computeExclusivePrefixSum( *this, begin, end );
+ Algorithms::VectorOperations< Device >::computeExclusivePrefixSum( *this, begin, end );
}
#ifdef TEMPLATE_EXPLICIT_INSTANTIATION
diff --git a/src/TNL/Containers/StaticVector_impl.h b/src/TNL/Containers/StaticVector_impl.h
index 251c1e7369e6c20e141ef3fa2fb675d217be4654..72b76bd29d62aa7881d237c5f3a80d5ac1be3503 100644
--- a/src/TNL/Containers/StaticVector_impl.h
+++ b/src/TNL/Containers/StaticVector_impl.h
@@ -181,8 +181,8 @@ StaticVector< Size, Real >::abs() const
}
-template< int Size, typename Real >
-StaticVector< Size, Real > operator * ( const Real& c, const StaticVector< Size, Real >& u )
+template< int Size, typename Real, typename Scalar >
+StaticVector< Size, Real > operator * ( const Scalar& c, const StaticVector< Size, Real >& u )
{
return u * c;
}
diff --git a/src/TNL/Containers/Vector.h b/src/TNL/Containers/Vector.h
index 0a73dd6988a0ecf687f0a2d6194d329b0cf7b7b9..01a6a3682f5d7779ec227732e0d682c5ef454109 100644
--- a/src/TNL/Containers/Vector.h
+++ b/src/TNL/Containers/Vector.h
@@ -11,22 +11,13 @@
#pragma once
#include <TNL/Containers/Array.h>
-#include <TNL/Functions/Domain.h>
namespace TNL {
-
-namespace Devices
-{
- class Host;
-}
-
namespace Containers {
-
-
template< typename Real = double,
- typename Device = Devices::Host,
- typename Index = int >
+ typename Device = Devices::Host,
+ typename Index = int >
class Vector : public Containers::Array< Real, Device, Index >
{
public:
diff --git a/src/TNL/Containers/Vector_impl.h b/src/TNL/Containers/Vector_impl.h
index b602d003204f3de8417c323f90bfe50ac71b06fc..b5a233faa6be107c2f7f776b7a66f58c998eaeae 100644
--- a/src/TNL/Containers/Vector_impl.h
+++ b/src/TNL/Containers/Vector_impl.h
@@ -10,7 +10,8 @@
#pragma once
-#include <TNL/Containers/VectorOperations.h>
+#include <TNL/Containers/Vector.h>
+#include <TNL/Containers/Algorithms/VectorOperations.h>
namespace TNL {
namespace Containers {
@@ -18,15 +19,16 @@ namespace Containers {
template< typename Real,
typename Device,
typename Index >
-Vector< Real, Device, Index >::Vector()
+Vector< Real, Device, Index >::
+Vector()
{
-
}
template< typename Real,
typename Device,
typename Index >
-Vector< Real, Device, Index >::Vector( const Index size )
+Vector< Real, Device, Index >::
+Vector( const Index size )
{
this->setSize( size );
}
@@ -35,18 +37,22 @@ Vector< Real, Device, Index >::Vector( const Index size )
template< typename Real,
typename Device,
typename Index >
-String Vector< Real, Device, Index >::getType()
+String
+Vector< Real, Device, Index >::
+getType()
{
return String( "Containers::Vector< " ) +
- TNL::getType< Real >() + ", " +
- Device::getDeviceType() + ", " +
- TNL::getType< Index >() + " >";
+ TNL::getType< Real >() + ", " +
+ Device::getDeviceType() + ", " +
+ TNL::getType< Index >() + " >";
};
template< typename Real,
typename Device,
typename Index >
-String Vector< Real, Device, Index >::getTypeVirtual() const
+String
+Vector< Real, Device, Index >::
+getTypeVirtual() const
{
return this->getType();
};
@@ -54,7 +60,9 @@ String Vector< Real, Device, Index >::getTypeVirtual() const
template< typename Real,
typename Device,
typename Index >
-String Vector< Real, Device, Index >::getSerializationType()
+String
+Vector< Real, Device, Index >::
+getSerializationType()
{
return HostType::getType();
};
@@ -62,7 +70,9 @@ String Vector< Real, Device, Index >::getSerializationType()
template< typename Real,
typename Device,
typename Index >
-String Vector< Real, Device, Index >::getSerializationTypeVirtual() const
+String
+Vector< Real, Device, Index >::
+getSerializationTypeVirtual() const
{
return this->getSerializationType();
};
@@ -70,27 +80,32 @@ String Vector< Real, Device, Index >::getSerializationTypeVirtual() const
template< typename Real,
typename Device,
typename Index >
-void Vector< Real, Device, Index >::addElement( const IndexType i,
- const RealType& value )
+void
+Vector< Real, Device, Index >::
+addElement( const IndexType i,
+ const RealType& value )
{
- VectorOperations< Device >::addElement( *this, i, value );
+ Algorithms::VectorOperations< Device >::addElement( *this, i, value );
}
template< typename Real,
typename Device,
typename Index >
-void Vector< Real, Device, Index >::addElement( const IndexType i,
- const RealType& value,
- const RealType& thisElementMultiplicator )
+void
+Vector< Real, Device, Index >::
+addElement( const IndexType i,
+ const RealType& value,
+ const RealType& thisElementMultiplicator )
{
- VectorOperations< Device >::addElement( *this, i, value, thisElementMultiplicator );
+ Algorithms::VectorOperations< Device >::addElement( *this, i, value, thisElementMultiplicator );
}
template< typename Real,
- typename Device,
- typename Index >
+ typename Device,
+ typename Index >
Vector< Real, Device, Index >&
- Vector< Real, Device, Index >::operator = ( const Vector< Real, Device, Index >& vector )
+Vector< Real, Device, Index >::
+operator = ( const Vector< Real, Device, Index >& vector )
{
Containers::Array< Real, Device, Index >::operator = ( vector );
return ( *this );
@@ -101,7 +116,8 @@ template< typename Real,
typename Index >
template< typename VectorT >
Vector< Real, Device, Index >&
- Vector< Real, Device, Index >::operator = ( const VectorT& vector )
+Vector< Real, Device, Index >::
+operator = ( const VectorT& vector )
{
Containers::Array< Real, Device, Index >::operator = ( vector );
return ( *this );
@@ -111,7 +127,9 @@ template< typename Real,
typename Device,
typename Index >
template< typename VectorT >
-bool Vector< Real, Device, Index >::operator == ( const VectorT& vector ) const
+bool
+Vector< Real, Device, Index >::
+operator == ( const VectorT& vector ) const
{
return Containers::Array< Real, Device, Index >::operator == ( vector );
}
@@ -120,7 +138,9 @@ template< typename Real,
typename Device,
typename Index >
template< typename VectorT >
-bool Vector< Real, Device, Index >::operator != ( const VectorT& vector ) const
+bool
+Vector< Real, Device, Index >::
+operator != ( const VectorT& vector ) const
{
return Containers::Array< Real, Device, Index >::operator != ( vector );
}
@@ -129,7 +149,9 @@ template< typename Real,
typename Device,
typename Index >
template< typename VectorT >
-Vector< Real, Device, Index >& Vector< Real, Device, Index >::operator -= ( const VectorT& vector )
+Vector< Real, Device, Index >&
+Vector< Real, Device, Index >::
+operator -= ( const VectorT& vector )
{
this->addVector( vector, -1.0 );
return *this;
@@ -139,7 +161,9 @@ template< typename Real,
typename Device,
typename Index >
template< typename VectorT >
-Vector< Real, Device, Index >& Vector< Real, Device, Index >::operator += ( const VectorT& vector )
+Vector< Real, Device, Index >&
+Vector< Real, Device, Index >::
+operator += ( const VectorT& vector )
{
this->addVector( vector );
return *this;
@@ -148,18 +172,22 @@ Vector< Real, Device, Index >& Vector< Real, Device, Index >::operator += ( cons
template< typename Real,
typename Device,
typename Index >
-Vector< Real, Device, Index >& Vector< Real, Device, Index >::operator *= ( const RealType& c )
+Vector< Real, Device, Index >&
+Vector< Real, Device, Index >::
+operator *= ( const RealType& c )
{
- VectorOperations< Device >::vectorScalarMultiplication( *this, c );
+ Algorithms::VectorOperations< Device >::vectorScalarMultiplication( *this, c );
return *this;
}
template< typename Real,
typename Device,
typename Index >
-Vector< Real, Device, Index >& Vector< Real, Device, Index >::operator /= ( const RealType& c )
+Vector< Real, Device, Index >&
+Vector< Real, Device, Index >::
+operator /= ( const RealType& c )
{
- VectorOperations< Device >::vectorScalarMultiplication( *this, 1.0 / c );
+ Algorithms::VectorOperations< Device >::vectorScalarMultiplication( *this, 1.0 / c );
return *this;
}
@@ -169,7 +197,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::max() const
{
- return VectorOperations< Device >::getVectorMax( *this );
+ return Algorithms::VectorOperations< Device >::getVectorMax( *this );
}
template< typename Real,
@@ -177,7 +205,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::min() const
{
- return VectorOperations< Device >::getVectorMin( *this );
+ return Algorithms::VectorOperations< Device >::getVectorMin( *this );
}
@@ -186,7 +214,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::absMax() const
{
- return VectorOperations< Device >::getVectorAbsMax( *this );
+ return Algorithms::VectorOperations< Device >::getVectorAbsMax( *this );
}
template< typename Real,
@@ -194,7 +222,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::absMin() const
{
- return VectorOperations< Device >::getVectorAbsMin( *this );
+ return Algorithms::VectorOperations< Device >::getVectorAbsMin( *this );
}
template< typename Real,
@@ -202,7 +230,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::lpNorm( const Real& p ) const
{
- return VectorOperations< Device >::getVectorLpNorm( *this, p );
+ return Algorithms::VectorOperations< Device >::getVectorLpNorm( *this, p );
}
@@ -211,7 +239,7 @@ template< typename Real,
typename Index >
Real Vector< Real, Device, Index >::sum() const
{
- return VectorOperations< Device >::getVectorSum( *this );
+ return Algorithms::VectorOperations< Device >::getVectorSum( *this );
}
@@ -221,7 +249,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceMax( const VectorT& v ) const
{
- return VectorOperations< Device >::getVectorDifferenceMax( *this, v );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceMax( *this, v );
}
@@ -231,7 +259,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceMin( const VectorT& v ) const
{
- return VectorOperations< Device >::getVectorDifferenceMin( *this, v );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceMin( *this, v );
}
@@ -241,7 +269,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceAbsMax( const VectorT& v ) const
{
- return VectorOperations< Device >::getVectorDifferenceAbsMax( *this, v );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceAbsMax( *this, v );
}
template< typename Real,
@@ -250,7 +278,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceAbsMin( const VectorT& v ) const
{
- return VectorOperations< Device >::getVectorDifferenceAbsMin( *this, v );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceAbsMin( *this, v );
}
template< typename Real,
@@ -259,7 +287,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceLpNorm( const VectorT& v, const Real& p ) const
{
- return VectorOperations< Device >::getVectorDifferenceLpNorm( *this, v, p );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceLpNorm( *this, v, p );
}
@@ -269,7 +297,7 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::differenceSum( const VectorT& v ) const
{
- return VectorOperations< Device >::getVectorDifferenceSum( *this, v );
+ return Algorithms::VectorOperations< Device >::getVectorDifferenceSum( *this, v );
}
@@ -278,7 +306,7 @@ template< typename Real,
typename Index >
void Vector< Real, Device, Index >::scalarMultiplication( const Real& alpha )
{
- VectorOperations< Device >::vectorScalarMultiplication( *this, alpha );
+ Algorithms::VectorOperations< Device >::vectorScalarMultiplication( *this, alpha );
}
@@ -288,18 +316,20 @@ template< typename Real,
template< typename VectorT >
Real Vector< Real, Device, Index >::scalarProduct( const VectorT& v ) const
{
- return VectorOperations< Device >::getScalarProduct( *this, v );
+ return Algorithms::VectorOperations< Device >::getScalarProduct( *this, v );
}
template< typename Real,
typename Device,
typename Index >
template< typename VectorT >
-void Vector< Real, Device, Index >::addVector( const VectorT& x,
- const Real& multiplicator,
- const Real& thisMultiplicator )
+void
+Vector< Real, Device, Index >::
+addVector( const VectorT& x,
+ const Real& multiplicator,
+ const Real& thisMultiplicator )
{
- VectorOperations< Device >::addVector( *this, x, multiplicator, thisMultiplicator );
+ Algorithms::VectorOperations< Device >::addVector( *this, x, multiplicator, thisMultiplicator );
}
template< typename Real,
@@ -314,7 +344,7 @@ addVectors( const VectorT& v1,
const Real& multiplicator2,
const Real& thisMultiplicator )
{
- VectorOperations< Device >::addVectors( *this, v1, multiplicator1, v2, multiplicator2, thisMultiplicator );
+ Algorithms::VectorOperations< Device >::addVectors( *this, v1, multiplicator1, v2, multiplicator2, thisMultiplicator );
}
template< typename Real,
@@ -322,16 +352,18 @@ template< typename Real,
typename Index >
void Vector< Real, Device, Index >::computePrefixSum()
{
- VectorOperations< Device >::computePrefixSum( *this, 0, this->getSize() );
+ Algorithms::VectorOperations< Device >::computePrefixSum( *this, 0, this->getSize() );
}
template< typename Real,
typename Device,
typename Index >
-void Vector< Real, Device, Index >::computePrefixSum( const IndexType begin,
- const IndexType end )
+void
+Vector< Real, Device, Index >::
+computePrefixSum( const IndexType begin,
+ const IndexType end )
{
- VectorOperations< Device >::computePrefixSum( *this, begin, end );
+ Algorithms::VectorOperations< Device >::computePrefixSum( *this, begin, end );
}
template< typename Real,
@@ -339,16 +371,18 @@ template< typename Real,
typename Index >
void Vector< Real, Device, Index >::computeExclusivePrefixSum()
{
- VectorOperations< Device >::computeExclusivePrefixSum( *this, 0, this->getSize() );
+ Algorithms::VectorOperations< Device >::computeExclusivePrefixSum( *this, 0, this->getSize() );
}
template< typename Real,
typename Device,
typename Index >
-void Vector< Real, Device, Index >::computeExclusivePrefixSum( const IndexType begin,
- const IndexType end )
+void
+Vector< Real, Device, Index >::
+computeExclusivePrefixSum( const IndexType begin,
+ const IndexType end )
{
- VectorOperations< Device >::computeExclusivePrefixSum( *this, begin, end );
+ Algorithms::VectorOperations< Device >::computeExclusivePrefixSum( *this, begin, end );
}
diff --git a/src/TNL/Curve.h b/src/TNL/Curve.h
index ace8d622be4afa36cbde901b3d8847a358294f1f..6704b8a98865aabf36a4b5b8ce971c32c1802ff6 100644
--- a/src/TNL/Curve.h
+++ b/src/TNL/Curve.h
@@ -13,7 +13,7 @@
#include <iomanip>
#include <fstream>
#include <cstring>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/Object.h>
#include <TNL/Math.h>
#include <TNL/Containers/StaticVector.h>
@@ -22,7 +22,8 @@
namespace TNL {
//! Basic structure for curves
-template< class T > class CurveElement
+template< class T >
+class CurveElement
{
public:
CurveElement() {};
@@ -71,7 +72,10 @@ template< class T > class CurveElement
bool separator;
};
-template< class T > class Curve : public Object, public List< CurveElement< T > >
+template< class T >
+class Curve
+ : public Object,
+ public Containers::List< CurveElement< T > >
{
public:
//! Basic contructor
@@ -94,20 +98,20 @@ template< class T > class Curve : public Object, public List< CurveElement< T >
//! Append new point
void Append( const T& vec, bool separator = false )
{
- List< CurveElement< T > > :: Append( CurveElement< T >( vec, separator ) );
+ Containers::List< CurveElement< T > > :: Append( CurveElement< T >( vec, separator ) );
};
//! Erase the curve
void Erase()
{
- List< CurveElement< T > >::reset();
+ Containers::List< CurveElement< T > >::reset();
};
//! Method for saving the object to a file as a binary data
bool save( File& file ) const
{
if( ! Object :: save( file ) ) return false;
- if( ! List< CurveElement< T > > :: DeepSave( file ) ) return false;
+ if( ! Containers::List< CurveElement< T > > :: DeepSave( file ) ) return false;
return true;
};
@@ -115,7 +119,7 @@ template< class T > class Curve : public Object, public List< CurveElement< T >
bool load( File& file )
{
if( ! Object :: load( file ) ) return false;
- if( ! List< CurveElement< T > > :: DeepLoad( file ) ) return false;
+ if( ! Containers::List< CurveElement< T > > :: DeepLoad( file ) ) return false;
return true;
};
diff --git a/src/TNL/Devices/Cuda.cu b/src/TNL/Devices/Cuda.cu
index 3d4ad065ddfa472ef50ca79bd3cd6fe43e49ce1b..e12517f560efa58213369db7b41c07a4750a9d2a 100644
--- a/src/TNL/Devices/Cuda.cu
+++ b/src/TNL/Devices/Cuda.cu
@@ -57,368 +57,371 @@ bool Cuda::checkDevice( const char* file_name, int line, cudaError error )
// 1
case cudaErrorMissingConfiguration:
std::cerr
- << "The device function being invoked (usually via ::cudaLaunch()) was not " << std::endl
- << "previously configured via the ::cudaConfigureCall() function. " << std::endl;
- break;
+ << "The device function being invoked (usually via ::cudaLaunch()) was not " << std::endl
+ << "previously configured via the ::cudaConfigureCall() function. " << std::endl;
+ break;
// 2
case cudaErrorMemoryAllocation:
std::cerr
- << "The API call failed because it was unable to allocate enough memory to " << std::endl
- << "perform the requested operation. " << std::endl;
- break;
+ << "The API call failed because it was unable to allocate enough memory to " << std::endl
+ << "perform the requested operation. " << std::endl;
+ break;
// 3
case cudaErrorInitializationError:
std::cerr
- << "The API call failed because the CUDA driver and runtime could not be " << std::endl
- << "initialized. " << std::endl;
- break;
+ << "The API call failed because the CUDA driver and runtime could not be " << std::endl
+ << "initialized. " << std::endl;
+ break;
// 4
case cudaErrorLaunchFailure:
std::cerr
- << "An exception occurred on the device while executing a kernel. Common " << std::endl
- << "causes include dereferencing an invalid device pointer and accessing " << std::endl
- << "out of bounds shared memory. The device cannot be used until " << std::endl
- << "::cudaThreadExit() is called. All existing device memory allocations " << std::endl
- << "are invalid and must be reconstructed if the program is to continue " << std::endl
- << "using CUDA. " << std::endl;
- break;
+ << "An exception occurred on the device while executing a kernel. Common " << std::endl
+ << "causes include dereferencing an invalid device pointer and accessing " << std::endl
+ << "out of bounds shared memory. The device cannot be used until " << std::endl
+ << "::cudaThreadExit() is called. All existing device memory allocations " << std::endl
+ << "are invalid and must be reconstructed if the program is to continue " << std::endl
+ << "using CUDA. " << std::endl;
+ break;
// 5
case cudaErrorPriorLaunchFailure:
std::cerr
- << "This indicated that a previous kernel launch failed. This was previously " << std::endl
- << "used for device emulation of kernel launches. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "This indicated that a previous kernel launch failed. This was previously " << std::endl
+ << "used for device emulation of kernel launches. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
// 6
case cudaErrorLaunchTimeout:
std::cerr
- << "This indicates that the device kernel took too long to execute. This can " << std::endl
- << "only occur if timeouts are enabled - see the device property " << std::endl
- << "ref ::cudaDeviceProp::kernelExecTimeoutEnabled \"kernelExecTimeoutEnabled\" " << std::endl
- << "for more information. The device cannot be used until ::cudaThreadExit() " << std::endl
- << "is called. All existing device memory allocations are invalid and must be " << std::endl
- << "reconstructed if the program is to continue using CUDA. " << std::endl;
- break;
+ << "This indicates that the device kernel took too long to execute. This can " << std::endl
+ << "only occur if timeouts are enabled - see the device property " << std::endl
+ << "ref ::cudaDeviceProp::kernelExecTimeoutEnabled \"kernelExecTimeoutEnabled\" " << std::endl
+ << "for more information. The device cannot be used until ::cudaThreadExit() " << std::endl
+ << "is called. All existing device memory allocations are invalid and must be " << std::endl
+ << "reconstructed if the program is to continue using CUDA. " << std::endl;
+ break;
// 7
case cudaErrorLaunchOutOfResources:
std::cerr
- << "This indicates that a launch did not occur because it did not have " << std::endl
- << "appropriate resources. Although this error is similar to " << std::endl
- << "::cudaErrorInvalidConfiguration, this error usually indicates that the " << std::endl
- << "user has attempted to pass too many arguments to the device kernel, or the " << std::endl
- << "kernel launch specifies too many threads for the kernel's register count. " << std::endl;
- break;
+ << "This indicates that a launch did not occur because it did not have " << std::endl
+ << "appropriate resources. Although this error is similar to " << std::endl
+ << "::cudaErrorInvalidConfiguration, this error usually indicates that the " << std::endl
+ << "user has attempted to pass too many arguments to the device kernel, or the " << std::endl
+ << "kernel launch specifies too many threads for the kernel's register count. " << std::endl;
+ break;
// 8
case cudaErrorInvalidDeviceFunction:
std::cerr
- << "The requested device function does not exist or is not compiled for the " << std::endl
- << "proper device architecture. " << std::endl;
- break;
+ << "The requested device function does not exist or is not compiled for the " << std::endl
+ << "proper device architecture. " << std::endl;
+ break;
// 9
case cudaErrorInvalidConfiguration:
std::cerr
- << "This indicates that a kernel launch is requesting resources that can " << std::endl
- << "never be satisfied by the current device. Requesting more shared memory " << std::endl
- << "per block than the device supports will trigger this error, as will " << std::endl
- << "requesting too many threads or blocks. See ::cudaDeviceProp for more " << std::endl
- << "device limitations. " << std::endl;
- break;
+ << "This indicates that a kernel launch is requesting resources that can " << std::endl
+ << "never be satisfied by the current device. Requesting more shared memory " << std::endl
+ << "per block than the device supports will trigger this error, as will " << std::endl
+ << "requesting too many threads or blocks. See ::cudaDeviceProp for more " << std::endl
+ << "device limitations. " << std::endl;
+ break;
// 10
case cudaErrorInvalidDevice:
std::cerr
- << "This indicates that the device ordinal supplied by the user does not " << std::endl
- << "correspond to a valid CUDA device. " << std::endl;
- break;
+ << "This indicates that the device ordinal supplied by the user does not " << std::endl
+ << "correspond to a valid CUDA device. " << std::endl;
+ break;
// 11
case cudaErrorInvalidValue:
std::cerr
- << "This indicates that one or more of the parameters passed to the API call " << std::endl
- << "is not within an acceptable range of values. " << std::endl;
- break;
+ << "This indicates that one or more of the parameters passed to the API call " << std::endl
+ << "is not within an acceptable range of values. " << std::endl;
+ break;
// 12
case cudaErrorInvalidPitchValue:
std::cerr
- << "This indicates that one or more of the pitch-related parameters passed " << std::endl
- << "to the API call is not within the acceptable range for pitch. " << std::endl;
- break;
+ << "This indicates that one or more of the pitch-related parameters passed " << std::endl
+ << "to the API call is not within the acceptable range for pitch. " << std::endl;
+ break;
// 13
case cudaErrorInvalidSymbol:
std::cerr
- << "This indicates that the symbol name/identifier passed to the API call " << std::endl
- << "is not a valid name or identifier. " << std::endl;
- break;
+ << "This indicates that the symbol name/identifier passed to the API call " << std::endl
+ << "is not a valid name or identifier. " << std::endl;
+ break;
// 14
case cudaErrorMapBufferObjectFailed:
- std::cerr
- << "This indicates that the buffer object could not be mapped. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that the buffer object could not be mapped. " << std::endl;
+ break;
// 15
case cudaErrorUnmapBufferObjectFailed:
std::cerr
- << "This indicates that the buffer object could not be unmapped. " << std::endl;
- break;
+ << "This indicates that the buffer object could not be unmapped. " << std::endl;
+ break;
// 16
case cudaErrorInvalidHostPointer:
std::cerr
- << "This indicates that at least one host pointer passed to the API call is " << std::endl
- << "not a valid host pointer. " << std::endl;
- break;
+ << "This indicates that at least one host pointer passed to the API call is " << std::endl
+ << "not a valid host pointer. " << std::endl;
+ break;
// 17
case cudaErrorInvalidDevicePointer:
std::cerr
- << "This indicates that at least one device pointer passed to the API call is " << std::endl
- << "not a valid device pointer. " << std::endl;
- break;
+ << "This indicates that at least one device pointer passed to the API call is " << std::endl
+ << "not a valid device pointer. " << std::endl;
+ break;
case cudaErrorInvalidTexture:
std::cerr
- << "This indicates that the texture passed to the API call is not a valid " << std::endl
- << "texture. " << std::endl;
- break;
+ << "This indicates that the texture passed to the API call is not a valid " << std::endl
+ << "texture. " << std::endl;
+ break;
case cudaErrorInvalidTextureBinding:
std::cerr
- << "This indicates that the texture binding is not valid. This occurs if you " << std::endl
- << "call ::cudaGetTextureAlignmentOffset() with an unbound texture. " << std::endl;
- break;
+ << "This indicates that the texture binding is not valid. This occurs if you " << std::endl
+ << "call ::cudaGetTextureAlignmentOffset() with an unbound texture. " << std::endl;
+ break;
case cudaErrorInvalidChannelDescriptor:
std::cerr
- << "This indicates that the channel descriptor passed to the API call is not " << std::endl
- << "valid. This occurs if the format is not one of the formats specified by " << std::endl
- << "::cudaChannelFormatKind, or if one of the dimensions is invalid. " << std::endl;
- break;
+ << "This indicates that the channel descriptor passed to the API call is not " << std::endl
+ << "valid. This occurs if the format is not one of the formats specified by " << std::endl
+ << "::cudaChannelFormatKind, or if one of the dimensions is invalid. " << std::endl;
+ break;
case cudaErrorInvalidMemcpyDirection:
std::cerr
- << "This indicates that the direction of the memcpy passed to the API call is " << std::endl
- << "not one of the types specified by ::cudaMemcpyKind. " << std::endl;
- break;
+ << "This indicates that the direction of the memcpy passed to the API call is " << std::endl
+ << "not one of the types specified by ::cudaMemcpyKind. " << std::endl;
+ break;
case cudaErrorAddressOfConstant:
std::cerr
- << "This indicated that the user has taken the address of a constant variable, " << std::endl
- << "which was forbidden up until the CUDA 3.1 release. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Variables in constant " << std::endl
- << "memory may now have their address taken by the runtime via " << std::endl
- << "::cudaGetSymbolAddress(). " << std::endl;
- break;
+ << "This indicated that the user has taken the address of a constant variable, " << std::endl
+ << "which was forbidden up until the CUDA 3.1 release. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Variables in constant " << std::endl
+ << "memory may now have their address taken by the runtime via " << std::endl
+ << "::cudaGetSymbolAddress(). " << std::endl;
+ break;
case cudaErrorTextureFetchFailed:
std::cerr
- << "This indicated that a texture fetch was not able to be performed. " << std::endl
- << "This was previously used for device emulation of texture operations. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "This indicated that a texture fetch was not able to be performed. " << std::endl
+ << "This was previously used for device emulation of texture operations. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
case cudaErrorTextureNotBound:
std::cerr
- << "This indicated that a texture was not bound for access. " << std::endl
- << "This was previously used for device emulation of texture operations. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "This indicated that a texture was not bound for access. " << std::endl
+ << "This was previously used for device emulation of texture operations. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
case cudaErrorSynchronizationError:
std::cerr
- << "This indicated that a synchronization operation had failed. " << std::endl
- << "This was previously used for some device emulation functions. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "This indicated that a synchronization operation had failed. " << std::endl
+ << "This was previously used for some device emulation functions. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
case cudaErrorInvalidFilterSetting:
std::cerr
- << "This indicates that a non-float texture was being accessed with linear " << std::endl
- << "filtering. This is not supported by CUDA. " << std::endl;
- break;
+ << "This indicates that a non-float texture was being accessed with linear " << std::endl
+ << "filtering. This is not supported by CUDA. " << std::endl;
+ break;
case cudaErrorInvalidNormSetting:
std::cerr
- << "This indicates that an attempt was made to read a non-float texture as a " << std::endl
- << "normalized float. This is not supported by CUDA. " << std::endl;
- break;
+ << "This indicates that an attempt was made to read a non-float texture as a " << std::endl
+ << "normalized float. This is not supported by CUDA. " << std::endl;
+ break;
case cudaErrorMixedDeviceExecution:
std::cerr
- << "Mixing of device and device emulation code was not allowed. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "Mixing of device and device emulation code was not allowed. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
case cudaErrorCudartUnloading:
std::cerr
- << "This indicated an issue with calling API functions during the unload " << std::endl
- << "process of the CUDA runtime in prior releases. " << std::endl
- << "This error return is deprecated as of CUDA 3.2. " << std::endl;
- break;
+ << "This indicated an issue with calling API functions during the unload " << std::endl
+ << "process of the CUDA runtime in prior releases. " << std::endl
+ << "This error return is deprecated as of CUDA 3.2. " << std::endl;
+ break;
case cudaErrorUnknown:
std::cerr
- << "This indicates that an unknown internal error has occurred. " << std::endl;
- break;
+ << "This indicates that an unknown internal error has occurred. " << std::endl;
+ break;
case cudaErrorNotYetImplemented:
std::cerr
- << "This indicates that the API call is not yet implemented. Production " << std::endl
- << "releases of CUDA will never return this error. " << std::endl;
- break;
+ << "This indicates that the API call is not yet implemented. Production " << std::endl
+ << "releases of CUDA will never return this error. " << std::endl;
+ break;
case cudaErrorMemoryValueTooLarge:
std::cerr
- << "This indicated that an emulated device pointer exceeded the 32-bit address " << std::endl
- << "range. " << std::endl
- << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
- << "removed with the CUDA 3.1 release. " << std::endl;
- break;
+ << "This indicated that an emulated device pointer exceeded the 32-bit address " << std::endl
+ << "range. " << std::endl
+ << "This error return is deprecated as of CUDA 3.1. Device emulation mode was " << std::endl
+ << "removed with the CUDA 3.1 release. " << std::endl;
+ break;
case cudaErrorInvalidResourceHandle:
std::cerr
- << "This indicates that a resource handle passed to the API call was not " << std::endl
- << "valid. Resource handles are opaque types like ::cudaStream_t and " << std::endl
- << "::cudaEvent_t. " << std::endl;
- break;
+ << "This indicates that a resource handle passed to the API call was not " << std::endl
+ << "valid. Resource handles are opaque types like ::cudaStream_t and " << std::endl
+ << "::cudaEvent_t. " << std::endl;
+ break;
case cudaErrorNotReady:
std::cerr
- << "This indicates that asynchronous operations issued previously have not " << std::endl
- << "completed yet. This result is not actually an error, but must be indicated " << std::endl
- << "differently than ::cudaSuccess (which indicates completion). Calls that " << std::endl
- << "may return this value include ::cudaEventQuery() and ::cudaStreamQuery(). " << std::endl;
- break;
+ << "This indicates that asynchronous operations issued previously have not " << std::endl
+ << "completed yet. This result is not actually an error, but must be indicated " << std::endl
+ << "differently than ::cudaSuccess (which indicates completion). Calls that " << std::endl
+ << "may return this value include ::cudaEventQuery() and ::cudaStreamQuery(). " << std::endl;
+ break;
case cudaErrorInsufficientDriver:
std::cerr
- << "This indicates that the installed NVIDIA CUDA driver is older than the " << std::endl
- << "CUDA runtime library. This is not a supported configuration. Users should " << std::endl
- << "install an updated NVIDIA display driver to allow the application to run. " << std::endl;
- break;
+ << "This indicates that the installed NVIDIA CUDA driver is older than the " << std::endl
+ << "CUDA runtime library. This is not a supported configuration. Users should " << std::endl
+ << "install an updated NVIDIA display driver to allow the application to run. " << std::endl;
+ break;
case cudaErrorSetOnActiveProcess:
std::cerr
- << "This indicates that the user has called ::cudaSetDevice(), " << std::endl
- << "::cudaSetValidDevices(), ::cudaSetDeviceFlags(), " << std::endl
- << "::cudaD3D9SetDirect3DDevice(), ::cudaD3D10SetDirect3DDevice, " << std::endl
- << "::cudaD3D11SetDirect3DDevice(), * or ::cudaVDPAUSetVDPAUDevice() after " << std::endl
- << "initializing the CUDA runtime by calling non-device management operations " << std::endl
- << "(allocating memory and launching kernels are examples of non-device " << std::endl
- << "management operations). This error can also be returned if using " << std::endl
- << "runtime/driver interoperability and there is an existing ::CUcontext " << std::endl
- << "active on the host thread. " << std::endl;
- break;
+ << "This indicates that the user has called ::cudaSetDevice(), " << std::endl
+ << "::cudaSetValidDevices(), ::cudaSetDeviceFlags(), " << std::endl
+ << "::cudaD3D9SetDirect3DDevice(), ::cudaD3D10SetDirect3DDevice, " << std::endl
+ << "::cudaD3D11SetDirect3DDevice(), * or ::cudaVDPAUSetVDPAUDevice() after " << std::endl
+ << "initializing the CUDA runtime by calling non-device management operations " << std::endl
+ << "(allocating memory and launching kernels are examples of non-device " << std::endl
+ << "management operations). This error can also be returned if using " << std::endl
+ << "runtime/driver interoperability and there is an existing ::CUcontext " << std::endl
+ << "active on the host thread. " << std::endl;
+ break;
case cudaErrorInvalidSurface:
std::cerr
- << "This indicates that the surface passed to the API call is not a valid " << std::endl
- << "surface. " << std::endl;
- break;
+ << "This indicates that the surface passed to the API call is not a valid " << std::endl
+ << "surface. " << std::endl;
+ break;
case cudaErrorNoDevice:
- std::cerr
- << "This indicates that no CUDA-capable devices were detected by the installed " << std::endl
- << "CUDA driver. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that no CUDA-capable devices were detected by the installed " << std::endl
+ << "CUDA driver. " << std::endl;
+ break;
case cudaErrorECCUncorrectable:
- std::cerr
- << "This indicates that an uncorrectable ECC error was detected during " << std::endl
- << "execution. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that an uncorrectable ECC error was detected during " << std::endl
+ << "execution. " << std::endl;
+ break;
case cudaErrorSharedObjectSymbolNotFound:
- std::cerr
- << "This indicates that a link to a shared object failed to resolve. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that a link to a shared object failed to resolve. " << std::endl;
+ break;
case cudaErrorSharedObjectInitFailed:
- std::cerr
- << "This indicates that initialization of a shared object failed. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that initialization of a shared object failed. " << std::endl;
+ break;
case cudaErrorUnsupportedLimit:
- std::cerr
- << "This indicates that the ::cudaLimit passed to the API call is not " << std::endl
- << "supported by the active device. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that the ::cudaLimit passed to the API call is not " << std::endl
+ << "supported by the active device. " << std::endl;
+ break;
case cudaErrorDuplicateVariableName:
- std::cerr
- << "This indicates that multiple global or constant variables (across separate " << std::endl
- << "CUDA source files in the application) share the same string name. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that multiple global or constant variables (across separate " << std::endl
+ << "CUDA source files in the application) share the same string name. " << std::endl;
+ break;
case cudaErrorDuplicateTextureName:
- std::cerr
- << "This indicates that multiple textures (across separate CUDA source " << std::endl
- << "files in the application) share the same string name. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that multiple textures (across separate CUDA source " << std::endl
+ << "files in the application) share the same string name. " << std::endl;
+ break;
case cudaErrorDuplicateSurfaceName:
- std::cerr
- << "This indicates that multiple surfaces (across separate CUDA source " << std::endl
- << "files in the application) share the same string name. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that multiple surfaces (across separate CUDA source " << std::endl
+ << "files in the application) share the same string name. " << std::endl;
+ break;
case cudaErrorDevicesUnavailable:
- std::cerr
- << "This indicates that all CUDA devices are busy or unavailable at the current " << std::endl
- << "time. Devices are often busy/unavailable due to use of " << std::endl
- << "::cudaComputeModeExclusive or ::cudaComputeModeProhibited. They can also " << std::endl
- << "be unavailable due to memory constraints on a device that already has " << std::endl
- << "active CUDA work being performed. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that all CUDA devices are busy or unavailable at the current " << std::endl
+ << "time. Devices are often busy/unavailable due to use of " << std::endl
+ << "::cudaComputeModeExclusive or ::cudaComputeModeProhibited. They can also " << std::endl
+ << "be unavailable due to memory constraints on a device that already has " << std::endl
+ << "active CUDA work being performed. " << std::endl;
+ break;
case cudaErrorInvalidKernelImage:
- std::cerr
- << "This indicates that the device kernel image is invalid. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that the device kernel image is invalid. " << std::endl;
+ break;
case cudaErrorNoKernelImageForDevice:
- std::cerr
- << "This indicates that there is no kernel image available that is suitable " << std::endl
- << "for the device. This can occur when a user specifies code generation " << std::endl
- << "options for a particular CUDA source file that do not include the " << std::endl
- << "corresponding device configuration. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that there is no kernel image available that is suitable " << std::endl
+ << "for the device. This can occur when a user specifies code generation " << std::endl
+ << "options for a particular CUDA source file that do not include the " << std::endl
+ << "corresponding device configuration. " << std::endl;
+ break;
case cudaErrorIncompatibleDriverContext:
- std::cerr
- << "This indicates that the current context is not compatible with this " << std::endl
- << "version of the CUDA Runtime. This can only occur if you are using CUDA " << std::endl
- << "Runtime/Driver interoperability and have created an existing Driver " << std::endl
- << "context using an older API. Please see \ref CUDART_DRIVER " << std::endl
- << "\"Interactions with the CUDA Driver API\" for more information. " << std::endl;
- break;
+ std::cerr
+ << "This indicates that the current context is not compatible with this " << std::endl
+ << "version of the CUDA Runtime. This can only occur if you are using CUDA " << std::endl
+ << "Runtime/Driver interoperability and have created an existing Driver " << std::endl
+ << "context using an older API. Please see \\ref CUDART_DRIVER " << std::endl
+ << "\"Interactions with the CUDA Driver API\" for more information. " << std::endl;
+ break;
case cudaErrorStartupFailure:
- std::cerr
- << "This indicates an internal startup failure in the CUDA runtime. " << std::endl;
- break;
+ std::cerr
+ << "This indicates an internal startup failure in the CUDA runtime. " << std::endl;
+ break;
case cudaErrorApiFailureBase:
- std::cerr
- << "Any unhandled CUDA driver error is added to this value and returned via " << std::endl
- << "the runtime. Production releases of CUDA should not return such errors. " << std::endl;
- break;
+ std::cerr
+ << "Any unhandled CUDA driver error is added to this value and returned via " << std::endl
+ << "the runtime. Production releases of CUDA should not return such errors. " << std::endl;
+ break;
+ default:
+ std::cerr << "(detailed description is not available)" << std::endl;
+ break;
}
throw EXIT_FAILURE;
return false;
diff --git a/src/TNL/Devices/Cuda.h b/src/TNL/Devices/Cuda.h
index dc3923f4af70ebbf4161b8542d08b2308b5fcf4f..4a09f0cd81c6028199bde20750034a28418f6e4c 100644
--- a/src/TNL/Devices/Cuda.h
+++ b/src/TNL/Devices/Cuda.h
@@ -82,6 +82,27 @@ class Cuda
#ifdef HAVE_CUDA
template< typename Index >
static __device__ Index getInterleaving( const Index index );
+
+ /****
+ * Declaration of variables for dynamic shared memory is difficult in
+ * templated functions. For example, the following does not work for
+ * different types T:
+ *
+ * template< typename T >
+ * void foo()
+ * {
+ * extern __shared__ T shx[];
+ * }
+ *
+ * This is because extern variables must be declared exactly once. In
+ * templated functions we need to have same variable name with different
+ * type, which causes the conflict. In CUDA samples they solve the problem
+ * using template specialization via classes, but using one base type and
+ * reinterpret_cast works too.
+ * See http://stackoverflow.com/a/19339004/4180822 for reference.
+ */
+ template< typename Element, size_t Alignment = sizeof( Element ) >
+ static __device__ Element* getSharedMemory();
#endif
#ifdef HAVE_CUDA
@@ -126,29 +147,6 @@ class Cuda
#define CudaSupportMissingMessage \
std::cerr << "The CUDA support is missing in the source file " << __FILE__ << " at line " << __LINE__ << ". Please set WITH_CUDA=yes in the install script. " << std::endl;
-
-// TODO: This would be nice in Cuda but C++ standard does not allow it.
-#ifdef HAVE_CUDA
- template< typename Element >
- struct getSharedMemory
- {
- __device__ operator Element*();
- };
-
- template<>
- struct getSharedMemory< double >
- {
- inline __device__ operator double*();
- };
-
- template<>
- struct getSharedMemory< long int >
- {
- inline __device__ operator long int*();
- };
-
-#endif
-
} // namespace Devices
} // namespace TNL
diff --git a/src/TNL/Devices/Cuda_impl.h b/src/TNL/Devices/Cuda_impl.h
index 9c0d252a5a9e59458ca6cb915e4543dba6cfddbe..580653b708ec95ba88e312c7a492eaafd5483757 100644
--- a/src/TNL/Devices/Cuda_impl.h
+++ b/src/TNL/Devices/Cuda_impl.h
@@ -140,25 +140,12 @@ __device__ Index Cuda::getInterleaving( const Index index )
return index + index / Cuda::getNumberOfSharedMemoryBanks();
}
-template< typename Element >
-__device__ getSharedMemory< Element >::operator Element*()
+template< typename Element, size_t Alignment >
+__device__ Element* Cuda::getSharedMemory()
{
- extern __shared__ int __sharedMemory[];
- return ( Element* ) __sharedMemory;
-};
-
-__device__ inline getSharedMemory< double >::operator double*()
-{
- extern __shared__ double __sharedMemoryDouble[];
- return ( double* ) __sharedMemoryDouble;
-};
-
-__device__ inline getSharedMemory< long int >::operator long int*()
-{
- extern __shared__ long int __sharedMemoryLongInt[];
- return ( long int* ) __sharedMemoryLongInt;
-};
-
+ extern __shared__ __align__ ( Alignment ) unsigned char __sdata[];
+ return reinterpret_cast< Element* >( __sdata );
+}
#endif /* HAVE_CUDA */
} // namespace Devices
diff --git a/src/TNL/Functions/MeshFunction.h b/src/TNL/Functions/MeshFunction.h
index f79ae2ed88f8542f76aef01a4dd42692ff790863..a3cf8a46c307312ff0e10127da53339c48ddb3e4 100644
--- a/src/TNL/Functions/MeshFunction.h
+++ b/src/TNL/Functions/MeshFunction.h
@@ -165,5 +165,3 @@ class MeshFunction :
} // namespace TNL
#include <TNL/Functions/MeshFunction_impl.h>
-#include <TNL/Functions/MeshFunctionGnuplotWriter_impl.h>
-#include <TNL/Functions/MeshFunctionVTKWriter_impl.h>
diff --git a/src/TNL/Functions/MeshFunctionGnuplotWriter.h b/src/TNL/Functions/MeshFunctionGnuplotWriter.h
index 15d715548d6304c88788219293d75f99d1270d5e..887498711e8f3e4422a7cc99994cfecc2301504d 100644
--- a/src/TNL/Functions/MeshFunctionGnuplotWriter.h
+++ b/src/TNL/Functions/MeshFunctionGnuplotWriter.h
@@ -175,3 +175,4 @@ class MeshFunctionGnuplotWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device
} // namespace Functions
} // namespace TNL
+#include <TNL/Functions/MeshFunctionGnuplotWriter_impl.h>
diff --git a/src/TNL/Functions/MeshFunctionGnuplotWriter_impl.h b/src/TNL/Functions/MeshFunctionGnuplotWriter_impl.h
index 5600bae1f891d6a8c8dd8db12b0b4c578fecc62f..65a555bacf7db21b26aaca4492068301bd16a255 100644
--- a/src/TNL/Functions/MeshFunctionGnuplotWriter_impl.h
+++ b/src/TNL/Functions/MeshFunctionGnuplotWriter_impl.h
@@ -10,6 +10,8 @@
#pragma once
+#include <TNL/Functions/MeshFunctionGnuplotWriter.h>
+
namespace TNL {
namespace Functions {
diff --git a/src/TNL/Functions/MeshFunctionVTKWriter.h b/src/TNL/Functions/MeshFunctionVTKWriter.h
index 239b62f3795b2b373f4725aab699516080778763..39651c27450c407639c1278f62220b82f51f240b 100644
--- a/src/TNL/Functions/MeshFunctionVTKWriter.h
+++ b/src/TNL/Functions/MeshFunctionVTKWriter.h
@@ -10,9 +10,13 @@
#pragma once
+#include <TNL/Meshes/Grid.h>
+
namespace TNL {
namespace Functions {
+template< typename, int, typename > class MeshFunction;
+
template< typename MeshFunction >
class MeshFunctionVTKWriter
{
@@ -207,3 +211,4 @@ class MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, Me
} // namespace Functions
} // namespace TNL
+#include <TNL/Functions/MeshFunctionVTKWriter_impl.h>
diff --git a/src/TNL/Functions/MeshFunctionVTKWriter_impl.h b/src/TNL/Functions/MeshFunctionVTKWriter_impl.h
index 0493b0a421c973c110a661364ce4aa9814fc30c8..b432671764c0545c64e5c8ea241c820525f20f05 100644
--- a/src/TNL/Functions/MeshFunctionVTKWriter_impl.h
+++ b/src/TNL/Functions/MeshFunctionVTKWriter_impl.h
@@ -10,6 +10,8 @@
#pragma once
+#include <TNL/Functions/MeshFunctionVTKWriter.h>
+
namespace TNL {
namespace Functions {
@@ -34,7 +36,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 1, MeshReal, Device, MeshIndex >, 1, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -61,7 +63,6 @@ write( const MeshFunctionType& function,
const RealType spaceStep = mesh.getSpaceSteps().x();
str << "POINTS " << mesh.getDimensions().x() + 1 << " float" << std::endl;
-
for (int i = 0; i <= mesh.getDimensions().x(); i++)
{
str << origin + i * spaceStep << " 0 0" << std::endl;
@@ -83,11 +84,9 @@ write( const MeshFunctionType& function,
str << "SCALARS cellFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Cell entity( mesh );
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
+ for( MeshIndex i = 0; i < mesh.template getEntitiesCount< typename MeshType::Cell >(); i++ )
{
+ typename MeshType::Cell entity = mesh.template getEntity< typename MeshType::Cell >( i );
entity.refresh();
str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
@@ -106,7 +105,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 1, MeshReal, Device, MeshIndex >, 0, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -133,7 +132,6 @@ write( const MeshFunctionType& function,
const RealType spaceStep = mesh.getSpaceSteps().x();
str << "POINTS " << mesh.getDimensions().x() + 1 << " float" << std::endl;
-
for (int i = 0; i < mesh.getDimensions().x() + 1; i++)
{
str << origin + i * spaceStep << " 0 0" << std::endl;
@@ -155,11 +153,9 @@ write( const MeshFunctionType& function,
str << "SCALARS VerticesFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Vertex entity( mesh );
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
+ for( MeshIndex i = 0; i < mesh.template getEntitiesCount< typename MeshType::Vertex >(); i++ )
{
+ typename MeshType::Vertex entity = mesh.template getEntity< typename MeshType::Vertex >( i );
entity.refresh();
str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
@@ -178,7 +174,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 2, MeshReal, Device, MeshIndex >, 2, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -205,9 +201,10 @@ write( const MeshFunctionType& function,
const RealType spaceStepX = mesh.getSpaceSteps().x();
const RealType originY = mesh.getOrigin().y();
const RealType spaceStepY = mesh.getSpaceSteps().y();
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ const MeshIndex entitiesCount = mesh.template getEntitiesCount< typename MeshType::Cell >();
- str << "POINTS " << (mesh.getDimensions().x() + 1) * (mesh.getDimensions().y() + 1) << " float" << std::endl;
-
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int j = 0; j < mesh.getDimensions().y() + 1; j++)
{
for (int i = 0; i < mesh.getDimensions().x() + 1; i++)
@@ -216,8 +213,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << mesh.getDimensions().x() * mesh.getDimensions().y() << " " <<
- mesh.getDimensions().x() * mesh.getDimensions().y() * 5 << std::endl;
+ str << std::endl << "CELLS " << entitiesCount << " " << entitiesCount * 5 << std::endl;
for (int j = 0; j < mesh.getDimensions().y(); j++)
{
for (int i = 0; i < mesh.getDimensions().x(); i++)
@@ -233,23 +229,17 @@ write( const MeshFunctionType& function,
str << "8 " << std::endl;
}
- str << std::endl << "CELL_DATA " << mesh.getDimensions().x() * mesh.getDimensions().y() << std::endl;
+ str << std::endl << "CELL_DATA " << entitiesCount << std::endl;
str << "SCALARS cellFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Cell entity( mesh );
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
+ for( MeshIndex i = 0; i < entitiesCount; i++ )
{
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
+ typename MeshType::Cell entity = mesh.template getEntity< typename MeshType::Cell >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
+
return true;
}
@@ -264,7 +254,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 2, MeshReal, Device, MeshIndex >, 1, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -293,9 +283,10 @@ write( const MeshFunctionType& function,
const RealType spaceStepX = mesh.getSpaceSteps().x();
const RealType originY = mesh.getOrigin().y();
const RealType spaceStepY = mesh.getSpaceSteps().y();
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ const MeshIndex entitiesCount = mesh.template getEntitiesCount< typename MeshType::Face >();
- str << "POINTS " << mesh.template getEntitiesCount< Vertex >() << " float" << std::endl;
-
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int j = 0; j < ( mesh.getDimensions().y() + 1); j++)
{
for (int i = 0; i < ( mesh.getDimensions().x() + 1 ); i++)
@@ -304,8 +295,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << mesh.template getEntitiesCount< Face >() << " " <<
- mesh.template getEntitiesCount< Face >() * 3 << std::endl;
+ str << std::endl << "CELLS " << entitiesCount << " " << entitiesCount * 3 << std::endl;
for (int j = 0; j < mesh.getDimensions().y(); j++)
{
for (int i = 0; i < ( mesh.getDimensions().x() + 1 ); i++)
@@ -322,49 +312,23 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELL_TYPES " << mesh.template getEntitiesCount< Face >() << std::endl;
- for (int i = 0; i < mesh.template getEntitiesCount< Face >(); i++)
+ str << std::endl << "CELL_TYPES " << entitiesCount << std::endl;
+ for (int i = 0; i < entitiesCount; i++)
{
str << "3" << std::endl;
}
- str << std::endl << "CELL_DATA " << mesh.template getEntitiesCount< Face >() << std::endl;
+ str << std::endl << "CELL_DATA " << entitiesCount << std::endl;
str << "SCALARS FaceslFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typedef typename MeshType::Face EntityType;
- typedef typename EntityType::EntityOrientationType EntityOrientation;
- EntityType entity( mesh );
-
- entity.setOrientation( EntityOrientation( 1.0, 0.0 ) );
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
+ for( MeshIndex i = 0; i < entitiesCount; i++ )
{
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
+ typename MeshType::Face entity = mesh.template getEntity< typename MeshType::Face >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
-
- entity.setOrientation( EntityOrientation( 0.0, 1.0 ) );
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
-
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
return true;
}
@@ -379,7 +343,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 2, MeshReal, Device, MeshIndex >, 0, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -407,10 +371,9 @@ write( const MeshFunctionType& function,
const RealType spaceStepX = mesh.getSpaceSteps().x();
const RealType originY = mesh.getOrigin().y();
const RealType spaceStepY = mesh.getSpaceSteps().y();
-
-
- str << "POINTS " << mesh.template getEntitiesCount< Vertex >() << " float" << std::endl;
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int j = 0; j < ( mesh.getDimensions().y() + 1); j++)
{
for (int i = 0; i < ( mesh.getDimensions().x() + 1 ); i++)
@@ -419,8 +382,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << mesh.template getEntitiesCount< Vertex >() << " " <<
- mesh.template getEntitiesCount< Vertex >() * 2 << std::endl;
+ str << std::endl << "CELLS " << verticesCount << " " << verticesCount * 2 << std::endl;
for (int j = 0; j < ( mesh.getDimensions().y() + 1 ); j++)
{
for (int i = 0; i < ( mesh.getDimensions().x() + 1 ); i++)
@@ -429,30 +391,23 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELL_TYPES " << mesh.template getEntitiesCount< Vertex >() << std::endl;
- for (int i = 0; i < mesh.template getEntitiesCount< Vertex >(); i++)
+ str << std::endl << "CELL_TYPES " << verticesCount << std::endl;
+ for (int i = 0; i < verticesCount; i++)
{
str << "1" << std::endl;
}
- str << std::endl << "CELL_DATA " << mesh.template getEntitiesCount< Vertex >() << std::endl;
+ str << std::endl << "CELL_DATA " << verticesCount << std::endl;
str << "SCALARS VerticesFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Vertex entity( mesh );
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
+ for( MeshIndex i = 0; i < verticesCount; i++ )
{
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
+ typename MeshType::Vertex entity = mesh.template getEntity< typename MeshType::Vertex >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
-
+
return true;
}
@@ -467,7 +422,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshIndex >, 3, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -496,11 +451,10 @@ write( const MeshFunctionType& function,
const RealType spaceStepY = mesh.getSpaceSteps().y();
const RealType originZ = mesh.getOrigin().z();
const RealType spaceStepZ = mesh.getSpaceSteps().z();
- const RealType entitiesCount = mesh.getDimensions().x() * mesh.getDimensions().y() * mesh.getDimensions().z();
-
- str << "POINTS " << (mesh.getDimensions().x()+1) * (mesh.getDimensions().y()+1) * (mesh.getDimensions().z()+1) <<
- " float" << std::endl;
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ const MeshIndex entitiesCount = mesh.template getEntitiesCount< typename MeshType::Cell >();
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int k = 0; k <= mesh.getDimensions().y(); k++)
{
for (int j = 0; j <= mesh.getDimensions().y(); j++)
@@ -543,24 +497,13 @@ write( const MeshFunctionType& function,
str << "SCALARS cellFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Cell entity( mesh );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() < mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
+ for( MeshIndex i = 0; i < entitiesCount; i++ )
{
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
+ typename MeshType::Cell entity = mesh.template getEntity< typename MeshType::Cell >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
+
return true;
}
@@ -575,7 +518,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshIndex >, 2, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -595,8 +538,6 @@ MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshInde
write( const MeshFunctionType& function,
std::ostream& str )
{
- typedef typename MeshType::template MeshEntity< 2 > Face;
- typedef typename MeshType::template MeshEntity< 3 > Cell;
writeHeader(function, str);
const MeshType& mesh = function.getMesh();
@@ -606,12 +547,10 @@ write( const MeshFunctionType& function,
const RealType spaceStepY = mesh.getSpaceSteps().y();
const RealType originZ = mesh.getOrigin().z();
const RealType spaceStepZ = mesh.getSpaceSteps().z();
- const RealType entitiesCount = mesh.template getEntitiesCount< Face >();
- const RealType pointsCount = mesh.template getEntitiesCount< Cell >();
-
- str << "POINTS " << pointsCount <<
- " float" << std::endl;
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ const MeshIndex entitiesCount = mesh.template getEntitiesCount< typename MeshType::Face >();
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int k = 0; k <= mesh.getDimensions().y(); k++)
{
for (int j = 0; j <= mesh.getDimensions().y(); j++)
@@ -624,8 +563,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << entitiesCount << " " <<
- entitiesCount * 5 << std::endl;
+ str << std::endl << "CELLS " << entitiesCount << " " << entitiesCount * 5 << std::endl;
for (int k = 0; k < mesh.getDimensions().z(); k++)
{
for (int j = 0; j < mesh.getDimensions().y(); j++)
@@ -678,67 +616,13 @@ write( const MeshFunctionType& function,
str << "SCALARS facesFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typedef typename MeshType::Face EntityType;
- typedef typename EntityType::EntityOrientationType EntityOrientation;
- EntityType entity( mesh );
-
- entity.setOrientation( EntityOrientation( 1.0, 0.0 , 0.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() < mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
- {
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
- }
-
- entity.setOrientation( EntityOrientation( 0.0, 1.0 , 0.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() < mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
- {
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
- }
-
- entity.setOrientation( EntityOrientation( 0.0, 0.0 , 1.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() <= mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
+ for( MeshIndex i = 0; i < entitiesCount; i++ )
{
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
+ typename MeshType::Face entity = mesh.template getEntity< typename MeshType::Face >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
-
+
return true;
}
@@ -753,7 +637,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshIndex >, 1, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -773,8 +657,6 @@ MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshInde
write( const MeshFunctionType& function,
std::ostream& str )
{
- typedef typename MeshType::template MeshEntity< 1 > Edge;
- typedef typename MeshType::template MeshEntity< 3 > Cell;
writeHeader(function, str);
const MeshType& mesh = function.getMesh();
@@ -784,12 +666,10 @@ write( const MeshFunctionType& function,
const RealType spaceStepY = mesh.getSpaceSteps().y();
const RealType originZ = mesh.getOrigin().z();
const RealType spaceStepZ = mesh.getSpaceSteps().z();
- const RealType entitiesCount = mesh.template getEntitiesCount< Edge >();
- const RealType pointsCount = mesh.template getEntitiesCount< Cell >();
-
- str << "POINTS " << pointsCount <<
- " float" << std::endl;
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
+ const MeshIndex entitiesCount = mesh.template getEntitiesCount< typename MeshType::Edge >();
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int k = 0; k <= mesh.getDimensions().y(); k++)
{
for (int j = 0; j <= mesh.getDimensions().y(); j++)
@@ -802,8 +682,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << entitiesCount << " " <<
- entitiesCount * 3 << std::endl;
+ str << std::endl << "CELLS " << entitiesCount << " " << entitiesCount * 3 << std::endl;
for (int k = 0; k <= mesh.getDimensions().z(); k++)
{
for (int j = 0; j <= mesh.getDimensions().y(); j++)
@@ -850,67 +729,13 @@ write( const MeshFunctionType& function,
str << "SCALARS edgesFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typedef typename MeshType::Face EntityType;
- typedef typename EntityType::EntityOrientationType EntityOrientation;
- EntityType entity( mesh );
-
- entity.setOrientation( EntityOrientation( 1.0, 0.0 , 0.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() <= mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
+ for( MeshIndex i = 0; i < entitiesCount; i++ )
{
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() < mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
- }
-
- entity.setOrientation( EntityOrientation( 0.0, 1.0 , 0.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() <= mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
- {
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() < mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
- }
-
- entity.setOrientation( EntityOrientation( 0.0, 0.0 , 1.0) );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() < mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
- {
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
+ typename MeshType::Edge entity = mesh.template getEntity< typename MeshType::Edge >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
-
+
return true;
}
@@ -925,7 +750,7 @@ template< typename MeshReal,
void
MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshIndex >, 0, Real > >::
writeHeader( const MeshFunctionType& function,
- std::ostream& str )
+ std::ostream& str )
{
const MeshType& mesh = function.getMesh();
const typename MeshType::VertexType& origin = mesh.getOrigin();
@@ -945,7 +770,6 @@ MeshFunctionVTKWriter< MeshFunction< Meshes::Grid< 3, MeshReal, Device, MeshInde
write( const MeshFunctionType& function,
std::ostream& str )
{
- typedef typename MeshType::template MeshEntity< 0 > Vertex;
writeHeader(function, str);
const MeshType& mesh = function.getMesh();
@@ -955,9 +779,9 @@ write( const MeshFunctionType& function,
const RealType spaceStepY = mesh.getSpaceSteps().y();
const RealType originZ = mesh.getOrigin().z();
const RealType spaceStepZ = mesh.getSpaceSteps().z();
+ const MeshIndex verticesCount = mesh.template getEntitiesCount< typename MeshType::Vertex >();
- str << "POINTS " << mesh.template getEntitiesCount< Vertex >() << " float" << std::endl;
-
+ str << "POINTS " << verticesCount << " float" << std::endl;
for (int k = 0; k <= mesh.getDimensions().y(); k++)
{
for (int j = 0; j <= mesh.getDimensions().y(); j++)
@@ -970,8 +794,7 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELLS " << mesh.template getEntitiesCount< Vertex >() << " " <<
- mesh.template getEntitiesCount< Vertex >() * 2 << std::endl;
+ str << std::endl << "CELLS " << verticesCount << " " << verticesCount * 2 << std::endl;
for (int k = 0; k < ( mesh.getDimensions().z() + 1 ); k++)
{
for (int j = 0; j < ( mesh.getDimensions().y() + 1 ); j++)
@@ -983,35 +806,23 @@ write( const MeshFunctionType& function,
}
}
- str << std::endl << "CELL_TYPES " << mesh.template getEntitiesCount< Vertex >() << std::endl;
- for (int i = 0; i < mesh.template getEntitiesCount< Vertex >(); i++)
+ str << std::endl << "CELL_TYPES " << verticesCount << std::endl;
+ for (int i = 0; i < verticesCount; i++)
{
str << "1" << std::endl;
}
- str << std::endl << "CELL_DATA " << mesh.template getEntitiesCount< Vertex >() << std::endl;
+ str << std::endl << "CELL_DATA " << verticesCount << std::endl;
str << "SCALARS verticesFunctionValues float 1" << std::endl;
str << "LOOKUP_TABLE default" << std::endl;
- typename MeshType::Vertex entity( mesh );
- for( entity.getCoordinates().z() = 0;
- entity.getCoordinates().z() <= mesh.getDimensions().z();
- entity.getCoordinates().z() ++ )
+ for( MeshIndex i = 0; i < verticesCount; i++ )
{
- for( entity.getCoordinates().y() = 0;
- entity.getCoordinates().y() <= mesh.getDimensions().y();
- entity.getCoordinates().y() ++ )
- {
- for( entity.getCoordinates().x() = 0;
- entity.getCoordinates().x() <= mesh.getDimensions().x();
- entity.getCoordinates().x() ++ )
- {
- entity.refresh();
- str << function.getData().getElement( entity.getIndex() ) << std::endl;
- }
- }
+ typename MeshType::Vertex entity = mesh.template getEntity< typename MeshType::Vertex >( i );
+ entity.refresh();
+ str << function.getData().getElement( entity.getIndex() ) << std::endl;
}
-
+
return true;
}
diff --git a/src/TNL/Images/DicomSeries.h b/src/TNL/Images/DicomSeries.h
index 1350cf4a280b7ba0ad89149f818a5476f29a98e1..b3ee1fb11c88f221512e6c27071f60fbdafddbc8 100644
--- a/src/TNL/Images/DicomSeries.h
+++ b/src/TNL/Images/DicomSeries.h
@@ -15,7 +15,7 @@
#pragma once
#include <TNL/Containers/Array.h>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/String.h>
#include <TNL/param-types.h>
#include <TNL/Images//Image.h>
@@ -106,7 +106,7 @@ class DicomSeries : public Image< int >
bool loadImage( const String& filePath, int number );
- List< String > fileList;
+ Containers::List< String > fileList;
Containers::Array<DicomHeader *,Devices::Host,int> dicomSeriesHeaders;
diff --git a/src/TNL/Images/DicomSeries_impl.h b/src/TNL/Images/DicomSeries_impl.h
index 5d0abfea4a386dd4fcb6b5fc1590ced960b56955..e9a3f064c78925c2dc333e43f56a50d8a86f0c4a 100644
--- a/src/TNL/Images/DicomSeries_impl.h
+++ b/src/TNL/Images/DicomSeries_impl.h
@@ -155,7 +155,7 @@ inline bool DicomSeries::retrieveFileList( const String& filePath)
String fileNamePrefix(fileName.getString(), 0, fileName.getLength() - separatorPosition);
struct dirent **dirp;
- List<String > files;
+ Containers::List<String > files;
//scan and sort directory
int ndirs = scandir(directoryPath.getString(), &dirp, filter, alphasort);
diff --git a/src/TNL/Matrices/CSR_impl.h b/src/TNL/Matrices/CSR_impl.h
index 0d7261a57351f67cf021b4485c531743ec3e9b30..d32355c7326f5fd8b98f749a660cc74521ffe425 100644
--- a/src/TNL/Matrices/CSR_impl.h
+++ b/src/TNL/Matrices/CSR_impl.h
@@ -629,7 +629,7 @@ void CSR< Real, Device, Index >::spmvCudaVectorized( const InVector& inVector,
const IndexType warpEnd,
const IndexType inWarpIdx ) const
{
- volatile Real* aux = Devices::getSharedMemory< Real >();
+ volatile Real* aux = Devices::Cuda::getSharedMemory< Real >();
for( IndexType row = warpStart; row < warpEnd; row++ )
{
aux[ threadIdx.x ] = 0.0;
diff --git a/src/TNL/Matrices/ChunkedEllpack_impl.h b/src/TNL/Matrices/ChunkedEllpack_impl.h
index 325ae861ab985f517879a3aeb4ea4076cedcb003..aed5b5ededb34f521271e749e5f31498274a82aa 100644
--- a/src/TNL/Matrices/ChunkedEllpack_impl.h
+++ b/src/TNL/Matrices/ChunkedEllpack_impl.h
@@ -1068,7 +1068,7 @@ __device__ void ChunkedEllpack< Real, Device, Index >::computeSliceVectorProduct
{
static_assert( std::is_same < DeviceType, Devices::Cuda >::value, "" );
- RealType* chunkProducts = Devices::getSharedMemory< RealType >();
+ RealType* chunkProducts = Devices::Cuda::getSharedMemory< RealType >();
ChunkedEllpackSliceInfo* sliceInfo = ( ChunkedEllpackSliceInfo* ) & chunkProducts[ blockDim.x ];
if( threadIdx.x == 0 )
diff --git a/src/TNL/Matrices/MatrixOperations.h b/src/TNL/Matrices/MatrixOperations.h
index 4cecc842596b3cbb0e7913c174d6e6794b159fbd..5ad1fcef3619395defd5d1104810e61fe69c99a6 100644
--- a/src/TNL/Matrices/MatrixOperations.h
+++ b/src/TNL/Matrices/MatrixOperations.h
@@ -90,16 +90,7 @@ GemvCudaKernel( const IndexType m,
IndexType elementIdx = blockIdx.x * blockDim.x + threadIdx.x;
const IndexType gridSize = blockDim.x * gridDim.x;
- // NOTE: Plain declaration such as
- // extern __shared__ RealType shx[];
- // won't work because extern variables must be declared exactly once.
- // In templated functions we need to have same variable name with
- // different type, which causes the conflict. In CUDA samples they
- // solve it using template specialization via classes, but using char
- // as the base type and reinterpret_cast works too.
- // See http://stackoverflow.com/a/19339004/4180822
- extern __shared__ __align__ ( 8 ) char __sdata[];
- RealType* shx = reinterpret_cast< RealType* >( __sdata );
+ RealType* shx = Devices::Cuda::getSharedMemory< RealType >();
if( threadIdx.x < n )
shx[ threadIdx.x ] = x[ threadIdx.x ];
@@ -161,7 +152,7 @@ public:
#ifdef HAVE_CUDA
Containers::Vector< RealType, Devices::Cuda, IndexType > xDevice;
xDevice.setSize( n );
- if( ! Containers::ArrayOperations< Devices::Cuda, Devices::Host >::copyMemory< RealType, RealType, IndexType >( xDevice.getData(), x, n ) )
+ if( ! Containers::Algorithms::ArrayOperations< Devices::Cuda, Devices::Host >::copyMemory< RealType, RealType, IndexType >( xDevice.getData(), x, n ) )
throw 1;
dim3 blockSize( 256 );
diff --git a/src/TNL/Matrices/MatrixReader_impl.h b/src/TNL/Matrices/MatrixReader_impl.h
index f7ada94abcdd9d7d7867604f38731f32b59cb7bb..87ddf93d3b7bfa665d5f0e441c6c74a99294857d 100644
--- a/src/TNL/Matrices/MatrixReader_impl.h
+++ b/src/TNL/Matrices/MatrixReader_impl.h
@@ -11,7 +11,7 @@
#pragma once
#include <iomanip>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/String.h>
#include <TNL/Containers/Vector.h>
#include <TNL/Timer.h>
@@ -162,7 +162,7 @@ template< typename Matrix >
bool MatrixReader< Matrix >::checkMtxHeader( const String& header,
bool& symmetric )
{
- List< String > parsedLine;
+ Containers::List< String > parsedLine;
header.parse( parsedLine );
if( parsedLine.getSize() < 5 )
return false;
@@ -208,7 +208,7 @@ bool MatrixReader< Matrix >::readMtxHeader( std::istream& file,
file.seekg( 0, std::ios::beg );
String line;
bool headerParsed( false );
- List< String > parsedLine;
+ Containers::List< String > parsedLine;
while( true )
{
line.getLine( file );
@@ -364,7 +364,7 @@ bool MatrixReader< Matrix >::parseMtxLineWithElement( const String& line,
IndexType& column,
RealType& value )
{
- List< String > parsedLine;
+ Containers::List< String > parsedLine;
line.parse( parsedLine );
if( parsedLine.getSize() != 3 )
{
diff --git a/src/TNL/Meshes/MeshDetails/traits/MeshSuperentityTraits.h b/src/TNL/Meshes/MeshDetails/traits/MeshSuperentityTraits.h
index 37c0ff6caaf4ee377b6dda679cbe93d40c844043..1f2739e3fd4da1b51b498e6b56e4c9bd8ae24790 100644
--- a/src/TNL/Meshes/MeshDetails/traits/MeshSuperentityTraits.h
+++ b/src/TNL/Meshes/MeshDetails/traits/MeshSuperentityTraits.h
@@ -12,7 +12,7 @@
#include <TNL/Containers/Array.h>
#include <TNL/Containers/ConstSharedArray.h>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/Meshes/MeshEntity.h>
#include <TNL/Meshes/MeshConfigBase.h>
#include <TNL/Meshes/Topologies/MeshEntityTopology.h>
@@ -60,7 +60,7 @@ class MeshSuperentityTraits
/****
* This is used by the mesh initializer.
*/
- typedef List< GlobalIndexType > GrowableContainerType;
+ typedef Containers::List< GlobalIndexType > GrowableContainerType;
};
diff --git a/src/TNL/Object.cpp b/src/TNL/Object.cpp
index 06a819137a4f1f9bcbe3603263a9fd2762a317e0..c5350b287a067fc961a074aba910e331fde1847c 100644
--- a/src/TNL/Object.cpp
+++ b/src/TNL/Object.cpp
@@ -10,8 +10,6 @@
#include <TNL/Object.h>
#include <TNL/Assert.h>
-#include <TNL/File.h>
-#include <TNL/List.h>
#include <iostream>
#include <fstream>
#include <cstring>
@@ -171,7 +169,7 @@ bool getObjectType( const String& fileName, String& type )
}
bool parseObjectType( const String& objectType,
- List< String >& parsedObjectType )
+ Containers::List< String >& parsedObjectType )
{
parsedObjectType.reset();
int objectTypeLength = objectType. getLength();
diff --git a/src/TNL/Object.h b/src/TNL/Object.h
index d05c7595a4a9ee0eaa11be299440a1d247122e22..ed0ac90de971773bf52b01bb8715b8f43e425908 100644
--- a/src/TNL/Object.h
+++ b/src/TNL/Object.h
@@ -12,13 +12,11 @@
#include <TNL/Devices/Cuda.h>
#include <TNL/String.h>
-
+#include <TNL/File.h>
+#include <TNL/Containers/List.h>
namespace TNL {
-class File;
-template< class T > class List;
-
//! This is basic class for all 'large' objects like matrices, meshes, grids, solvers etc.
/*!
* Objects like numerical grids, meshes, matrices large vectors etc.
@@ -92,6 +90,6 @@ bool getObjectType( File& file, String& type );
bool getObjectType( const String& file_name, String& type );
bool parseObjectType( const String& objectType,
- List< String >& parsedObjectType );
+ Containers::List< String >& parsedObjectType );
} // namespace TNL
diff --git a/src/TNL/Solvers/Linear/CWYGMRES_impl.h b/src/TNL/Solvers/Linear/CWYGMRES_impl.h
index 988e7b9ddc0a844fc12ed2d1927a29b999fc0b80..5273f001bdd09294c4d976d40ef516022b2b5318 100644
--- a/src/TNL/Solvers/Linear/CWYGMRES_impl.h
+++ b/src/TNL/Solvers/Linear/CWYGMRES_impl.h
@@ -434,8 +434,8 @@ hauseholder_apply_trunc( HostVector& out,
// here we duplicate the upper (m+1)x(m+1) submatrix of Y on host for fast access
RealType* host_yi = &YL[ i * (restarting + 1) ];
RealType host_z[ i + 1 ];
- if( ! Containers::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( host_yi, y_i.getData(), restarting + 1 ) ||
- ! Containers::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( host_z, z.getData(), i + 1 ) )
+ if( ! Containers::Algorithms::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( host_yi, y_i.getData(), restarting + 1 ) ||
+ ! Containers::Algorithms::ArrayOperations< Devices::Host, Devices::Cuda >::copyMemory< RealType, RealType, IndexType >( host_z, z.getData(), i + 1 ) )
{
std::cerr << "Failed to copy part of device vectors y_i or z to host buffer." << std::endl;
throw 1;
diff --git a/src/TNL/Solvers/MeshTypeResolver.h b/src/TNL/Solvers/MeshTypeResolver.h
index c0dae65549302b7cfc9b8d110a282c9072aafadc..604b2dfbe70e4286a6ba5bcb3d00fa39cf7874cc 100644
--- a/src/TNL/Solvers/MeshTypeResolver.h
+++ b/src/TNL/Solvers/MeshTypeResolver.h
@@ -51,20 +51,20 @@ class MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >
protected:
static bool resolveMeshDimensions( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for disabled dimensions
template< int MeshDimensions,
typename = typename std::enable_if< ! ConfigTagDimensions<ConfigTag,MeshDimensions>::enabled >::type,
typename = void >
static bool resolveMeshRealType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for enabled dimensions
template< int MeshDimensions,
typename = typename std::enable_if< ConfigTagDimensions<ConfigTag,MeshDimensions>::enabled >::type >
static bool resolveMeshRealType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for disabled real types
template< int MeshDimensions,
@@ -72,14 +72,14 @@ class MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >
typename = typename std::enable_if< ! ConfigTagReal<ConfigTag, MeshRealType>::enabled >::type,
typename = void >
static bool resolveMeshIndexType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for enabled real types
template< int MeshDimensions,
typename MeshRealType,
typename = typename std::enable_if< ConfigTagReal<ConfigTag, MeshRealType>::enabled >::type >
static bool resolveMeshIndexType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for disabled index types
template< int MeshDimensions,
@@ -88,7 +88,7 @@ class MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >
typename = typename std::enable_if< ! ConfigTagIndex<ConfigTag, MeshIndexType>::enabled >::type,
typename = void >
static bool resolveMeshType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
// Overload for enabled index types
template< int MeshDimensions,
@@ -96,7 +96,7 @@ class MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >
typename MeshIndexType,
typename = typename std::enable_if< ConfigTagIndex<ConfigTag, MeshIndexType>::enabled >::type >
static bool resolveMeshType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
@@ -115,7 +115,7 @@ class MeshTypeResolverDimensionsSupportChecker< Dimensions, true, MeshTypeResolv
public:
static bool checkDimensions( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
};
template< int Dimensions, typename MeshTypeResolver >
@@ -124,7 +124,7 @@ class MeshTypeResolverDimensionsSupportChecker< Dimensions, false, MeshTypeResol
public:
static bool checkDimensions( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType );
+ const Containers::List< String >& parsedMeshType );
};*/
} // namespace Solvers
diff --git a/src/TNL/Solvers/MeshTypeResolver_impl.h b/src/TNL/Solvers/MeshTypeResolver_impl.h
index 2ef00775bac59eba58b6ecb672c7cf3ba6b17d04..d1b1b6d4263387a71f1192c2391b1c8f5217c3d0 100644
--- a/src/TNL/Solvers/MeshTypeResolver_impl.h
+++ b/src/TNL/Solvers/MeshTypeResolver_impl.h
@@ -58,8 +58,8 @@ bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::ru
std::cerr << "I am not able to detect the mesh type from the file " << meshFileName << "." << std::endl;
return EXIT_FAILURE;
}
- std::cout << meshType << " detected in " << meshFileName << " file." << std::endl;
- List< String > parsedMeshType;
+ std::cout << meshType << " detected in " << meshFileName << " file." << std::endl;
+ Containers::List< String > parsedMeshType;
if( ! parseObjectType( meshType, parsedMeshType ) )
{
std::cerr << "Unable to parse the mesh type " << meshType << "." << std::endl;
@@ -73,8 +73,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
typename Device,
typename Index,
typename ConfigTag >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshDimensions( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshDimensions( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
int dimensions = atoi( parsedMeshType[ 1 ].getString() );
@@ -94,8 +96,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
typename Index,
typename ConfigTag >
template< int MeshDimensions, typename, typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshRealType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshRealType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
std::cerr << "Mesh dimension " << MeshDimensions << " is not supported." << std::endl;
return false;
@@ -107,8 +111,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
typename Index,
typename ConfigTag >
template< int MeshDimensions, typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshRealType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshRealType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
if( parsedMeshType[ 2 ] == "float" )
return resolveMeshIndexType< MeshDimensions, float >( parameters, parsedMeshType );
@@ -128,8 +134,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
template< int MeshDimensions,
typename MeshRealType,
typename, typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshIndexType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshIndexType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
std::cerr << "The type '" << parsedMeshType[ 4 ] << "' is not allowed for real type." << std::endl;
return false;
@@ -143,8 +151,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
template< int MeshDimensions,
typename MeshRealType,
typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshIndexType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshIndexType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
if( parsedMeshType[ 4 ] == "short int" )
return resolveMeshType< MeshDimensions, MeshRealType, short int >( parameters, parsedMeshType );
@@ -165,8 +175,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
typename MeshRealType,
typename MeshIndexType,
typename, typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
std::cerr << "The type '" << parsedMeshType[ 4 ] << "' is not allowed for indexing type." << std::endl;
return false;
@@ -181,8 +193,10 @@ template< template< typename Real, typename Device, typename Index, typename Mes
typename MeshRealType,
typename MeshIndexType,
typename >
-bool MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::resolveMeshType( const Config::ParameterContainer& parameters,
- const List< String >& parsedMeshType )
+bool
+MeshTypeResolver< ProblemSetter, Real, Device, Index, ConfigTag, true >::
+resolveMeshType( const Config::ParameterContainer& parameters,
+ const Containers::List< String >& parsedMeshType )
{
if( parsedMeshType[ 0 ] == "Meshes::Grid" )
{
diff --git a/src/TNL/String.cpp b/src/TNL/String.cpp
index 723f024a22e623facf03a9994bfe55f81ac6d6c9..a2bf9ff9c9583d48162729ff3d0093cfd6bab9e9 100644
--- a/src/TNL/String.cpp
+++ b/src/TNL/String.cpp
@@ -13,7 +13,7 @@
#include <assert.h>
#include <TNL/String.h>
#include <TNL/Assert.h>
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/File.h>
#include <TNL/Math.h>
#ifdef HAVE_MPI
@@ -421,7 +421,7 @@ bool String :: getLine( std::istream& stream )
return true;
}
-int String :: parse( List< String >& list, const char separator ) const
+int String :: parse( Containers::List< String >& list, const char separator ) const
{
list.reset();
String copy( *this );
diff --git a/src/TNL/String.h b/src/TNL/String.h
index 45d94d13b954f93c276320600b442cc872ba9df4..254c283e2cc259af4056fddfb0446ef82e4af5cc 100644
--- a/src/TNL/String.h
+++ b/src/TNL/String.h
@@ -18,8 +18,10 @@
namespace TNL {
-template< class T > class List;
class File;
+namespace Containers {
+ template< class T > class List;
+}
//! Class for managing strings
class String
@@ -147,7 +149,7 @@ class String
bool getLine( std::istream& stream );
//! Parse the string into list of strings w.r.t. given separator.
- int parse( List< String >& list, const char separator = ' ' ) const;
+ int parse( Containers::List< String >& list, const char separator = ' ' ) const;
friend std::ostream& operator << ( std::ostream& stream, const String& str );
};
diff --git a/src/TNL/legacy/benchmarks/matrix-solvers-benchmark.h b/src/TNL/legacy/benchmarks/matrix-solvers-benchmark.h
index c3784b64e7b2e621da62b02775974bbf0ef7f68c..33153a50ff0e192db596897a7b701ace99171f68 100644
--- a/src/TNL/legacy/benchmarks/matrix-solvers-benchmark.h
+++ b/src/TNL/legacy/benchmarks/matrix-solvers-benchmark.h
@@ -357,7 +357,7 @@ int main( int argc, char* argv[] )
std::cerr << "Unable to detect object type in " << inputFile << std::endl;
return EXIT_FAILURE;
}
- List< String > parsedObjectType;
+ Containers::List< String > parsedObjectType;
parseObjectType( objectType,
parsedObjectType );
String objectClass = parsedObjectType[ 0 ];
diff --git a/src/Tools/tnl-dicom-reader.cpp b/src/Tools/tnl-dicom-reader.cpp
index f7e3ba24322a139d6556eafb7fbc498e354e3f2c..5b38b123092ccadf5c1d0fab1df842e7d258640a 100644
--- a/src/Tools/tnl-dicom-reader.cpp
+++ b/src/Tools/tnl-dicom-reader.cpp
@@ -38,7 +38,7 @@ bool processDicomFiles( const Config::ParameterContainer& parameters )
bool processDicomSeries( const Config::ParameterContainer& parameters )
{
- const List< String >& dicomSeriesNames = parameters.getParameter< List< String > >( "dicom-series" );
+ const Containers::List< String >& dicomSeriesNames = parameters.getParameter< Containers::List< String > >( "dicom-series" );
String meshFile = parameters.getParameter< String >( "mesh-file" );
bool verbose = parameters.getParameter< bool >( "verbose" );
diff --git a/src/Tools/tnl-diff.cpp b/src/Tools/tnl-diff.cpp
index 3e1155ce722d0c40142e99889addbc26365b6c56..76c2e792bc5e7b2a8f09304ac038334506c7e586 100644
--- a/src/Tools/tnl-diff.cpp
+++ b/src/Tools/tnl-diff.cpp
@@ -15,7 +15,7 @@
void setupConfig( Config::ConfigDescription& config )
{
config.addEntry< String >( "mesh", "Input mesh file.", "mesh.tnl" );
- config.addRequiredEntry< List< String > >( "input-files", "The first set of the input files." );
+ config.addRequiredEntry< Containers::List< String > >( "input-files", "The first set of the input files." );
config.addEntry< String >( "output-file", "File for the output data.", "tnl-diff.log" );
config.addEntry< String >( "mode", "Mode 'couples' compares two subsequent files. Mode 'sequence' compares the input files against the first one. 'halves' compares the files from the and the second half of the intput files.", "couples" );
config.addEntryEnum< String >( "couples" );
@@ -56,7 +56,7 @@ int main( int argc, char* argv[] )
return EXIT_FAILURE;
}
std::cout << meshType << " detected in " << meshFile << " file." << std::endl;
- List< String > parsedMeshType;
+ Containers::List< String > parsedMeshType;
if( ! parseObjectType( meshType, parsedMeshType ) )
{
std::cerr << "Unable to parse the mesh type " << meshType << "." << std::endl;
diff --git a/src/Tools/tnl-diff.h b/src/Tools/tnl-diff.h
index 0d6e1a0f7b2012093e90bc03244c88b1c3770468..73036199d0dcb70c3d32ad1fcfae7d1f4e68318e 100644
--- a/src/Tools/tnl-diff.h
+++ b/src/Tools/tnl-diff.h
@@ -24,7 +24,7 @@ template< typename MeshPointer, typename Element, typename Real, typename Index
bool computeDifferenceOfMeshFunctions( const MeshPointer& meshPointer, const Config::ParameterContainer& parameters )
{
bool verbose = parameters. getParameter< bool >( "verbose" );
- List< String > inputFiles = parameters. getParameter< List< String > >( "input-files" );
+ Containers::List< String > inputFiles = parameters. getParameter< Containers::List< String > >( "input-files" );
String mode = parameters. getParameter< String >( "mode" );
String outputFileName = parameters. getParameter< String >( "output-file" );
double snapshotPeriod = parameters. getParameter< double >( "snapshot-period" );
@@ -161,7 +161,7 @@ template< typename MeshPointer, typename Element, typename Real, typename Index
bool computeDifferenceOfVectors( const MeshPointer& meshPointer, const Config::ParameterContainer& parameters )
{
bool verbose = parameters. getParameter< bool >( "verbose" );
- List< String > inputFiles = parameters. getParameter< List< String > >( "input-files" );
+ Containers::List< String > inputFiles = parameters. getParameter< Containers::List< String > >( "input-files" );
String mode = parameters. getParameter< String >( "mode" );
String outputFileName = parameters. getParameter< String >( "output-file" );
double snapshotPeriod = parameters. getParameter< double >( "snapshot-period" );
@@ -307,7 +307,7 @@ bool computeDifference( const MeshPointer& meshPointer, const String& objectType
template< typename MeshPointer, typename Element, typename Real >
bool setIndexType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
String indexType;
@@ -335,8 +335,8 @@ bool setIndexType( const MeshPointer& meshPointer,
template< typename MeshPointer >
bool setTupleType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
- const List< String >& parsedElementType,
+ const Containers::List< String >& parsedObjectType,
+ const Containers::List< String >& parsedElementType,
const Config::ParameterContainer& parameters )
{
int dimensions = atoi( parsedElementType[ 1 ].getString() );
@@ -386,7 +386,7 @@ bool setTupleType( const MeshPointer& meshPointer,
template< typename MeshPointer >
bool setElementType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
String elementType;
@@ -410,7 +410,7 @@ bool setElementType( const MeshPointer& meshPointer,
return setIndexType< MeshPointer, double, double >( meshPointer, inputFileName, parsedObjectType, parameters );
if( elementType == "long double" )
return setIndexType< MeshPointer, long double, long double >( meshPointer, inputFileName, parsedObjectType, parameters );
- List< String > parsedElementType;
+ Containers::List< String > parsedElementType;
if( ! parseObjectType( elementType, parsedElementType ) )
{
std::cerr << "Unable to parse object type " << elementType << "." << std::endl;
@@ -427,7 +427,7 @@ template< typename Mesh >
bool processFiles( const Config::ParameterContainer& parameters )
{
int verbose = parameters. getParameter< int >( "verbose");
- List< String > inputFiles = parameters. getParameter< List< String > >( "input-files" );
+ Containers::List< String > inputFiles = parameters. getParameter< Containers::List< String > >( "input-files" );
String& inputFile = inputFiles[ 0 ];
/****
@@ -454,7 +454,7 @@ bool processFiles( const Config::ParameterContainer& parameters )
if( verbose )
std::cout << objectType << " detected ... ";
- List< String > parsedObjectType;
+ Containers::List< String > parsedObjectType;
if( ! parseObjectType( objectType, parsedObjectType ) )
{
std::cerr << "Unable to parse object type " << objectType << "." << std::endl;
diff --git a/src/Tools/tnl-image-converter.cpp b/src/Tools/tnl-image-converter.cpp
index 3d8f59bc9b6365a375662a739006b1d1ce42570e..0cece9cf7836e410ff55e089568d6f0873f05692 100644
--- a/src/Tools/tnl-image-converter.cpp
+++ b/src/Tools/tnl-image-converter.cpp
@@ -37,7 +37,7 @@ void configSetup( Config::ConfigDescription& config )
bool processImages( const Config::ParameterContainer& parameters )
{
- const List< String >& inputImages = parameters.getParameter< List< String > >( "input-images" );
+ const Containers::List< String >& inputImages = parameters.getParameter< Containers::List< String > >( "input-images" );
String meshFile = parameters.getParameter< String >( "mesh-file" );
bool verbose = parameters.getParameter< bool >( "verbose" );
@@ -133,7 +133,7 @@ bool processImages( const Config::ParameterContainer& parameters )
bool processTNLFiles( const Config::ParameterContainer& parameters )
{
- const List< String >& inputFiles = parameters.getParameter< List< String > >( "input-files" );
+ const Containers::List< String >& inputFiles = parameters.getParameter< Containers::List< String > >( "input-files" );
const String& imageFormat = parameters.getParameter< String >( "image-format" );
String meshFile = parameters.getParameter< String >( "mesh-file" );
bool verbose = parameters.getParameter< bool >( "verbose" );
diff --git a/src/Tools/tnl-init.cpp b/src/Tools/tnl-init.cpp
index e61afd2f627cbfc7f65f59e0f48d304c3167db72..90870bad68daa3761c20c44b2d04e24ad22c5677 100644
--- a/src/Tools/tnl-init.cpp
+++ b/src/Tools/tnl-init.cpp
@@ -61,7 +61,7 @@ int main( int argc, char* argv[] )
return EXIT_FAILURE;
}
std::cout << meshType << " detected in " << meshFile << " file." << std::endl;
- List< String > parsedMeshType;
+ Containers::List< String > parsedMeshType;
if( ! parseObjectType( meshType, parsedMeshType ) )
{
std::cerr << "Unable to parse the mesh type " << meshType << "." << std::endl;
diff --git a/src/Tools/tnl-init.h b/src/Tools/tnl-init.h
index ffe12126468cfca85bb05e4e0e4ca2585c69498c..98f0e461751a955f73a92bee419f29917bc7b3a5 100644
--- a/src/Tools/tnl-init.h
+++ b/src/Tools/tnl-init.h
@@ -201,7 +201,7 @@ bool resolveRealType( const Config::ParameterContainer& parameters )
template< int Dimensions, typename RealType, typename IndexType >
-bool resolveMesh( const List< String >& parsedMeshType,
+bool resolveMesh( const Containers::List< String >& parsedMeshType,
const Config::ParameterContainer& parameters )
{
std::cout << "+ -> Setting mesh type to " << parsedMeshType[ 0 ] << " ... " << std::endl;
@@ -216,7 +216,7 @@ bool resolveMesh( const List< String >& parsedMeshType,
}
template< int Dimensions, typename RealType >
-bool resolveIndexType( const List< String >& parsedMeshType,
+bool resolveIndexType( const Containers::List< String >& parsedMeshType,
const Config::ParameterContainer& parameters )
{
std::cout << "+ -> Setting index type to " << parsedMeshType[ 4 ] << " ... " << std::endl;
@@ -230,7 +230,7 @@ bool resolveIndexType( const List< String >& parsedMeshType,
}
template< int Dimensions >
-bool resolveRealType( const List< String >& parsedMeshType,
+bool resolveRealType( const Containers::List< String >& parsedMeshType,
const Config::ParameterContainer& parameters )
{
std::cout << "+ -> Setting real type to " << parsedMeshType[ 2 ] << " ... " << std::endl;
@@ -246,7 +246,7 @@ bool resolveRealType( const List< String >& parsedMeshType,
return false;
}
-bool resolveMeshType( const List< String >& parsedMeshType,
+bool resolveMeshType( const Containers::List< String >& parsedMeshType,
const Config::ParameterContainer& parameters )
{
std::cout << "+ -> Setting dimensions to " << parsedMeshType[ 1 ] << " ... " << std::endl;
diff --git a/src/Tools/tnl-view.cpp b/src/Tools/tnl-view.cpp
index f7d974821f6d702b6a23afa6bfe33e82d8d1703a..3f4e19304eda56866e0350e888c728aa8457b87f 100644
--- a/src/Tools/tnl-view.cpp
+++ b/src/Tools/tnl-view.cpp
@@ -71,7 +71,7 @@ int main( int argc, char* argv[] )
return EXIT_FAILURE;
}
std::cout << meshType << " detected in " << meshFile << " file." << std::endl;
- List< String > parsedMeshType;
+ Containers::List< String > parsedMeshType;
if( ! parseObjectType( meshType, parsedMeshType ) )
{
std::cerr << "Unable to parse the mesh type " << meshType << "." << std::endl;
diff --git a/src/Tools/tnl-view.h b/src/Tools/tnl-view.h
index c6754cd0c5dfba5a7c5614c066c9d775b285892c..60c366c4aa581f4cb107d12b1357aef61e966440 100644
--- a/src/Tools/tnl-view.h
+++ b/src/Tools/tnl-view.h
@@ -83,7 +83,7 @@ template< typename MeshPointer,
int EntityDimensions >
bool setMeshEntityType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
if( parsedObjectType[ 3 ] == "float" )
@@ -100,7 +100,7 @@ template< typename MeshReal,
typename MeshIndex >
bool setMeshEntityDimensions( const SharedPointer< Meshes::Grid< 1, MeshReal, Devices::Host, MeshIndex > >& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
typedef Meshes::Grid< 1, MeshReal, Devices::Host, MeshIndex > Mesh;
@@ -124,7 +124,7 @@ template< typename MeshReal,
typename MeshIndex >
bool setMeshEntityDimensions( const SharedPointer< Meshes::Grid< 2, MeshReal, Devices::Host, MeshIndex > >& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
typedef Meshes::Grid< 2, MeshReal, Devices::Host, MeshIndex > Mesh;
@@ -151,7 +151,7 @@ template< typename MeshReal,
typename MeshIndex >
bool setMeshEntityDimensions( const SharedPointer< Meshes::Grid< 3, MeshReal, Devices::Host, MeshIndex > >& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
typedef Meshes::Grid< 3, MeshReal, Devices::Host, MeshIndex > Mesh;
@@ -180,7 +180,7 @@ bool setMeshEntityDimensions( const SharedPointer< Meshes::Grid< 3, MeshReal, De
template< typename MeshPointer >
bool setMeshFunction( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
if( parsedObjectType[ 1 ] != meshPointer->getSerializationType() )
@@ -195,7 +195,7 @@ bool setMeshFunction( const MeshPointer& meshPointer,
template< typename MeshPointer, typename Element, typename Real, typename Index, int Dimensions >
bool convertObject( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
int verbose = parameters. getParameter< int >( "verbose");
@@ -248,7 +248,7 @@ bool convertObject( const MeshPointer& meshPointer,
template< typename MeshPointer, typename Element, typename Real, typename Index >
bool setDimensions( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
int dimensions( 0 );
@@ -276,7 +276,7 @@ bool setDimensions( const MeshPointer& meshPointer,
template< typename MeshPointer, typename Element, typename Real >
bool setIndexType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
String indexType;
@@ -300,8 +300,8 @@ bool setIndexType( const MeshPointer& meshPointer,
template< typename MeshPointer >
bool setTupleType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
- const List< String >& parsedElementType,
+ const Containers::List< String >& parsedObjectType,
+ const Containers::List< String >& parsedElementType,
const Config::ParameterContainer& parameters )
{
int dimensions = atoi( parsedElementType[ 1 ].getString() );
@@ -351,7 +351,7 @@ bool setTupleType( const MeshPointer& meshPointer,
template< typename MeshPointer >
bool setElementType( const MeshPointer& meshPointer,
const String& inputFileName,
- const List< String >& parsedObjectType,
+ const Containers::List< String >& parsedObjectType,
const Config::ParameterContainer& parameters )
{
String elementType;
@@ -373,16 +373,15 @@ bool setElementType( const MeshPointer& meshPointer,
return setIndexType< MeshPointer, double, double >( meshPointer, inputFileName, parsedObjectType, parameters );
if( elementType == "long double" )
return setIndexType< MeshPointer, long double, long double >( meshPointer, inputFileName, parsedObjectType, parameters );
- List< String > parsedElementType;
+ Containers::List< String > parsedElementType;
if( ! parseObjectType( elementType, parsedElementType ) )
{
std::cerr << "Unable to parse object type " << elementType << "." << std::endl;
return false;
}
- // FIXME: this does not compile for an unknown reason
-// if( parsedElementType[ 0 ] == "Containers::StaticVector" ||
-// parsedElementType[ 0 ] == "tnlStaticVector" ) // TODO: remove deprecated type names
-// return setTupleType< MeshPointer >( meshPointer, inputFileName, parsedObjectType, parsedElementType, parameters );
+ if( parsedElementType[ 0 ] == "Containers::StaticVector" ||
+ parsedElementType[ 0 ] == "tnlStaticVector" ) // TODO: remove deprecated type names
+ return setTupleType< MeshPointer >( meshPointer, inputFileName, parsedObjectType, parsedElementType, parameters );
std::cerr << "Unknown element type " << elementType << "." << std::endl;
return false;
@@ -406,7 +405,7 @@ bool processFiles( const Config::ParameterContainer& parameters )
meshPointer->writeMesh( "mesh.asy", "asymptote" );
bool checkOutputFile = parameters. getParameter< bool >( "check-output-file" );
- List< String > inputFiles = parameters. getParameter< List< String > >( "input-files" );
+ Containers::List< String > inputFiles = parameters. getParameter< Containers::List< String > >( "input-files" );
bool error( false );
//#ifdef HAVE_OPENMP
//#pragma omp parallel for
@@ -440,7 +439,7 @@ bool processFiles( const Config::ParameterContainer& parameters )
if( verbose )
std::cout << objectType << " detected ... ";
- List< String > parsedObjectType;
+ Containers::List< String > parsedObjectType;
if( ! parseObjectType( objectType, parsedObjectType ) )
{
std::cerr << "Unable to parse object type " << objectType << "." << std::endl;
diff --git a/src/Tools/tnlcurve2gnuplot.cpp b/src/Tools/tnlcurve2gnuplot.cpp
index baea729a13a36d63fc0d7bd86ac9287e91c35af4..16f7f372f37ea7881506aac6d7894036676a0256 100644
--- a/src/Tools/tnlcurve2gnuplot.cpp
+++ b/src/Tools/tnlcurve2gnuplot.cpp
@@ -36,9 +36,9 @@ int main( int argc, char* argv[] )
return 1;
}
- List< String > input_files = parameters. getParameter< List< String > >( "input-files" );
- List< String > output_files;
- if( ! parameters. getParameter< List< String > >( "output-files", output_files ) )
+ Containers::List< String > input_files = parameters. getParameter< Containers::List< String > >( "input-files" );
+ Containers::List< String > output_files;
+ if( ! parameters. getParameter< Containers::List< String > >( "output-files", output_files ) )
std::cout << "No output files were given." << std::endl;
int output_step( 1 );
parameters. getParameter< int >( "output-step", output_step );
diff --git a/src/UnitTests/Containers/ArrayOperationsTest.h b/src/UnitTests/Containers/ArrayOperationsTest.h
index 796f5b6e69f651c4755c55e1b0696550b14ac481..d11489d4d561abaf77acadcbe31870cb5ce63a6f 100644
--- a/src/UnitTests/Containers/ArrayOperationsTest.h
+++ b/src/UnitTests/Containers/ArrayOperationsTest.h
@@ -10,7 +10,7 @@
#pragma once
-#include <TNL/Containers/ArrayOperations.h>
+#include <TNL/Containers/Algorithms/ArrayOperations.h>
#include <TNL/Devices/Cuda.h>
#ifdef HAVE_GTEST
@@ -19,6 +19,7 @@
using namespace TNL;
using namespace TNL::Containers;
+using namespace TNL::Containers::Algorithms;
int getTestSize()
{
diff --git a/src/UnitTests/Containers/VectorOperationsTest.h b/src/UnitTests/Containers/VectorOperationsTest.h
index 6f7d7d46ce3d89b77309c34343af12e60af5b90f..08d18b1704ca0c23ee32928551b55a5feb96c772 100644
--- a/src/UnitTests/Containers/VectorOperationsTest.h
+++ b/src/UnitTests/Containers/VectorOperationsTest.h
@@ -19,6 +19,7 @@
using namespace TNL;
using namespace TNL::Containers;
+using namespace TNL::Containers::Algorithms;
#ifdef HAVE_GTEST
diff --git a/src/UnitTests/ListTest.cpp b/src/UnitTests/ListTest.cpp
index d14883f82f0f187000fe9ad17e9a9a01f361c761..6774d29f578e911d42e919d2dd8ff0b9f81a1610 100644
--- a/src/UnitTests/ListTest.cpp
+++ b/src/UnitTests/ListTest.cpp
@@ -12,7 +12,7 @@
#include "gtest/gtest.h"
#endif
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
using namespace TNL;
diff --git a/tests/benchmarks/spmv.h b/tests/benchmarks/spmv.h
index a197a15a73a0df9d2096dd6fc866b0a2b51e7036..2856b012f26aaf3dd079129bb04a29e93b3efdec 100644
--- a/tests/benchmarks/spmv.h
+++ b/tests/benchmarks/spmv.h
@@ -2,7 +2,7 @@
#include "benchmarks.h"
-#include <TNL/List.h>
+#include <TNL/Containers/List.h>
#include <TNL/Matrices/CSR.h>
#include <TNL/Matrices/Ellpack.h>
#include <TNL/Matrices/SlicedEllpack.h>
@@ -102,7 +102,7 @@ benchmarkSpMV( Benchmark & benchmark,
CudaVector deviceVector, deviceVector2;
// create benchmark group
- List< String > parsedType;
+ Containers::List< String > parsedType;
parseObjectType( HostMatrix::getType(), parsedType );
benchmark.createHorizontalGroup( parsedType[ 0 ], 2 );