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

Temporarily introduce a SIMT_loop to test out approaches prior to making a global change to

Browse Source 
accelerator_loop
This commit is contained in:
Peter Boyle 2019-06-08 13:44:27 +01:00
parent ad2c433574
commit c933ac2248
1 changed files with 4 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
Grid/qcd/action/fermion/implementation/WilsonKernelsGpuImplementation.h
View File

@ -36,13 +36,6 @@ NAMESPACE_BEGIN(Grid);
////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////
// Gpu implementation; thread loop is implicit ; move to header // Gpu implementation; thread loop is implicit ; move to header
////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////
accelerator_inline void synchronise(void)
{
#ifdef __CUDA_ARCH__
__syncthreads();
#endif
return;
}
accelerator_inline int get_my_lanes(int Nsimd) accelerator_inline int get_my_lanes(int Nsimd)
{ {
#ifdef __CUDA_ARCH__ #ifdef __CUDA_ARCH__
@ -281,10 +274,12 @@ GPU_EMPTY(GparityWilsonImplDF);
}); });
#define HOST_CALL(A) \ #define HOST_CALL(A) \
accelerator_loopN( ss, Ls*Nsite, { \ const uint64_t nsimd = Simd::Nsimd(); \
const uint64_t NN = Nsite*Ls; \
SIMT_loop( ss, NN, nsimd, { \
int sF = ss; \ int sF = ss; \
int sU = ss/Ls; \ int sU = ss/Ls; \
WilsonKernels<Impl>::A(st_v,U_v,st.CommBuf(),sF,sU,in_v,out_v); \ WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,in_v,out_v); \
}); });
template <class Impl> template <class Impl>