BiEll sym. (32ed06d9) · Commits · TNL / tnl-dev

src/implementation/matrices/tnlBiEllpackMatrix_impl.h

+0 −1

Original line number	Diff line number	Diff line
		@@ -1093,7 +1093,6 @@ void tnlBiEllpackMatrix< Real, Device, Index, StripSize >::spmvCuda( const InVec
		{
		if( this->columnIndexes[ elementPtr ] < this->getColumns() )
		temp[ threadIdx.x ] += inVector[ this->columnIndexes[ elementPtr ] ] * this->values[ elementPtr ];
		outVector.add( this->columnIndexes[ elementPtr ], inVector[ this->columnIndexes[ elementPtr ] ] * this->values[ elementPtr ] );
		elementPtr += this->warpSize;
		}
		IndexType bisection2 = this->warpSize;

src/implementation/matrices/tnlBiEllpackSymMatrix_impl.h

0 → 100644

+1634 −0

File added.

Preview size limit exceeded, changes collapsed.

src/matrices/tnlBiEllpackSymMatrix.h

0 → 100644

+188 −0

Original line number	Diff line number	Diff line
		#ifndef tnlBiEllpackSymMatrix_H_
		#define tnlBiEllpackSymMatrix_H_

		#include <matrices/tnlSparseMatrix.h>
		#include <core/vectors/tnlVector.h>

		template< typename Device >
		class tnlBiEllpackSymMatrixDeviceDependentCode;

		template< typename Real, typename Device = tnlCuda, typename Index = int, int StripSize = 32 >
		class tnlBiEllpackSymMatrix : public tnlSparseMatrix< Real, Device, Index >
		{
		public:
		typedef Real RealType;
		typedef Device DeviceType;
		typedef Index IndexType;
		typedef typename tnlSparseMatrix< RealType, DeviceType, IndexType >::RowLengthsVector RowLengthsVector;
		typedef typename tnlSparseMatrix< RealType, DeviceType, IndexType >::ValuesVector ValuesVector;
		typedef typename tnlSparseMatrix< RealType, DeviceType, IndexType >::ColumnIndexesVector ColumnIndexesVector;
		typedef tnlBiEllpackSymMatrix< Real, Device, Index > ThisType;
		typedef tnlBiEllpackSymMatrix< Real, tnlHost, Index > HostType;
		typedef tnlBiEllpackSymMatrix< Real, tnlCuda, Index > CudaType;

		tnlBiEllpackSymMatrix();

		static tnlString getType();

		tnlString getTypeVirtual() const;

		bool setDimensions( const IndexType rows,
		const IndexType columns );

		bool setRowLengths( const RowLengthsVector& rowLengths );

		IndexType getRowLength( const IndexType row ) const;

		template< typename Real2,
		typename Device2,
		typename Index2 >
		bool setLike( const tnlBiEllpackSymMatrix< Real2, Device2, Index2, StripSize >& matrix );

		void getRowLengths( tnlVector< IndexType, DeviceType, IndexType >& rowLengths ) const;

		bool setElement( const IndexType row,
		const IndexType column,
		const RealType& value );

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		bool setElementFast( const IndexType row,
		const IndexType column,
		const RealType& value );

		bool addElement( const IndexType row,
		const IndexType column,
		const RealType& value,
		const RealType& thisElementMultiplicator = 1.0 );

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		bool addElementFast( const IndexType row,
		const IndexType column,
		const RealType& value,
		const RealType& thisElementMultiplicator = 1.0 );

		bool setRow( const IndexType row,
		const IndexType* columns,
		const RealType* values,
		const IndexType numberOfElements );

		bool addRow( const IndexType row,
		const IndexType* columns,
		const RealType* values,
		const IndexType numberOfElements,
		const RealType& thisElementMultiplicator = 1.0 );

		RealType getElement( const IndexType row,
		const IndexType column ) const;

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		RealType getElementFast( const IndexType row,
		const IndexType column ) const;

		void getRow( const IndexType row,
		IndexType* columns,
		RealType* values ) const;

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		IndexType getGroupLength( const IndexType strip,
		const IndexType group ) const;

		template< typename InVector,
		typename OutVector >
		void vectorProduct( const InVector& inVector,
		OutVector& outVector ) const;

		template< typename InVector,
		typename OutVector >
		void vectorProductHost( const InVector& inVector,
		OutVector& outVector ) const;

		void setVirtualRows(const IndexType rows);

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		IndexType getNumberOfGroups( const IndexType row ) const;

		bool vectorProductTest() const;

		void reset();

		bool save( tnlFile& file ) const;

		bool load( tnlFile& file );

		bool save( const tnlString& fileName ) const;

		bool load( const tnlString& fileName );

		void print( ostream& str ) const;

		void performRowBubbleSort( tnlVector< Index, Device, Index >& tempRowLengths );
		void computeColumnSizes( tnlVector< Index, Device, Index >& tempRowLengths );

		// void verifyRowLengths( const typename tnlBiEllpackSymMatrix< Real, Device, Index, StripSize >::RowLengthsVector& rowLengths );

		template< typename InVector,
		typename OutVector >
		#ifdef HAVE_CUDA
		__device__
		#endif
		void spmvCuda( const InVector& inVector,
		OutVector& outVector,
		/*const IndexType warpStart,
		const IndexType inWarpIdx*/
		int globalIdx ) const;

		#ifdef HAVE_CUDA
		__device__ __host__
		#endif
		IndexType getStripLength( const IndexType strip ) const;

		#ifdef HAVE_CUDA
		__device__
		#endif
		void performRowBubbleSortCudaKernel( const typename tnlBiEllpackSymMatrix< Real, Device, Index, StripSize >::RowLengthsVector& rowLengths,
		const IndexType strip );

		#ifdef HAVE_CUDA
		__device__
		#endif
		void computeColumnSizesCudaKernel( const typename tnlBiEllpackSymMatrix< Real, Device, Index, StripSize >::RowLengthsVector& rowLengths,
		const IndexType numberOfStrips,
		const IndexType strip );

		#ifdef HAVE_CUDA
		__device__
		#endif
		IndexType power( const IndexType number,
		const IndexType exponent ) const;

		typedef tnlBiEllpackSymMatrixDeviceDependentCode< DeviceType > DeviceDependentCode;
		friend class tnlBiEllpackSymMatrixDeviceDependentCode< DeviceType >;

		private:

		IndexType warpSize;

		IndexType logWarpSize;

		IndexType virtualRows;

		tnlVector< Index, Device, Index > rowPermArray;

		tnlVector< Index, Device, Index > groupPointers;

		};

		#include <implementation/matrices/tnlBiEllpackSymMatrix_impl.h>

		#endif

tests/long-time-unit-tests/matrix-formats-test.h

+20 −4

Original line number	Diff line number	Diff line
		@@ -30,6 +30,8 @@
		#include <matrices/tnlSlicedEllpackMatrix.h>
		#include <matrices/tnlChunkedEllpackMatrix.h>
		#include <matrices/tnlCSRMatrix.h>
		#include <matrices/tnlBiEllpackMatrix.h>
		#include <matrices/tnlBiEllpackSymMatrix.h>

		void setupConfig( tnlConfigDescription& config )
		{
		@@ -41,6 +43,8 @@ void setupConfig( tnlConfigDescription& config )
		config.addEntryEnum< tnlString >( "sliced-ellpack" );
		config.addEntryEnum< tnlString >( "chunked-ellpack" );
		config.addEntryEnum< tnlString >( "csr" );
		config.addEntryEnum< tnlString >( "bi-ell" );
		config.addEntryEnum< tnlString >( "bi-ell-sym" );
		config.addEntry< bool >( "hard-test", "Comparison against the dense matrix.", false );
		config.addEntry< bool >( "multiplication-test", "Matrix-vector multiplication test.", false );
		config.addEntry< bool >( "verbose", "Verbose mode." );
		@@ -66,16 +70,16 @@ bool testMatrix( const tnlParameterContainer& parameters )
		}
		if( ! tnlMatrixReader< Matrix >::readMtxFile( file, matrix, verbose ) )
		return false;
		if( ! tnlMatrixReader< Matrix >::verifyMtxFile( file, matrix, verbose ) )
		return false;
		//if( ! tnlMatrixReader< Matrix >::verifyMtxFile( file, matrix, verbose ) )
		// return false;
		if( parameters.GetParameter< bool >( "hard-test" ) )
		{
		typedef tnlDenseMatrix< RealType, DeviceType, IndexType > DenseMatrix;
		DenseMatrix denseMatrix;
		if( ! tnlMatrixReader< DenseMatrix >::readMtxFile( file, denseMatrix, verbose ) )
		return false;
		if( ! tnlMatrixReader< DenseMatrix >::verifyMtxFile( file, denseMatrix, verbose ) )
		return false;
		//if( ! tnlMatrixReader< DenseMatrix >::verifyMtxFile( file, denseMatrix, verbose ) )
		// return false;
		//matrix.print( cout );
		//denseMatrix.print( cout );
		for( IndexType i = 0; i < matrix.getRows(); i++ )
		@@ -171,6 +175,18 @@ int main( int argc, char* argv[] )
		return EXIT_FAILURE;
		return EXIT_SUCCESS;
		}
		if( matrixFormat == "bi-ell" )
		{
		if( !testMatrix< tnlBiEllpackMatrix< double, tnlHost, int > >( parameters ) )
		return EXIT_FAILURE;
		return EXIT_SUCCESS;
		}
		if( matrixFormatt == "bi-ell-sym" )
		{
		if( !testMatrix< tnlBiEllpackSymMatrix< double, tnlHost, int > >( parameters ) )
		return EXIT_FAILURE;
		return EXIT_SUCCESS;
		}
		cerr << "Uknown matrix format " << matrixFormat << "." << endl;
		return EXIT_FAILURE;
		}