portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2025-04-09 21:50:45 +01:00
Code Releases Activity

Clean acceleerator barrier

Browse Source
This commit is contained in:
Peter Boyle 2019-06-15 12:51:45 +01:00
parent 8e394d3bf9
commit d763c303c5
1 changed files with 1 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
Grid/qcd/action/fermion/WilsonCompressor.h
View File

@ -68,20 +68,12 @@ public:
/*****************************************************/ /*****************************************************/
/* Compress includes precision change if mpi data is not same */ /* Compress includes precision change if mpi data is not same */
/*****************************************************/ /*****************************************************/
#if 0
accelerator_inline void Compress(SiteHalfSpinor *buf,Integer o,const SiteSpinor &in) {
SiteHalfSpinor tmp;
projector::Proj(tmp,in,mu,dag);
vstream(buf[o],tmp);
}
#else
template<class _SiteHalfSpinor, class _SiteSpinor> template<class _SiteHalfSpinor, class _SiteSpinor>
accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) { accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) {
_SiteHalfSpinor tmp; _SiteHalfSpinor tmp;
projector::Proj(tmp,in,mu,dag); projector::Proj(tmp,in,mu,dag);
vstream(buf[o],tmp); vstream(buf[o],tmp);
} }
#endif
/*****************************************************/ /*****************************************************/
/* Exchange includes precision change if mpi data is not same */ /* Exchange includes precision change if mpi data is not same */
@ -405,9 +397,7 @@ public:
this->face_table_computed=1; this->face_table_computed=1;
assert(this->u_comm_offset==this->_unified_buffer_size); assert(this->u_comm_offset==this->_unified_buffer_size);
this->halogtime+=usecond(); this->halogtime+=usecond();
#ifdef GRID_NVCC accelerator_barrier();
cudaDeviceSynchronize();
#endif
} }
}; };