gallo/atrip
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Lint and tidy up Equations

Browse Source
This commit is contained in:
Alejandro Gallo 2022-09-08 13:51:49 +02:00
parent 368c5619cc
commit 2c5a4620ca
1 changed files with 81 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

151
src/atrip/Equations.cxx
View File

@ -77,7 +77,7 @@ namespace cuda {
to[i] = maybeConjugateScalar<F>(from[i]); to[i] = maybeConjugateScalar<F>(from[i]);
} }
} }
template <typename F> template <typename F>
__global__ __global__
@ -132,30 +132,29 @@ namespace cuda {
#endif #endif
#if defined(HAVE_CUDA) #if defined(HAVE_CUDA)
#define LIMS_KS \ #define LIMS_KS() \
size_t \ size_t kmin = blockIdx.x * blockDim.x + threadIdx.x, \
kmin = blockIdx.x * blockDim.x + threadIdx.x, \ k = kmin, \
k = kmin, \ idx = kmin * size * size * size \
idx = kmin * size * size * size \ ; \
; \
k < (kmin < size) ? kmin + 1 : size k < (kmin < size) ? kmin + 1 : size
#else #else
#define LIMS_KS size_t k=0, idx=0; k < size #define LIMS_KS size_t k=0, idx=0; k < size
#endif #endif
#define _IJK_(i, j, k) i + j*size + k*size*size #define _IJK_(i, j, k) i + j*size + k*size*size
#define _REORDER_BODY_(...) \ #define _REORDER_BODY_(...) \
for (LIMS_KS ; k++) \ for (LIMS_KS() ; k++) \
for (size_t j = 0; j < size; j++) \ for (size_t j = 0; j < size; j++) \
for (size_t i = 0; i < size; i++, idx++) { \ for (size_t i = 0; i < size; i++, idx++) { \
__VA_ARGS__ \ __VA_ARGS__ \
} }
#define _MAKE_REORDER_(_enum, ...) \ #define _MAKE_REORDER_(_enum, ...) \
template <typename F> \ template <typename F> \
__MAYBE_GLOBAL__ \ __MAYBE_GLOBAL__ \
void reorder(reorder_proxy< F, _enum > p, \ void reorder(reorder_proxy< F, _enum > p, \
size_t size, F* to, F* from) { \ size_t size, F* to, F* from) { \
_REORDER_BODY_(__VA_ARGS__) \ _REORDER_BODY_(__VA_ARGS__) \
} }
#if defined(HAVE_CUDA) #if defined(HAVE_CUDA)
#define GO(__TO, __FROM) cuda::sum_in_place<F>(&__TO, &__FROM); #define GO(__TO, __FROM) cuda::sum_in_place<F>(&__TO, &__FROM);
#else #else
@ -166,7 +165,7 @@ namespace cuda {
template <typename F, reordering_t R> template <typename F, reordering_t R>
__MAYBE_GLOBAL__ \ __MAYBE_GLOBAL__ \
void reorder(reorder_proxy<F, R> proxy, void reorder(reorder_proxy<F, R> proxy,
size_t size, F* to, F* from); size_t size, F* to, F* from);
_MAKE_REORDER_(IJK, GO(to[idx], from[_IJK_(i, j, k)])) _MAKE_REORDER_(IJK, GO(to[idx], from[_IJK_(i, j, k)]))
_MAKE_REORDER_(IKJ, GO(to[idx], from[_IJK_(i, k, j)])) _MAKE_REORDER_(IKJ, GO(to[idx], from[_IJK_(i, k, j)]))
@ -174,7 +173,7 @@ namespace cuda {
_MAKE_REORDER_(JKI, GO(to[idx], from[_IJK_(j, k, i)])) _MAKE_REORDER_(JKI, GO(to[idx], from[_IJK_(j, k, i)]))
_MAKE_REORDER_(KIJ, GO(to[idx], from[_IJK_(k, i, j)])) _MAKE_REORDER_(KIJ, GO(to[idx], from[_IJK_(k, i, j)]))
_MAKE_REORDER_(KJI, GO(to[idx], from[_IJK_(k, j, i)])) _MAKE_REORDER_(KJI, GO(to[idx], from[_IJK_(k, j, i)]))
#undef LIMS_KS #undef LIMS_KS
#undef _MAKE_REORDER #undef _MAKE_REORDER
@ -446,50 +445,59 @@ double getEnergySame
#if defined(ATRIP_USE_DGEMM) #if defined(ATRIP_USE_DGEMM)
#if defined(HAVE_CUDA) #if defined(HAVE_CUDA)
#define REORDER(__II, __JJ, __KK) \ #define REORDER(__II, __JJ, __KK) \
reorder<<< \ reorder<<<bs, ths>>>(reorder_proxy<DataFieldType<F>, \
bs, ths \ __II ## __JJ ## __KK >{}, \
>>>(reorder_proxy<DataFieldType<F>, __II ## __JJ ## __KK >{}, \ No, Tijk, _t_buffer);
No, Tijk, _t_buffer); #define DGEMM_PARTICLES(__A, __B) \
#define DGEMM_PARTICLES(__A, __B) \ atrip::xgemm<F>("T", \
atrip::xgemm<F>("T", \ "N", \
"N", \ (int const*)&NoNo, \
(int const*)&NoNo, \ (int const*)&No, \
(int const*)&No, \ (int const*)&Nv, \
(int const*)&Nv, \ &one, \
&one, \ (DataFieldType<F>*)__A, \
(DataFieldType<F>*)__A, \ (int const*)&Nv, \
(int const*)&Nv, \ (DataFieldType<F>*)__B, \
(DataFieldType<F>*)__B, \ (int const*)&Nv, \
(int const*)&Nv, \ &zero, \
&zero, \ _t_buffer, \
_t_buffer, \ (int const*)&NoNo);
(int const*)&NoNo); #define DGEMM_HOLES(__A, __B, __TRANSB) \
#define DGEMM_HOLES(__A, __B, __TRANSB) \ atrip::xgemm<F>("N", \
atrip::xgemm<F>("N", \ __TRANSB, \
__TRANSB, \ (int const*)&NoNo, \
(int const*)&NoNo, \ (int const*)&No, \
(int const*)&No, \ (int const*)&No, \
(int const*)&No, \ &m_one, \
&m_one, \ __A, \
__A, \ (int const*)&NoNo, \
(int const*)&NoNo, \ (DataFieldType<F>*)__B, \
(DataFieldType<F>*)__B, \ (int const*)&No, \
(int const*)&No, \ &zero, \
&zero, \ _t_buffer, \
_t_buffer, \ (int const*)&NoNo \
(int const*)&NoNo \ );
);
#define MAYBE_CONJ(_conj, _buffer) \ #define MAYBE_CONJ(_conj, _buffer) \
cuda::maybeConjugate<<< \ cuda::maybeConjugate<<< \
Atrip::kernelDimensions.ooo.blocks, \ Atrip::kernelDimensions.ooo.blocks, \
Atrip::kernelDimensions.ooo.threads \ Atrip::kernelDimensions.ooo.threads \
>>>((DataFieldType<F>*)_conj, (DataFieldType<F>*)_buffer, NoNoNo); >>>((DataFieldType<F>*)_conj, (DataFieldType<F>*)_buffer, NoNoNo);
// END CUDA ////////////////////////////////////////////////////////////////////
#else #else
// NONCUDA //////////////////////////////////////////////////////////////////////
#define REORDER(__II, __JJ, __KK) \
reorder(reorder_proxy<DataFieldType<F>, __II ## __JJ ## __KK >{}, \ // NONCUDA /////////////////////////////////////////////////////////////////////
No, Tijk, _t_buffer);
#define REORDER(__II, __JJ, __KK) \
reorder(reorder_proxy<DataFieldType<F>, \
__II ## __JJ ## __KK >{}, \
No, Tijk, _t_buffer);
#define DGEMM_PARTICLES(__A, __B) \ #define DGEMM_PARTICLES(__A, __B) \
atrip::xgemm<F>("T", \ atrip::xgemm<F>("T", \
"N", \ "N", \
@ -502,7 +510,7 @@ double getEnergySame
__B, \ __B, \
(int const*)&Nv, \ (int const*)&Nv, \
&zero, \ &zero, \
_t_buffer, \ _t_buffer, \
(int const*)&NoNo \ (int const*)&NoNo \
); );
#define DGEMM_HOLES(__A, __B, __TRANSB) \ #define DGEMM_HOLES(__A, __B, __TRANSB) \
@ -517,8 +525,8 @@ double getEnergySame
__B, \ __B, \
(int const*)&No, \ (int const*)&No, \
&zero, \ &zero, \
_t_buffer, \ _t_buffer, \
(int const*)&NoNo \ (int const*)&NoNo \
); );
#define MAYBE_CONJ(_conj, _buffer) \ #define MAYBE_CONJ(_conj, _buffer) \
for (size_t __i = 0; __i < NoNoNo; ++__i) \ for (size_t __i = 0; __i < NoNoNo; ++__i) \
@ -550,9 +558,10 @@ double getEnergySame
// Set Tijk to zero // Set Tijk to zero
#ifdef HAVE_CUDA #ifdef HAVE_CUDA
WITH_CHRONO("double:reorder", WITH_CHRONO("double:reorder",
cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk, NoNoNo); cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk,
// synchronize all initializations to zero NoNoNo);
) // synchronize all initializations to zero
)
#else #else
WITH_CHRONO("double:reorder", WITH_CHRONO("double:reorder",
for (size_t k = 0; k < NoNoNo; k++) { for (size_t k = 0; k < NoNoNo; k++) {
@ -640,6 +649,8 @@ double getEnergySame
{ // free resources { // free resources
#ifdef HAVE_CUDA #ifdef HAVE_CUDA
// we need to synchronize here since we need
// the Tijk for next process in the pipeline
cuCtxSynchronize(); cuCtxSynchronize();
cuMemFree((CUdeviceptr)_vhhh); cuMemFree((CUdeviceptr)_vhhh);
cuMemFree((CUdeviceptr)_t_buffer); cuMemFree((CUdeviceptr)_t_buffer);