DelayedUpdateSYCL< T, T_FP > Class Template Reference

implements delayed update on Intel GPU using SYCL More...

Collaboration diagram for DelayedUpdateSYCL< T, T_FP >:

Public Member Functions
	DelayedUpdateSYCL ()
	default constructor More...
	~DelayedUpdateSYCL ()
void	resize (int norb, int delay)
	resize the internal storage More...
template<typename TREAL >
void	invert_transpose (const Matrix< T > &logdetT, Matrix< T > &Ainv, std::complex< TREAL > &log_value)
	compute the inverse of the transpose of matrix A and its determinant value in log More...
void	initializeInv (const Matrix< T > &Ainv)
	initialize internal objects when Ainv is refreshed More...
int	getDelayCount () const
template<typename VVT >
void	getInvRow (const Matrix< T > &Ainv, int rowchanged, VVT &invRow)
	compute the row of up-to-date Ainv More...
template<typename VVT , typename RATIOT >
void	acceptRow (Matrix< T > &Ainv, int rowchanged, const VVT &psiV, const RATIOT ratio_new)
	accept a move with the update delayed More...
void	updateInvMat (Matrix< T > &Ainv, bool transfer_to_host=true)
	update the full Ainv and reset delay_count More...

Private Member Functions
void	clearDelayCount ()
	reset delay count to 0 More...

Private Attributes
Matrix< T >	U
Matrix< T >	Binv
Matrix< T >	V
Matrix< T, SYCLAllocator< T > >	temp_gpu
Matrix< T, SYCLAllocator< T > >	U_gpu
	GPU copy of U, V, Binv, Ainv. More...
Matrix< T, SYCLAllocator< T > >	V_gpu
Matrix< T, SYCLAllocator< T > >	Binv_gpu
Matrix< T, SYCLAllocator< T > >	Ainv_gpu
Vector< T >	p
Vector< int, SYCLHostAllocator< int > >	delay_list
int	delay_count
	current number of delays, increase one for each acceptance, reset to 0 after updating Ainv More...
syclSolverInverter< T_FP >	sycl_inverter_
PrefetchedRange	prefetched_range
Matrix< T >	Ainv_buffer
sycl::queue	m_queue_

Detailed Description

template<typename T, typename T_FP>
class qmcplusplus::DelayedUpdateSYCL< T, T_FP >

implements delayed update on Intel GPU using SYCL

Template Parameters

T	base precision for most computation
T_FP	high precision for matrix inversion, T_FP >= T

Definition at line 35 of file DelayedUpdateSYCL.h.

Constructor & Destructor Documentation

DelayedUpdateSYCL()

DelayedUpdateSYCL ( )

inline

default constructor

Definition at line 73 of file DelayedUpdateSYCL.h.

References qmcplusplus::createSYCLInOrderQueueOnDefaultDevice(), and DelayedUpdateSYCL< T, T_FP >::m_queue_.

: delay_count(0) { m_queue_ = createSYCLInOrderQueueOnDefaultDevice(); }

~DelayedUpdateSYCL()

~DelayedUpdateSYCL ( )

inline

Definition at line 75 of file DelayedUpdateSYCL.h.

References qmcplusplus::syclSolver::freeBuffer().

{ syclSolver::freeBuffer(); }

Member Function Documentation

acceptRow()

void acceptRow ( Matrix< T > &  Ainv, int  rowchanged, const VVT &  psiV, const RATIOT  ratio_new  )

inline

accept a move with the update delayed

Parameters

Ainv	inverse matrix
rowchanged	the row id corresponding to the proposed electron
psiV	new orbital values

Before delay_count reaches the maximum delay, only Binv is updated with a recursive algorithm

Definition at line 161 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_buffer, DelayedUpdateSYCL< T, T_FP >::Binv, Matrix< T, Alloc >::cols(), BLAS::cone, qmcplusplus::syclBLAS::copy_n(), BLAS::czero, Matrix< T, Alloc >::data(), Vector< T, Alloc >::data(), DelayedUpdateSYCL< T, T_FP >::delay_count, DelayedUpdateSYCL< T, T_FP >::delay_list, BLAS::gemv(), BLAS::ger(), PrefetchedRange::getOffset(), DelayedUpdateSYCL< T, T_FP >::p, DelayedUpdateSYCL< T, T_FP >::prefetched_range, Matrix< T, Alloc >::rows(), DelayedUpdateSYCL< T, T_FP >::U, DelayedUpdateSYCL< T, T_FP >::updateInvMat(), and DelayedUpdateSYCL< T, T_FP >::V.

  {
    // update Binv from delay_count to delay_count+1
    constexpr T cone(1);
    constexpr T czero(0);
    const int norb     = Ainv.rows();
    const int lda_Binv = Binv.cols();
    std::copy_n(Ainv_buffer[prefetched_range.getOffset(rowchanged)], norb, V[delay_count]);
    std::copy_n(psiV.data(), norb, U[delay_count]);
    delay_list[delay_count] = rowchanged;
    // the new Binv is [[X Y] [Z sigma]]
    BLAS::gemv('T', norb, delay_count + 1, -cone, V.data(), norb, psiV.data(), 1, czero, p.data(), 1);
    // sigma
    const T sigma                  = static_cast<T>(RATIOT(1) / ratio_new);
    Binv[delay_count][delay_count] = sigma;
    // Y
    BLAS::gemv('T', delay_count, delay_count, sigma, Binv.data(), lda_Binv, p.data(), 1, czero,
               Binv.data() + delay_count, lda_Binv);
    // X
    BLAS::ger(delay_count, delay_count, cone, Binv[delay_count], 1, Binv.data() + delay_count, lda_Binv, Binv.data(),
              lda_Binv);
    // Z
    for (int i = 0; i < delay_count; i++)
      Binv[delay_count][i] *= sigma;
    delay_count++;
    // update Ainv when maximal delay is reached
    if (delay_count == lda_Binv)
      updateInvMat(Ainv, false);
  }

clearDelayCount()

void clearDelayCount ( )

inlineprivate

reset delay count to 0

Definition at line 65 of file DelayedUpdateSYCL.h.

References PrefetchedRange::clear(), DelayedUpdateSYCL< T, T_FP >::delay_count, and DelayedUpdateSYCL< T, T_FP >::prefetched_range.

Referenced by DelayedUpdateSYCL< T, T_FP >::initializeInv(), DelayedUpdateSYCL< T, T_FP >::invert_transpose(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

  {
    delay_count = 0;
    prefetched_range.clear();
  }

getDelayCount()

int getDelayCount ( ) const

inline

Definition at line 121 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::delay_count.

{ return delay_count; }

getInvRow()

void getInvRow ( const Matrix< T > &  Ainv, int  rowchanged, VVT &  invRow  )

inline

compute the row of up-to-date Ainv

Parameters

Ainv	inverse matrix
rowchanged	the row id corresponding to the proposed electron

Definition at line 128 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_buffer, DelayedUpdateSYCL< T, T_FP >::Ainv_gpu, DelayedUpdateSYCL< T, T_FP >::Binv, PrefetchedRange::checkRange(), Matrix< T, Alloc >::cols(), BLAS::cone, qmcplusplus::syclBLAS::copy_n(), BLAS::czero, Matrix< T, Alloc >::data(), Vector< T, Alloc >::data(), DelayedUpdateSYCL< T, T_FP >::delay_count, BLAS::gemv(), PrefetchedRange::getOffset(), DelayedUpdateSYCL< T, T_FP >::m_queue_, omptarget::min(), DelayedUpdateSYCL< T, T_FP >::p, DelayedUpdateSYCL< T, T_FP >::prefetched_range, Matrix< T, Alloc >::rows(), PrefetchedRange::setRange(), Matrix< T, Alloc >::size(), DelayedUpdateSYCL< T, T_FP >::U, and DelayedUpdateSYCL< T, T_FP >::V.

  {
    if (!prefetched_range.checkRange(rowchanged))
    {
      const int last_row = std::min(rowchanged + Ainv_buffer.rows(), Ainv.rows());
      m_queue_.memcpy(Ainv_buffer.data(), Ainv_gpu[rowchanged], invRow.size() * (last_row - rowchanged) * sizeof(T))
          .wait();
      prefetched_range.setRange(rowchanged, last_row);
    }

    // save AinvRow to new_AinvRow
    std::copy_n(Ainv_buffer[prefetched_range.getOffset(rowchanged)], invRow.size(), invRow.data());
    if (delay_count > 0)
    {
      constexpr T cone(1);
      constexpr T czero(0);
      const int norb     = Ainv.rows();
      const int lda_Binv = Binv.cols();
      // multiply V (NxK) Binv(KxK) U(KxN) AinvRow right to the left
      BLAS::gemv('T', norb, delay_count, cone, U.data(), norb, invRow.data(), 1, czero, p.data(), 1);
      BLAS::gemv('N', delay_count, delay_count, -cone, Binv.data(), lda_Binv, p.data(), 1, czero, Binv[delay_count], 1);
      BLAS::gemv('N', norb, delay_count, cone, V.data(), norb, Binv[delay_count], 1, cone, invRow.data(), 1);
    }
  }

initializeInv()

void initializeInv ( const Matrix< T > &  Ainv )

inline

initialize internal objects when Ainv is refreshed

Parameters

Ainv	inverse matrix

Definition at line 114 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_gpu, DelayedUpdateSYCL< T, T_FP >::clearDelayCount(), Matrix< T, Alloc >::data(), DelayedUpdateSYCL< T, T_FP >::m_queue_, and Matrix< T, Alloc >::size().

  {
    // must be blocking due to potential consumption of Ainv_gpu
    m_queue_.memcpy(Ainv_gpu.data(), Ainv.data(), Ainv.size() * sizeof(T)).wait();
    clearDelayCount();
  }

invert_transpose()

void invert_transpose ( const Matrix< T > &  logdetT, Matrix< T > &  Ainv, std::complex< TREAL > &  log_value  )

inline

compute the inverse of the transpose of matrix A and its determinant value in log

Template Parameters

TREAL

real type

Definition at line 104 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_gpu, DelayedUpdateSYCL< T, T_FP >::clearDelayCount(), DelayedUpdateSYCL< T, T_FP >::m_queue_, and DelayedUpdateSYCL< T, T_FP >::sycl_inverter_.

  {
    clearDelayCount();

    sycl_inverter_.invert_transpose(logdetT, Ainv, Ainv_gpu, log_value, m_queue_);
  }

resize()

void resize ( int  norb, int  delay  )

inline

resize the internal storage

Parameters

norb	number of electrons/orbitals
delay,maximum	delay 0<delay<=norb

Definition at line 81 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_buffer, DelayedUpdateSYCL< T, T_FP >::Ainv_gpu, DelayedUpdateSYCL< T, T_FP >::Binv, DelayedUpdateSYCL< T, T_FP >::Binv_gpu, DelayedUpdateSYCL< T, T_FP >::delay_list, omptarget::min(), DelayedUpdateSYCL< T, T_FP >::p, Matrix< T, Alloc >::resize(), Vector< T, Alloc >::resize(), DelayedUpdateSYCL< T, T_FP >::temp_gpu, DelayedUpdateSYCL< T, T_FP >::U, DelayedUpdateSYCL< T, T_FP >::U_gpu, DelayedUpdateSYCL< T, T_FP >::V, and DelayedUpdateSYCL< T, T_FP >::V_gpu.

  {
    //tempMat.resize(norb, delay);
    V.resize(delay, norb);
    U.resize(delay, norb);
    p.resize(delay);
    Binv.resize(delay, delay);
    // prefetch 8% more rows corresponding to roughly 96% acceptance ratio
    Ainv_buffer.resize(std::min(static_cast<int>(delay * 1.08), norb), norb);

    temp_gpu.resize(norb, delay);
    delay_list.resize(delay);
    U_gpu.resize(delay, norb);
    V_gpu.resize(delay, norb);
    Binv_gpu.resize(delay, delay);
    //delay_list_gpu.resize(delay);
    Ainv_gpu.resize(norb, norb);
  }

updateInvMat()

void updateInvMat ( Matrix< T > &  Ainv, bool  transfer_to_host = true  )

inline

update the full Ainv and reset delay_count

Parameters

Ainv	inverse matrix

Definition at line 194 of file DelayedUpdateSYCL.h.

References DelayedUpdateSYCL< T, T_FP >::Ainv_gpu, qmcplusplus::applyW_stageV_sycl(), DelayedUpdateSYCL< T, T_FP >::Binv, DelayedUpdateSYCL< T, T_FP >::Binv_gpu, DelayedUpdateSYCL< T, T_FP >::clearDelayCount(), Matrix< T, Alloc >::cols(), BLAS::cone, BLAS::czero, Matrix< T, Alloc >::data(), DelayedUpdateSYCL< T, T_FP >::delay_count, DelayedUpdateSYCL< T, T_FP >::delay_list, qmcplusplus::syclBLAS::gemm(), DelayedUpdateSYCL< T, T_FP >::m_queue_, Matrix< T, Alloc >::rows(), Matrix< T, Alloc >::size(), DelayedUpdateSYCL< T, T_FP >::temp_gpu, DelayedUpdateSYCL< T, T_FP >::U, DelayedUpdateSYCL< T, T_FP >::U_gpu, and DelayedUpdateSYCL< T, T_FP >::V_gpu.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow().

  {
    // update the inverse matrix
    if (delay_count > 0)
    {
      constexpr T cone(1);
      constexpr T czero(0);
      const int norb     = Ainv.rows();
      const int lda_Binv = Binv.cols();

      m_queue_.memcpy(U_gpu.data(), U.data(), norb * delay_count * sizeof(T));
      m_queue_.memcpy(Binv_gpu.data(), Binv.data(), lda_Binv * delay_count * sizeof(T));

      syclBLAS::gemm(m_queue_, 'T', 'N', delay_count, norb, norb, cone, U_gpu.data(), norb, Ainv_gpu.data(), norb,
                     czero, temp_gpu.data(), lda_Binv);

      applyW_stageV_sycl(m_queue_, delay_list.data(), delay_count, temp_gpu.data(), norb, temp_gpu.cols(), V_gpu.data(),
                         Ainv_gpu.data());

      syclBLAS::gemm(m_queue_, 'N', 'N', norb, delay_count, delay_count, cone, V_gpu.data(), norb, Binv_gpu.data(),
                     lda_Binv, czero, U_gpu.data(), norb);

#ifdef SYCL_BLOCKING
      syclBLAS::gemm(m_queue_, 'N', 'N', norb, norb, delay_count, -cone, U_gpu.data(), norb, temp_gpu.data(), lda_Binv,
                     cone, Ainv_gpu.data(), norb)
          .wait();
#else
      syclBLAS::gemm(m_queue_, 'N', 'N', norb, norb, delay_count, -cone, U_gpu.data(), norb, temp_gpu.data(), lda_Binv,
                     cone, Ainv_gpu.data(), norb);
#endif

      clearDelayCount();
    }

    // transfer Ainv_gpu to Ainv and wait till completion
    if (transfer_to_host)
      m_queue_.memcpy(Ainv.data(), Ainv_gpu.data(), Ainv.size() * sizeof(T)).wait();
  }

Member Data Documentation

Ainv_buffer

Matrix<T> Ainv_buffer

private

Definition at line 60 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), and DelayedUpdateSYCL< T, T_FP >::resize().

Ainv_gpu

Matrix<T, SYCLAllocator<T> > Ainv_gpu

private

Definition at line 47 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::getInvRow(), DelayedUpdateSYCL< T, T_FP >::initializeInv(), DelayedUpdateSYCL< T, T_FP >::invert_transpose(), DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

Binv

Matrix<T> Binv

private

Definition at line 39 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

Binv_gpu

Matrix<T, SYCLAllocator<T> > Binv_gpu

private

Definition at line 46 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

delay_count

int delay_count

private

current number of delays, increase one for each acceptance, reset to 0 after updating Ainv

Definition at line 53 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::clearDelayCount(), DelayedUpdateSYCL< T, T_FP >::getDelayCount(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

delay_list

Vector<int, SYCLHostAllocator<int> > delay_list

private

Definition at line 51 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

m_queue_

sycl::queue m_queue_

private

Definition at line 62 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::DelayedUpdateSYCL(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), DelayedUpdateSYCL< T, T_FP >::initializeInv(), DelayedUpdateSYCL< T, T_FP >::invert_transpose(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

p

Vector<T> p

private

Definition at line 49 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), and DelayedUpdateSYCL< T, T_FP >::resize().

prefetched_range

PrefetchedRange prefetched_range

private

Definition at line 58 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::clearDelayCount(), and DelayedUpdateSYCL< T, T_FP >::getInvRow().

sycl_inverter_

syclSolverInverter<T_FP> sycl_inverter_

private

Definition at line 55 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::invert_transpose().

temp_gpu

Matrix<T, SYCLAllocator<T> > temp_gpu

private

Definition at line 42 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

U

Matrix<T> U

private

Definition at line 38 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

U_gpu

Matrix<T, SYCLAllocator<T> > U_gpu

private

GPU copy of U, V, Binv, Ainv.

Definition at line 44 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

V

Matrix<T> V

private

Definition at line 40 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::acceptRow(), DelayedUpdateSYCL< T, T_FP >::getInvRow(), and DelayedUpdateSYCL< T, T_FP >::resize().

V_gpu

Matrix<T, SYCLAllocator<T> > V_gpu

private

Definition at line 45 of file DelayedUpdateSYCL.h.

Referenced by DelayedUpdateSYCL< T, T_FP >::resize(), and DelayedUpdateSYCL< T, T_FP >::updateInvMat().

---

The documentation for this class was generated from the following file:

• /home/pk7/projects/qmc/for_cron_doxygen/qmcpack/src/QMCWaveFunctions/Fermion/DelayedUpdateSYCL.h