Commit 5f8e4c10 authored by Illia Kolesnik's avatar Illia Kolesnik Committed by Tomáš Oberhuber
Browse files

Changed number of elements per warp for CSR Adaptive, CSR Multivector and Improved CSR Light

parent 01ed593c

src/TNL/Matrices/Legacy/CSR.h

+5 −2
Original line number Diff line number Diff line
@@ -115,8 +115,11 @@ public: 
   static constexpr Index THREADS_VECTOR = 128;

   static constexpr Index THREADS_LIGHT = 128;



   /* Max length of row to process one warp */

   static constexpr Index MAX_ELEMENTS_PER_WARP = 1024;

   /* Max length of row to process one warp for CSR Light, MultiVector */

   static constexpr Index MAX_ELEMENTS_PER_WARP = 384;



   /* Max length of row to process one warp for CSR Adaptive */

   static constexpr Index MAX_ELEMENTS_PER_WARP_ADAPT = 512;



   /* How many shared memory use per block in CSR Adaptive kernel */

   static constexpr Index SHARED_PER_BLOCK = 24576;

src/TNL/Matrices/Legacy/CSR_impl.h

+3 −3
Original line number Diff line number Diff line
@@ -143,7 +143,7 @@ Index findLimit(const Index start, 
            type = Type::STREAM;

            return current;

         } else {                  // one long row

            if (sum <= 2 * matrix.MAX_ELEMENTS_PER_WARP)

            if (sum <= 2 * matrix.MAX_ELEMENTS_PER_WARP_ADAPT)

               type = Type::VECTOR;

            else

               type = Type::LONG;
@@ -1765,7 +1765,7 @@ void SpMVCSRAdaptivePrepare( const Real *inVector, 
      SpMVCSRAdaptive< Real, Index, warpSize,

            matrix.WARPS,

            matrix.SHARED_PER_WARP, 

            matrix.MAX_ELEMENTS_PER_WARP >

            matrix.MAX_ELEMENTS_PER_WARP_ADAPT >

         <<<blocks, threads>>>(

               inVector,

               outVector,