Fix zeroing

include/atrip/SliceUnion.hpp

@@ -465,7 +465,7 @@ template <typename F=double>
 
 #if defined(HAVE_CUDA)
       LOG(1,"Atrip") << "warming communication up " << slices.size() << "\n";
-      WITH_CHRONO("cuda:warmup",
+      WITH_CHRONO("cuda:warmup",
                   int nRanks=Atrip::np, requestCount=0;
                   int nSends=sliceBuffers.size()*nRanks;
                   MPI_Request *requests = (MPI_Request*) malloc(nSends*2 * sizeof(MPI_Request));

src/atrip/Atrip.cxx

@@ -202,7 +202,7 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
   _CHECK_CUDA_SUCCESS("Zijk",
                       cuMemAlloc(&Zijk, sizeof(F) * No * No * No));
 #else
-  std::vector<F> &Tai = _Tai, &epsi = _epsi, &epsa = _epsa;
+  DataPtr<F> Tai = _Tai.data(), epsi = _epsi.data(), epsa = _epsa.data();
   Zijk = (DataFieldType<F>*)malloc(No*No*No * sizeof(DataFieldType<F>));
   Tijk = (DataFieldType<F>*)malloc(No*No*No * sizeof(DataFieldType<F>));
 #endif
@@ -693,7 +693,7 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
                                       (DataFieldType<F>*)Tai,
 #else
       singlesContribution<F>(No, Nv, abc[0], abc[1], abc[2],
-                             Tai.data(),
+                             Tai,
 #endif
                              (DataFieldType<F>*)abhh.unwrapSlice(Slice<F>::AB,
                                                                  abc),
@@ -733,18 +733,30 @@ Atrip::Output Atrip::run(Atrip::Input<F> const& in) {
                                 1, 1, // for cuda
                                 _epsabc,
                                 No,
+#if defined(HAVE_CUDA)
                                 (DataFieldType<F>*)epsi,
                                 (DataFieldType<F>*)Tijk,
                                 (DataFieldType<F>*)Zijk,
+#else
+                                epsi,
+                                Tijk,
+                                Zijk,
+#endif
                                 tupleEnergy);
                   } else {
                     ACC_FUNCALL(getEnergySame<DataFieldType<F>>,
                                 1, 1, // for cuda
                                 _epsabc,
                                 No,
+#if defined(HAVE_CUDA)
                                 (DataFieldType<F>*)epsi,
                                 (DataFieldType<F>*)Tijk,
                                 (DataFieldType<F>*)Zijk,
+#else
+                                epsi,
+                                Tijk,
+                                Zijk,
+#endif
                                 tupleEnergy);
                   })
 

src/atrip/Equations.cxx

@@ -533,11 +533,15 @@ void getEnergySame
 //                        cuMemAlloc((CUdeviceptr*)&_vhhh,
 //                                   NoNoNo * sizeof(DataFieldType<F>)));
 //                )
-//    const size_t
+#if !defined(ATRIP_ONLY_DGEMM)
-//      bs = Atrip::kernelDimensions.ooo.blocks,
+    // we still have to zero this
-//      ths = Atrip::kernelDimensions.ooo.threads;
+    const size_t
-    //cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_t_buffer, NoNoNo);
+      bs = Atrip::kernelDimensions.ooo.blocks,
-    //cuda::zeroing<<<bs, ths>>>((DataFieldType<F>*)_vhhh, NoNoNo);
+      ths = Atrip::kernelDimensions.ooo.threads;
+    acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)_t_buffer, NoNoNo);
+    acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)_vhhh, NoNoNo);
+#endif
+
 #else
     DataFieldType<F>* _t_buffer = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F));
     DataFieldType<F>* _vhhh = (DataFieldType<F>*)malloc(NoNoNo * sizeof(F));
@@ -554,12 +558,14 @@ void getEnergySame
                 acc::zeroing<<<bs, ths>>>((DataFieldType<F>*)Tijk,
                                            NoNoNo);
                 )
-#else
+#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) && !defined(ATRIP_ONLY_DGEMM) */
+#endif /* !defined(HAVE_CUDA) */
 
 
 #if defined(ATRIP_ONLY_DGEMM)