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

Implement zeroing of tensors through memcpy and cuMemcpy

Browse Source
This commit is contained in:
Alejandro Gallo 2023-01-26 01:49:55 +01:00
parent 4e2d1143e5
commit 93cba3c3ab
4 changed files with 24 additions and 60 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
include/atrip/Operations.hpp
View File

@ -24,15 +24,6 @@ namespace acc {
// cuda kernels // cuda kernels
template <typename F>
__MAYBE_GLOBAL__
void zeroing(F* a, size_t n) {
F zero = {0};
for (size_t i = 0; i < n; i++) {
a[i] = zero;
}
}
//// ////
template <typename F> template <typename F>
__MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__ __MAYBE_DEVICE__ __MAYBE_HOST__ __INLINE__

4
src/atrip/Atrip.cxx
View File

@ -160,9 +160,9 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
LOG(0,"Atrip") << "ooo blocks: " LOG(0,"Atrip") << "ooo blocks: "
<< Atrip::kernelDimensions.ooo.blocks << "\n"; << Atrip::kernelDimensions.ooo.blocks << "\n";
LOG(0,"Atrip") << "ooo threads per block: " LOG(0,"Atrip") << "ooo threads per block: "
<< Atrip::kernelDimensions.ooo.threads << "\n"; << Atrip::kernelDimensions.ooo.threads << "\n";
#endif #endif
// allocate the three scratches, see piecuch // allocate the three scratches, see piecuch

5
src/atrip/Complex.cxx
View File

@ -21,11 +21,6 @@ namespace atrip {
template <> double maybeConjugate(const double a) { return a; } template <> double maybeConjugate(const double a) { return a; }
template <> Complex maybeConjugate(const Complex a) { return std::conj(a); } template <> Complex maybeConjugate(const Complex a) { return std::conj(a); }
#if defined(HAVE_CUDA)
#endif
namespace traits { namespace traits {
template <typename F> bool isComplex() { return false; } template <typename F> bool isComplex() { return false; }
template <> bool isComplex<double>() { return false; } template <> bool isComplex<double>() { return false; }

66
src/atrip/Equations.cxx
View File

@ -13,6 +13,8 @@
// limitations under the License. // limitations under the License.
// [[file:~/cuda/atrip/atrip.org::*Prolog][Prolog:2]] // [[file:~/cuda/atrip/atrip.org::*Prolog][Prolog:2]]
#include <cstring>
#include<atrip/Equations.hpp> #include<atrip/Equations.hpp>
#include<atrip/CUDA.hpp> #include<atrip/CUDA.hpp>
@ -580,13 +582,8 @@ void getEnergySame
) )
#define MAYBE_CONJ(_conj, _buffer) \ #define MAYBE_CONJ(_conj, _buffer) \
do { \ do { \
acc::maybeConjugate<<< \ acc::maybeConjugate<<<1, 1 \
\ >>>((DataFieldType<F>*)_conj, \
Atrip::kernelDimensions.ooo.blocks, \
\
Atrip::kernelDimensions.ooo.threads \
\
>>>((DataFieldType<F>*)_conj, \
(DataFieldType<F>*)_buffer, \ (DataFieldType<F>*)_buffer, \
NoNoNo); \ NoNoNo); \
} while (0) } while (0)
@ -648,50 +645,31 @@ void getEnergySame
F one{1.0}, m_one{-1.0}, zero{0.0}; F one{1.0}, m_one{-1.0}, zero{0.0};
const size_t NoNoNo = No*NoNo; const size_t NoNoNo = No*NoNo;
// Zeroing vectors
#ifdef HAVE_CUDA #ifdef HAVE_CUDA
// DataFieldType<F>* _t_buffer;
// DataFieldType<F>* _vhhh;
// WITH_CHRONO("double:cuda:alloc",
// _CHECK_CUDA_SUCCESS("Allocating _t_buffer",
// cuMemAlloc((CUdeviceptr*)&_t_buffer,
// NoNoNo * sizeof(DataFieldType<F>)));
// _CHECK_CUDA_SUCCESS("Allocating _vhhh",
// cuMemAlloc((CUdeviceptr*)&_vhhh,
// NoNoNo * sizeof(DataFieldType<F>)));
// )
#if !defined(ATRIP_ONLY_DGEMM) #if !defined(ATRIP_ONLY_DGEMM)
// we still have to zero this {
const size_t const size_t elements = NoNoNo * sizeof(DataFieldType<F>)/4;
bs = Atrip::kernelDimensions.ooo.blocks, WITH_CHRONO("double:zeroing",
ths = Atrip::kernelDimensions.ooo.threads; _CHECK_CUDA_SUCCESS("Zeroing Tijk",
acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)_t_buffer, NoNoNo); cuMemsetD32_v2((CUdeviceptr)Tijk, 0x00, elements));
acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)_vhhh, NoNoNo); _CHECK_CUDA_SUCCESS("Zeroing t buffer",
cuMemsetD32_v2((CUdeviceptr)_t_buffer, 0x00, elements));
_CHECK_CUDA_SUCCESS("Zeroing vhhh buffer",
cuMemsetD32_v2((CUdeviceptr)_vhhh, 0x00, elements));
)
}
#endif #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));
DataFieldType<F> zero_h{0.0}; std::memset((void*)_t_buffer, 0x00, NoNoNo * sizeof(DataFieldType<F>));
for (size_t i=0; i < NoNoNo; i++) { std::memset((void*)_vhhh, 0x00, NoNoNo * sizeof(DataFieldType<F>));
_t_buffer[i] = zero_h; std::memset((void*)Tijk, 0x00, NoNoNo * sizeof(DataFieldType<F>));
_vhhh[i] = zero_h; #endif /* HAVE_CUDA */
}
#endif
// Set Tijk to zero
#if defined(HAVE_CUDA) && !defined(ATRIP_ONLY_DGEMM)
WITH_CHRONO("double:reorder",
acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk,
NoNoNo);
)
#endif
#if !defined(HAVE_CUDA)
WITH_CHRONO("double:reorder",
for (size_t k = 0; k < NoNoNo; k++) {
Tijk[k] = DataFieldType<F>{0.0};
})
#endif /* !defined(HAVE_CUDA) */
#if defined(ATRIP_ONLY_DGEMM) #if defined(ATRIP_ONLY_DGEMM)