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

Add support for only calculating DGEMM parts ATRIP_ONLY_DGEMM

Browse Source
This commit is contained in:
Alejandro Gallo 2022-10-08 16:02:49 +02:00
parent 452c0fe001
commit 6871372cac
2 changed files with 66 additions and 38 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/atrip/Atrip.cxx
View File

@ -646,6 +646,9 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
// COMPUTE SINGLES %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% {{{1 // COMPUTE SINGLES %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% {{{1
OCD_Barrier(universe); OCD_Barrier(universe);
#if defined(ATRIP_ONLY_DGEMM)
if (false)
#endif
if (!isFakeTuple(i)) { if (!isFakeTuple(i)) {
WITH_CHRONO("oneshot-unwrap", WITH_CHRONO("oneshot-unwrap",
WITH_CHRONO("unwrap", WITH_CHRONO("unwrap",
@ -678,6 +681,9 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
// COMPUTE ENERGY %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% {{{1 // COMPUTE ENERGY %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% {{{1
#if defined(ATRIP_ONLY_DGEMM)
if (false)
#endif
if (!isFakeTuple(i)) { if (!isFakeTuple(i)) {
double tupleEnergy(0.); double tupleEnergy(0.);

96
src/atrip/Equations.cxx
View File

@ -156,7 +156,6 @@ namespace cuda {
* in order to have an argument in the signature of * in order to have an argument in the signature of
* the function that helps the compiler know which * the function that helps the compiler know which
* instantiation it should take. * instantiation it should take.
*
*/ */
template <typename F, reordering_t R> template <typename F, reordering_t R>
struct reorder_proxy {}; struct reorder_proxy {};
@ -436,22 +435,20 @@ double getEnergySame
, DataFieldType<F>* Tijk_ , DataFieldType<F>* Tijk_
) { ) {
const size_t a = abc[0], b = abc[1], c = abc[2] const size_t NoNo = No*No;
, NoNo = No*No
;
DataFieldType<F>* Tijk = (DataFieldType<F>*)Tijk_; DataFieldType<F>* Tijk = (DataFieldType<F>*)Tijk_;
#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<<<bs, ths>>>(reorder_proxy< \ reorder<<<bs, ths>>>(reorder_proxy< \
DataFieldType<F>, \ DataFieldType<F>, \
__II ## __JJ ## __KK \ __II ## __JJ ## __KK \
>{}, \ >{}, \
No, \ No, \
Tijk, \ Tijk, \
_t_buffer); _t_buffer)
#define DGEMM_PARTICLES(__A, __B) \ #define DGEMM_PARTICLES(__A, __B) \
atrip::xgemm<F>("T", \ atrip::xgemm<F>("T", \
"N", \ "N", \
@ -481,11 +478,18 @@ double getEnergySame
_t_buffer, \ _t_buffer, \
(int const*)&NoNo \ (int const*)&NoNo \
) )
#define MAYBE_CONJ(_conj, _buffer) \ #define MAYBE_CONJ(_conj, _buffer) \
cuda::maybeConjugate<<< \ do { \
Atrip::kernelDimensions.ooo.blocks, \ cuda::maybeConjugate<<< \
Atrip::kernelDimensions.ooo.threads \ \
>>>((DataFieldType<F>*)_conj, (DataFieldType<F>*)_buffer, NoNoNo); Atrip::kernelDimensions.ooo.blocks, \
\
Atrip::kernelDimensions.ooo.threads \
\
>>>((DataFieldType<F>*)_conj, \
(DataFieldType<F>*)_buffer, \
NoNoNo); \
} while (0)
// END CUDA //////////////////////////////////////////////////////////////////// // END CUDA ////////////////////////////////////////////////////////////////////
@ -500,7 +504,9 @@ double getEnergySame
#define REORDER(__II, __JJ, __KK) \ #define REORDER(__II, __JJ, __KK) \
reorder(reorder_proxy<DataFieldType<F>, \ reorder(reorder_proxy<DataFieldType<F>, \
__II ## __JJ ## __KK >{}, \ __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", \
@ -531,9 +537,13 @@ double getEnergySame
_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) \ do { \
_conj[__i] = maybeConjugate<F>(_buffer[__i]); for (size_t __i = 0; __i < NoNoNo; ++__i) { \
_conj[__i] \
= maybeConjugate<F>(_buffer[__i]); \
} \
} while (0)
#endif #endif
F one{1.0}, m_one{-1.0}, zero{0.0}; F one{1.0}, m_one{-1.0}, zero{0.0};
@ -552,8 +562,12 @@ double getEnergySame
const size_t const size_t
bs = Atrip::kernelDimensions.ooo.blocks, bs = Atrip::kernelDimensions.ooo.blocks,
ths = Atrip::kernelDimensions.ooo.threads; ths = Atrip::kernelDimensions.ooo.threads;
#if !defined(ATRIP_ONLY_DGEMM)
cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_t_buffer, NoNoNo); cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_t_buffer, NoNoNo);
cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_vhhh, NoNoNo); cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_vhhh, NoNoNo);
#endif
#else #else
DataFieldType<F>* _t_buffer = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F)); DataFieldType<F>* _t_buffer = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F));
DataFieldType<F>* _vhhh = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F)); DataFieldType<F>* _vhhh = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F));
@ -565,7 +579,7 @@ double getEnergySame
#endif #endif
// Set Tijk to zero // Set Tijk to zero
#ifdef HAVE_CUDA #if defined(HAVE_CUDA) && !defined(ATRIP_ONLY_DGEMM)
WITH_CHRONO("double:reorder", WITH_CHRONO("double:reorder",
cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk, cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk,
NoNoNo); NoNoNo);
@ -577,43 +591,51 @@ double getEnergySame
}) })
#endif #endif
#if defined(ATRIP_ONLY_DGEMM)
#undef MAYBE_CONJ
#undef REORDER
#define MAYBE_CONJ(a, b) do {} while(0)
#define REORDER(i, j, k) do {} while(0)
#endif
// HOLES // HOLES
WITH_CHRONO("doubles:holes", WITH_CHRONO("doubles:holes",
{ {
// VhhhC[i + k*No + L*NoNo] * TABhh[L + j*No]; H1 // VhhhC[i + k*No + L*NoNo] * TABhh[L + j*No]; H1
MAYBE_CONJ(_vhhh, VhhhC) MAYBE_CONJ(_vhhh, VhhhC);
WITH_CHRONO("doubles:holes:1", WITH_CHRONO("doubles:holes:1",
DGEMM_HOLES(_vhhh, TABhh, "N"); DGEMM_HOLES(_vhhh, TABhh, "N");
REORDER(I, K, J) REORDER(I, K, J);
) )
// VhhhC[j + k*No + L*NoNo] * TABhh[i + L*No]; H0 // VhhhC[j + k*No + L*NoNo] * TABhh[i + L*No]; H0
WITH_CHRONO("doubles:holes:2", WITH_CHRONO("doubles:holes:2",
DGEMM_HOLES(_vhhh, TABhh, "T"); DGEMM_HOLES(_vhhh, TABhh, "T");
REORDER(J, K, I) REORDER(J, K, I);
) )
// VhhhB[i + j*No + L*NoNo] * TAChh[L + k*No]; H5 // VhhhB[i + j*No + L*NoNo] * TAChh[L + k*No]; H5
MAYBE_CONJ(_vhhh, VhhhB) MAYBE_CONJ(_vhhh, VhhhB);
WITH_CHRONO("doubles:holes:3", WITH_CHRONO("doubles:holes:3",
DGEMM_HOLES(_vhhh, TAChh, "N"); DGEMM_HOLES(_vhhh, TAChh, "N");
REORDER(I, J, K) REORDER(I, J, K);
) )
// VhhhB[k + j*No + L*NoNo] * TAChh[i + L*No]; H3 // VhhhB[k + j*No + L*NoNo] * TAChh[i + L*No]; H3
WITH_CHRONO("doubles:holes:4", WITH_CHRONO("doubles:holes:4",
DGEMM_HOLES(_vhhh, TAChh, "T"); DGEMM_HOLES(_vhhh, TAChh, "T");
REORDER(K, J, I) REORDER(K, J, I);
) )
// VhhhA[j + i*No + L*NoNo] * TBChh[L + k*No]; H1 // VhhhA[j + i*No + L*NoNo] * TBChh[L + k*No]; H1
MAYBE_CONJ(_vhhh, VhhhA) MAYBE_CONJ(_vhhh, VhhhA);
WITH_CHRONO("doubles:holes:5", WITH_CHRONO("doubles:holes:5",
DGEMM_HOLES(_vhhh, TBChh, "N"); DGEMM_HOLES(_vhhh, TBChh, "N");
REORDER(J, I, K) REORDER(J, I, K);
) )
// VhhhA[k + i*No + L*NoNo] * TBChh[j + L*No]; H4 // VhhhA[k + i*No + L*NoNo] * TBChh[j + L*No]; H4
WITH_CHRONO("doubles:holes:6", WITH_CHRONO("doubles:holes:6",
DGEMM_HOLES(_vhhh, TBChh, "T"); DGEMM_HOLES(_vhhh, TBChh, "T");
REORDER(K, I, J) REORDER(K, I, J);
) )
} }
) )
@ -625,32 +647,32 @@ double getEnergySame
// TAphh[E + i*Nv + j*NoNv] * VBCph[E + k*Nv]; P0 // TAphh[E + i*Nv + j*NoNv] * VBCph[E + k*Nv]; P0
WITH_CHRONO("doubles:particles:1", WITH_CHRONO("doubles:particles:1",
DGEMM_PARTICLES(TAphh, VBCph); DGEMM_PARTICLES(TAphh, VBCph);
REORDER(I, J, K) REORDER(I, J, K);
) )
// TAphh[E + i*Nv + k*NoNv] * VCBph[E + j*Nv]; P3 // TAphh[E + i*Nv + k*NoNv] * VCBph[E + j*Nv]; P3
WITH_CHRONO("doubles:particles:2", WITH_CHRONO("doubles:particles:2",
DGEMM_PARTICLES(TAphh, VCBph); DGEMM_PARTICLES(TAphh, VCBph);
REORDER(I, K, J) REORDER(I, K, J);
) )
// TCphh[E + k*Nv + i*NoNv] * VABph[E + j*Nv]; P5 // TCphh[E + k*Nv + i*NoNv] * VABph[E + j*Nv]; P5
WITH_CHRONO("doubles:particles:3", WITH_CHRONO("doubles:particles:3",
DGEMM_PARTICLES(TCphh, VABph); DGEMM_PARTICLES(TCphh, VABph);
REORDER(K, I, J) REORDER(K, I, J);
) )
// TCphh[E + k*Nv + j*NoNv] * VBAph[E + i*Nv]; P2 // TCphh[E + k*Nv + j*NoNv] * VBAph[E + i*Nv]; P2
WITH_CHRONO("doubles:particles:4", WITH_CHRONO("doubles:particles:4",
DGEMM_PARTICLES(TCphh, VBAph); DGEMM_PARTICLES(TCphh, VBAph);
REORDER(K, J, I) REORDER(K, J, I);
) )
// TBphh[E + j*Nv + i*NoNv] * VACph[E + k*Nv]; P1 // TBphh[E + j*Nv + i*NoNv] * VACph[E + k*Nv]; P1
WITH_CHRONO("doubles:particles:5", WITH_CHRONO("doubles:particles:5",
DGEMM_PARTICLES(TBphh, VACph); DGEMM_PARTICLES(TBphh, VACph);
REORDER(J, I, K) REORDER(J, I, K);
) )
// TBphh[E + j*Nv + k*NoNv] * VCAph[E + i*Nv]; P4 // TBphh[E + j*Nv + k*NoNv] * VCAph[E + i*Nv]; P4
WITH_CHRONO("doubles:particles:6", WITH_CHRONO("doubles:particles:6",
DGEMM_PARTICLES(TBphh, VCAph); DGEMM_PARTICLES(TBphh, VCAph);
REORDER(J, K, I) REORDER(J, K, I);
) )
} }
) )