Compare commits
4 Commits
cuda
...
4e2d1143e5
| Author | SHA1 | Date | |
|---|---|---|---|
| 4e2d1143e5 | |||
|
|
933d556c84 | ||
| c7e3fa45bd | |||
|
|
2b8b3bd421 |
@@ -38,6 +38,8 @@ namespace acc {
|
|||||||
__MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__
|
__MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__
|
||||||
F maybeConjugateScalar(const F &a) { return a; }
|
F maybeConjugateScalar(const F &a) { return a; }
|
||||||
|
|
||||||
|
// TODO: instantiate for std::complex<double>
|
||||||
|
|
||||||
#if defined(HAVE_CUDA)
|
#if defined(HAVE_CUDA)
|
||||||
template <>
|
template <>
|
||||||
__MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__
|
__MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__
|
||||||
|
|||||||
@@ -200,7 +200,7 @@ template <typename F=double>
|
|||||||
: Slice<F>::Fetch
|
: Slice<F>::Fetch
|
||||||
;
|
;
|
||||||
if (blank.info.state == Slice<F>::SelfSufficient) {
|
if (blank.info.state == Slice<F>::SelfSufficient) {
|
||||||
#if defined(HAVE_CUDA)
|
#if defined(HAVE_CUDA) && !defined(ATRIP_SOURCES_IN_GPU)
|
||||||
const size_t _size = sizeof(F) * sliceSize;
|
const size_t _size = sizeof(F) * sliceSize;
|
||||||
// TODO: this is code duplication with downstairs
|
// TODO: this is code duplication with downstairs
|
||||||
if (freePointers.size() == 0) {
|
if (freePointers.size() == 0) {
|
||||||
@@ -221,7 +221,6 @@ template <typename F=double>
|
|||||||
(void*)SOURCES_DATA(sources[from.source]),
|
(void*)SOURCES_DATA(sources[from.source]),
|
||||||
sizeof(F) * sliceSize));
|
sizeof(F) * sliceSize));
|
||||||
))
|
))
|
||||||
|
|
||||||
#else
|
#else
|
||||||
blank.data = SOURCES_DATA(sources[from.source]);
|
blank.data = SOURCES_DATA(sources[from.source]);
|
||||||
#endif
|
#endif
|
||||||
@@ -388,6 +387,22 @@ template <typename F=double>
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
static size_t
|
||||||
|
getSize(const std::vector<size_t> sliceLength,
|
||||||
|
const std::vector<size_t> paramLength,
|
||||||
|
const size_t np,
|
||||||
|
const MPI_Comm global_world) {
|
||||||
|
const RankMap<F> rankMap(paramLength, np, global_world);
|
||||||
|
const size_t
|
||||||
|
nSources = rankMap.nSources(),
|
||||||
|
sliceSize = std::accumulate(sliceLength.begin(),
|
||||||
|
sliceLength.end(),
|
||||||
|
1UL,
|
||||||
|
std::multiplies<size_t>());
|
||||||
|
return nSources * sliceSize;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
// CONSTRUCTOR
|
// CONSTRUCTOR
|
||||||
SliceUnion( std::vector<typename Slice<F>::Type> sliceTypes_
|
SliceUnion( std::vector<typename Slice<F>::Type> sliceTypes_
|
||||||
, std::vector<size_t> sliceLength_
|
, std::vector<size_t> sliceLength_
|
||||||
@@ -574,12 +589,15 @@ template <typename F=double>
|
|||||||
// TODO: do it through the slice class
|
// TODO: do it through the slice class
|
||||||
slice.info.state = Slice<F>::Dispatched;
|
slice.info.state = Slice<F>::Dispatched;
|
||||||
#if defined(HAVE_CUDA) && defined(ATRIP_SOURCES_IN_GPU)
|
#if defined(HAVE_CUDA) && defined(ATRIP_SOURCES_IN_GPU)
|
||||||
# if !defined(ATRIP_CUDA_AWARE_MPI)
|
# if !defined(ATRIP_CUDA_AWARE_MPI)
|
||||||
# error "You need CUDA aware MPI to have slices on the GPU"
|
# error "You need CUDA aware MPI to have slices on the GPU"
|
||||||
# endif
|
# endif
|
||||||
MPI_Irecv((void*)slice.data,
|
MPI_Irecv((void*)slice.data,
|
||||||
|
#elif defined(HAVE_CUDA) && !defined(ATRIP_SOURCES_IN_GPU)
|
||||||
|
slice.mpi_data = (F*)malloc(sizeof(F) * slice.size);
|
||||||
|
MPI_Irecv(slice.mpi_data,
|
||||||
#else
|
#else
|
||||||
MPI_Irecv(slice.data,
|
MPI_Irecv((void*)slice.data,
|
||||||
#endif
|
#endif
|
||||||
slice.size,
|
slice.size,
|
||||||
traits::mpi::datatypeOf<F>(),
|
traits::mpi::datatypeOf<F>(),
|
||||||
|
|||||||
@@ -235,11 +235,54 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
|
|||||||
MPI_Comm_size(child_comm, &child_size);
|
MPI_Comm_size(child_comm, &child_size);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// a, b, c, d, e, f and P => Nv
|
||||||
|
// H => No
|
||||||
|
// total_source_sizes contains a list of the number of elements
|
||||||
|
// in all sources of every tensor union, therefore nSlices * sliceSize
|
||||||
|
const std::vector<size_t> total_source_sizes = {
|
||||||
|
// ABPH
|
||||||
|
SliceUnion<F>::getSize({Nv, No}, {Nv, Nv}, (size_t)np, universe),
|
||||||
|
// ABHH
|
||||||
|
SliceUnion<F>::getSize({No, No}, {Nv, Nv}, (size_t)np, universe),
|
||||||
|
// TABHH
|
||||||
|
SliceUnion<F>::getSize({No, No}, {Nv, Nv}, (size_t)np, universe),
|
||||||
|
// TAPHH
|
||||||
|
SliceUnion<F>::getSize({Nv, No, No}, {Nv}, (size_t)np, universe),
|
||||||
|
// HHHA
|
||||||
|
SliceUnion<F>::getSize({No, No, No}, {Nv}, (size_t)np, universe),
|
||||||
|
};
|
||||||
|
|
||||||
|
const size_t
|
||||||
|
total_source_size = sizeof(DataFieldType<F>)
|
||||||
|
* std::accumulate(total_source_sizes.begin(),
|
||||||
|
total_source_sizes.end(),
|
||||||
|
0UL);
|
||||||
|
|
||||||
|
#if defined(HAVE_CUDA)
|
||||||
|
DataPtr<F> all_sources_pointer;
|
||||||
|
cuMemAlloc(&all_sources_pointer, total_source_size);
|
||||||
|
#else
|
||||||
|
DataPtr<F>
|
||||||
|
all_sources_pointer = (DataPtr<F>)malloc(total_source_size);
|
||||||
|
#endif
|
||||||
|
size_t _source_pointer_idx = 0;
|
||||||
|
|
||||||
// BUILD SLICES PARAMETRIZED BY NV x NV =============================={{{1
|
// BUILD SLICES PARAMETRIZED BY NV x NV =============================={{{1
|
||||||
WITH_CHRONO("nv-nv-slices",
|
WITH_CHRONO("nv-nv-slices",
|
||||||
LOG(0,"Atrip") << "building NV x NV slices\n";
|
LOG(0,"Atrip") << "building NV x NV slices\n";
|
||||||
|
// TODO
|
||||||
|
// DataPtr<F> offseted_pointer = all_sources_pointer
|
||||||
|
// * total_source_sizes[_source_pointer_idx++];
|
||||||
ABPH<F> abph(*in.Vppph, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
ABPH<F> abph(*in.Vppph, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
||||||
|
|
||||||
|
// TODO
|
||||||
|
// DataPtr<F> offseted_pointer = all_sources_pointer
|
||||||
|
// * total_source_sizes[_source_pointer_idx++];
|
||||||
ABHH<F> abhh(*in.Vpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
ABHH<F> abhh(*in.Vpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
||||||
|
|
||||||
|
// TODO
|
||||||
|
// DataPtr<F> offseted_pointer = all_sources_pointer
|
||||||
|
// * total_source_sizes[_source_pointer_idx++];
|
||||||
TABHH<F> tabhh(*in.Tpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
TABHH<F> tabhh(*in.Tpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
||||||
)
|
)
|
||||||
|
|
||||||
@@ -251,7 +294,13 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
|
|||||||
// BUILD SLICES PARAMETRIZED BY NV ==================================={{{1
|
// BUILD SLICES PARAMETRIZED BY NV ==================================={{{1
|
||||||
WITH_CHRONO("nv-slices",
|
WITH_CHRONO("nv-slices",
|
||||||
LOG(0,"Atrip") << "building NV slices\n";
|
LOG(0,"Atrip") << "building NV slices\n";
|
||||||
|
// TODO
|
||||||
|
// DataPtr<F> offseted_pointer = all_sources_pointer
|
||||||
|
// * total_source_sizes[_source_pointer_idx++];
|
||||||
TAPHH<F> taphh(*in.Tpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
TAPHH<F> taphh(*in.Tpphh, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
||||||
|
// TODO
|
||||||
|
// DataPtr<F> offseted_pointer = all_sources_pointer
|
||||||
|
// * total_source_sizes[_source_pointer_idx++];
|
||||||
HHHA<F> hhha(*in.Vhhhp, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
HHHA<F> hhha(*in.Vhhhp, (size_t)No, (size_t)Nv, (size_t)np, child_comm, universe);
|
||||||
)
|
)
|
||||||
|
|
||||||
@@ -903,5 +952,5 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
|
|||||||
}
|
}
|
||||||
// instantiate
|
// instantiate
|
||||||
template Atrip::Output Atrip::run(Atrip::Input<double> const& in);
|
template Atrip::Output Atrip::run(Atrip::Input<double> const& in);
|
||||||
template Atrip::Output Atrip::run(Atrip::Input<Complex> const& in);
|
// template Atrip::Output Atrip::run(Atrip::Input<Complex> const& in);
|
||||||
// Main:1 ends here
|
// Main:1 ends here
|
||||||
|
|||||||
@@ -25,11 +25,8 @@ namespace atrip {
|
|||||||
|
|
||||||
#if defined(HAVE_CUDA)
|
#if defined(HAVE_CUDA)
|
||||||
#define FOR_K() \
|
#define FOR_K() \
|
||||||
for (size_t kmin = blockIdx.x * blockDim.x + threadIdx.x, \
|
const size_t k = blockIdx.x * blockDim.x + threadIdx.x; \
|
||||||
k = kmin, \
|
size_t idx = 0;
|
||||||
idx = kmin * size * size * size; \
|
|
||||||
k < (kmin < size) ? kmin + 1 : size; \
|
|
||||||
k++)
|
|
||||||
#else
|
#else
|
||||||
#define FOR_K() for (size_t k=0, idx=0; k < size; k++)
|
#define FOR_K() for (size_t k=0, idx=0; k < size; k++)
|
||||||
#endif
|
#endif
|
||||||
@@ -102,6 +99,7 @@ namespace atrip {
|
|||||||
# define MIN(a, b) std::min((a), (b))
|
# define MIN(a, b) std::min((a), (b))
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
#if defined(ATRIP_NEW_ENERGY)
|
||||||
|
|
||||||
// [[file:~/cuda/atrip/atrip.org::*Energy][Energy:2]]
|
// [[file:~/cuda/atrip/atrip.org::*Energy][Energy:2]]
|
||||||
template <typename F>
|
template <typename F>
|
||||||
@@ -250,6 +248,131 @@ void getEnergySame
|
|||||||
}
|
}
|
||||||
// Energy:2 ends here
|
// Energy:2 ends here
|
||||||
|
|
||||||
|
#else
|
||||||
|
|
||||||
|
// [[file:~/cuda/atrip/atrip.org::*Energy][Energy:2]]
|
||||||
|
template <typename F>
|
||||||
|
__MAYBE_GLOBAL__
|
||||||
|
void getEnergyDistinct
|
||||||
|
( F const epsabc
|
||||||
|
, size_t const No
|
||||||
|
, F* const epsi
|
||||||
|
, F* const Tijk
|
||||||
|
, F* const Zijk
|
||||||
|
, double* _energy
|
||||||
|
) {
|
||||||
|
constexpr size_t blockSize=16;
|
||||||
|
F energy(0.);
|
||||||
|
for (size_t kk=0; kk<No; kk+=blockSize){
|
||||||
|
const size_t kend( MIN(No, kk+blockSize) );
|
||||||
|
for (size_t jj(kk); jj<No; jj+=blockSize){
|
||||||
|
const size_t jend( MIN( No, jj+blockSize) );
|
||||||
|
for (size_t ii(jj); ii<No; ii+=blockSize){
|
||||||
|
const size_t iend( MIN( No, ii+blockSize) );
|
||||||
|
for (size_t k(kk); k < kend; k++){
|
||||||
|
const F ek(epsi[k]);
|
||||||
|
const size_t jstart = jj > k ? jj : k;
|
||||||
|
for (size_t j(jstart); j < jend; j++){
|
||||||
|
F const ej(epsi[j]);
|
||||||
|
F const facjk = j == k ? F(0.5) : F(1.0);
|
||||||
|
size_t istart = ii > j ? ii : j;
|
||||||
|
for (size_t i(istart); i < iend; i++){
|
||||||
|
const F
|
||||||
|
ei(epsi[i])
|
||||||
|
, facij = i == j ? F(0.5) : F(1.0)
|
||||||
|
, denominator(epsabc - ei - ej - ek)
|
||||||
|
, U(Zijk[i + No*j + No*No*k])
|
||||||
|
, V(Zijk[i + No*k + No*No*j])
|
||||||
|
, W(Zijk[j + No*i + No*No*k])
|
||||||
|
, X(Zijk[j + No*k + No*No*i])
|
||||||
|
, Y(Zijk[k + No*i + No*No*j])
|
||||||
|
, Z(Zijk[k + No*j + No*No*i])
|
||||||
|
, A(acc::maybeConjugateScalar<F>(Tijk[i + No*j + No*No*k]))
|
||||||
|
, B(acc::maybeConjugateScalar<F>(Tijk[i + No*k + No*No*j]))
|
||||||
|
, C(acc::maybeConjugateScalar<F>(Tijk[j + No*i + No*No*k]))
|
||||||
|
, D(acc::maybeConjugateScalar<F>(Tijk[j + No*k + No*No*i]))
|
||||||
|
, E(acc::maybeConjugateScalar<F>(Tijk[k + No*i + No*No*j]))
|
||||||
|
, _F(acc::maybeConjugateScalar<F>(Tijk[k + No*j + No*No*i]))
|
||||||
|
, value
|
||||||
|
= 3.0 * ( A * U
|
||||||
|
+ B * V
|
||||||
|
+ C * W
|
||||||
|
+ D * X
|
||||||
|
+ E * Y
|
||||||
|
+ _F * Z )
|
||||||
|
+ ( ( U + X + Y )
|
||||||
|
- 2.0 * ( V + W + Z )
|
||||||
|
) * ( A + D + E )
|
||||||
|
+ ( ( V + W + Z )
|
||||||
|
- 2.0 * ( U + X + Y )
|
||||||
|
) * ( B + C + _F )
|
||||||
|
;
|
||||||
|
energy += 2.0 * value / denominator * facjk * facij;
|
||||||
|
} // i
|
||||||
|
} // j
|
||||||
|
} // k
|
||||||
|
} // ii
|
||||||
|
} // jj
|
||||||
|
} // kk
|
||||||
|
*_energy = acc::real(energy);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
template <typename F>
|
||||||
|
__MAYBE_GLOBAL__
|
||||||
|
void getEnergySame
|
||||||
|
( F const epsabc
|
||||||
|
, size_t const No
|
||||||
|
, F* const epsi
|
||||||
|
, F* const Tijk
|
||||||
|
, F* const Zijk
|
||||||
|
, double* _energy
|
||||||
|
) {
|
||||||
|
constexpr size_t blockSize = 16;
|
||||||
|
F energy = F(0.);
|
||||||
|
for (size_t kk=0; kk<No; kk+=blockSize){
|
||||||
|
const size_t kend( MIN( kk+blockSize, No) );
|
||||||
|
for (size_t jj(kk); jj<No; jj+=blockSize){
|
||||||
|
const size_t jend( MIN( jj+blockSize, No) );
|
||||||
|
for (size_t ii(jj); ii<No; ii+=blockSize){
|
||||||
|
const size_t iend( MIN( ii+blockSize, No) );
|
||||||
|
for (size_t k(kk); k < kend; k++){
|
||||||
|
const F ek(epsi[k]);
|
||||||
|
const size_t jstart = jj > k ? jj : k;
|
||||||
|
for(size_t j(jstart); j < jend; j++){
|
||||||
|
const F facjk( j == k ? F(0.5) : F(1.0));
|
||||||
|
const F ej(epsi[j]);
|
||||||
|
const size_t istart = ii > j ? ii : j;
|
||||||
|
for(size_t i(istart); i < iend; i++){
|
||||||
|
const F
|
||||||
|
ei(epsi[i])
|
||||||
|
, facij ( i==j ? F(0.5) : F(1.0))
|
||||||
|
, denominator(epsabc - ei - ej - ek)
|
||||||
|
, U(Zijk[i + No*j + No*No*k])
|
||||||
|
, V(Zijk[j + No*k + No*No*i])
|
||||||
|
, W(Zijk[k + No*i + No*No*j])
|
||||||
|
, A(acc::maybeConjugateScalar<F>(Tijk[i + No*j + No*No*k]))
|
||||||
|
, B(acc::maybeConjugateScalar<F>(Tijk[j + No*k + No*No*i]))
|
||||||
|
, C(acc::maybeConjugateScalar<F>(Tijk[k + No*i + No*No*j]))
|
||||||
|
, value
|
||||||
|
= F(3.0) * ( A * U
|
||||||
|
+ B * V
|
||||||
|
+ C * W
|
||||||
|
)
|
||||||
|
- ( A + B + C ) * ( U + V + W )
|
||||||
|
;
|
||||||
|
energy += F(2.0) * value / denominator * facjk * facij;
|
||||||
|
} // i
|
||||||
|
} // j
|
||||||
|
} // k
|
||||||
|
} // ii
|
||||||
|
} // jj
|
||||||
|
} // kk
|
||||||
|
*_energy = acc::real(energy);
|
||||||
|
}
|
||||||
|
// Energy:2 ends here
|
||||||
|
#endif /* defined(ATRIP_NEW_ENERGY) */
|
||||||
|
|
||||||
// [[file:~/cuda/atrip/atrip.org::*Energy][Energy:3]]
|
// [[file:~/cuda/atrip/atrip.org::*Energy][Energy:3]]
|
||||||
// instantiate double
|
// instantiate double
|
||||||
template
|
template
|
||||||
@@ -274,6 +397,8 @@ void getEnergySame
|
|||||||
, DataFieldType<double>* energy
|
, DataFieldType<double>* energy
|
||||||
);
|
);
|
||||||
|
|
||||||
|
// TODO: put this back in
|
||||||
|
#if defined(ATRIP_WITH_COMPLEX)
|
||||||
// instantiate Complex
|
// instantiate Complex
|
||||||
template
|
template
|
||||||
__MAYBE_GLOBAL__
|
__MAYBE_GLOBAL__
|
||||||
@@ -297,6 +422,7 @@ void getEnergySame
|
|||||||
, DataFieldType<double>* energy
|
, DataFieldType<double>* energy
|
||||||
);
|
);
|
||||||
// Energy:3 ends here
|
// Energy:3 ends here
|
||||||
|
#endif
|
||||||
|
|
||||||
// [[file:~/cuda/atrip/atrip.org::*Singles%20contribution][Singles contribution:2]]
|
// [[file:~/cuda/atrip/atrip.org::*Singles%20contribution][Singles contribution:2]]
|
||||||
template <typename F> __MAYBE_GLOBAL__
|
template <typename F> __MAYBE_GLOBAL__
|
||||||
@@ -416,7 +542,7 @@ void 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<<<bs, ths>>>(reorder_proxy< \
|
reorder<<<1, No>>>(reorder_proxy< \
|
||||||
DataFieldType<F>, \
|
DataFieldType<F>, \
|
||||||
__II ## __JJ ## __KK \
|
__II ## __JJ ## __KK \
|
||||||
>{}, \
|
>{}, \
|
||||||
|
|||||||
Reference in New Issue
Block a user