Improve memory usage when calculating normalization integrals

.gitignore

@@ -30,4 +30,17 @@ Tutorials/Dalitz/doc/*.aux
 Tutorials/Dalitz/doc/*.log
 Tutorials/Dalitz/doc/*.toc
 Tutorials/Dalitz/doc/*.synctex.gz
-
+UnitTests/AmpToolsInterfaceTest
+UnitTests/AmpToolsInterfaceTest_GPU
+UnitTests/AmpToolsInterfaceTestMPI
+UnitTests/AmpToolsInterfaceTestMPI_GPU
+UnitTests/configurationInfoTest
+UnitTests/configurationInfoTest_GPU
+UnitTests/fitResultsTest
+UnitTests/fitResultsTest_GPU
+UnitTests/fitResultsTestMPI
+UnitTests/fitResultsTestMPI_GPU
+UnitTests/parserTest
+UnitTests/parserTest_GPU
+UnitTests/writeModels
+UnitTests/writeModels_GPU

AmpTools/GPUManager/CUDA-Complex.cuh

@@ -37,6 +37,8 @@
 #ifndef __CUDA_COMPLEX__H__
 #define __CUDA_COMPLEX__H__
 
+#include <math.h>
+
 #include "GPUManager/GPUCustomTypes.h"
 
 #ifdef __CUDACC__

AmpTools/GPUManager/GPUAmpProductKernel.cu

@@ -38,7 +38,7 @@
 
 __global__ void
 amp_kernel( GDouble* pfDevAmps, GDouble* pfDevVVStar, GDouble* pfDevWeights, 
-            int nAmps, int nEvents, double* pdDevRes )
+            int nAmps, unsigned int nEvents, double* pdDevRes )
 {
 	int i = threadIdx.x + GPU_BLOCK_SIZE_X * threadIdx.y + 
             ( blockIdx.x + blockIdx.y * gridDim.x ) * GPU_BLOCK_SIZE_SQ;
@@ -100,7 +100,7 @@ amp_kernel( GDouble* pfDevAmps, GDouble* pfDevVVStar, GDouble* pfDevWeights,
 
 extern "C" void GPU_ExecAmpKernel( dim3 dimGrid, dim3 dimBlock, 
      GDouble* pfDevAmps, GDouble* pfDevVVStar, GDouble* pfDevWeights,
-     int nAmps, int nEvents, double* pdDevRes )
+     int nAmps, unsigned int nEvents, double* pdDevRes )
 {
 	amp_kernel<<< dimGrid, dimBlock >>>( pfDevAmps, pfDevVVStar, pfDevWeights, 
                                        nAmps, nEvents, pdDevRes );

AmpTools/GPUManager/GPUCustomTypes.h

@@ -50,12 +50,15 @@ using namespace std;
 
 // a few helpers to make operations on GPU a little more user friendly
 
-// standard arguments to kernel launches
+// standard arguments to ampliutde kernel launches
 #define GPU_AMP_PROTO GDouble* pfDevData, GDouble* pfDevUserVars, \
-                      WCUComplex* pcDevAmp, int* piDevPerm, int iNParticles, int iNEvents
-#define GPU_AMP_ARGS  pfDevData, pfDevUserVars, pcDevAmp, piDevPerm, iNParticles, iNEvents
+                      WCUComplex* pcDevAmp, int* piDevPerm, unsigned int iNParticles, \
+                      unsigned int iNEvents, unsigned int startEvent
+#define GPU_AMP_ARGS  pfDevData, pfDevUserVars, pcDevAmp, piDevPerm, iNParticles, \
+                      iNEvents, startEvent                   
 
-// how to get the event we are working with:
+// how to get the event we are working with -- this is within a chunk and is typically
+// set to iEvent by the user
 #define GPU_THIS_EVENT threadIdx.x + GPU_BLOCK_SIZE_X * threadIdx.y + \
                       ( blockIdx.x + blockIdx.y * gridDim.x ) * GPU_BLOCK_SIZE_SQ
 
@@ -65,11 +68,13 @@ using namespace std;
 #define GPU_PY 2
 #define GPU_PZ 3
 
-#define GPU_KIN(id,val) pfDevData[ ( piDevPerm[id] * 4 + val ) * iNEvents + iEvent]
+// both the data and user vars are always stored for the entire data set so the indices
+// need to be adjusted by startEvent when doing a chunked amplitude calculation
+#define GPU_KIN(id,val) pfDevData[ ( piDevPerm[id] * 4 + val ) * iNEvents + iEvent + startEvent ]
 #define GPU_P4(id) { GPU_KIN(id,GPU_E)  , GPU_KIN(id,GPU_PX), \
                      GPU_KIN(id,GPU_PY) , GPU_KIN(id,GPU_PZ)}
 
-#define GPU_UVARS(val) pfDevUserVars[val*iNEvents+iEvent]
+#define GPU_UVARS(val) pfDevUserVars[val*iNEvents + iEvent + startEvent]
 
 #define COPY_P4(a1,a2) GDouble a2[4]; for( int zzz = 0; zzz < 4; ++zzz ) a2[zzz] = a1[zzz];
 

AmpTools/GPUManager/GPUFactPermKernel.cu

@@ -37,29 +37,29 @@
 #include "GPUCustomTypes.h"
 
 __global__ void
-fact_perm_kernel( GDouble* pfDevAmps, GDouble* pcDevCalcAmp, int nFact,
-                  int nPerm, int nEvents )
+fact_perm_kernel( GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact,
+                  int nPerm, unsigned int nEvents )
 {
 
 	int i = threadIdx.x + GPU_BLOCK_SIZE_X * threadIdx.y + 
             ( blockIdx.x + blockIdx.y * gridDim.x ) * GPU_BLOCK_SIZE_SQ;
 
   for( int iPerm = 0; iPerm < nPerm; ++iPerm ){
 
-    int offsetP = 2*nEvents*iPerm + 2*i;
+    unsigned int offsetP = 2*nEvents*iPerm + 2*i;
 
-    double ampRe = pcDevCalcAmp[offsetP];
-    double ampIm = pcDevCalcAmp[offsetP+1];
+    GDouble ampRe = pcDevAmpFact[offsetP];
+    GDouble ampIm = pcDevAmpFact[offsetP+1];
 
     for( int iFactor = 1; iFactor < nFact; ++iFactor ){
 
-       int offsetF = 2*nEvents*nPerm*iFactor + offsetP;
+       unsigned int offsetF = 2*nEvents*nPerm*iFactor + offsetP;
 
-       double re = ampRe;
-       double im = ampIm;
+       GDouble re = ampRe;
+       GDouble im = ampIm;
 
-       ampRe = re * pcDevCalcAmp[offsetF] - im * pcDevCalcAmp[offsetF+1];
-       ampIm = re * pcDevCalcAmp[offsetF+1] + im * pcDevCalcAmp[offsetF];
+       ampRe = re * pcDevAmpFact[offsetF] - im * pcDevAmpFact[offsetF+1];
+       ampIm = re * pcDevAmpFact[offsetF+1] + im * pcDevAmpFact[offsetF];
     }
 
     pfDevAmps[2*i]   += ampRe;
@@ -71,8 +71,8 @@ fact_perm_kernel( GDouble* pfDevAmps, GDouble* pcDevCalcAmp, int nFact,
 }
 
 extern "C" void GPU_ExecFactPermKernel( dim3 dimGrid, dim3 dimBlock,
-    GDouble* pfDevAmps, GDouble* pcDevCalcAmp, int nFact, int nPerm, int nEvents )
+    GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact, int nPerm, unsigned int nEvents )
 {
-  fact_perm_kernel<<< dimGrid, dimBlock >>>( pfDevAmps, pcDevCalcAmp, 
+  fact_perm_kernel<<< dimGrid, dimBlock >>>( pfDevAmps, pcDevAmpFact, 
                                              nFact, nPerm, nEvents );
 }

AmpTools/GPUManager/GPUKernel.h

@@ -41,16 +41,16 @@
 
 extern "C" void GPU_ExecAmpKernel(dim3 dimGrid,dim3 dimBlock,GDouble* pfDevAmps,
                                   GDouble* pfDevVVStar, GDouble* pfDevWeights,
-                                  int nAmps, int nEvents, double* pfDevRes);
+                                  int nAmps, unsigned int nEvents, double* pfDevRes);
 
 extern "C" void GPU_ExecFactPermKernel( dim3 dimGrid, dim3 dimBlock,
-                                        GDouble* pfDevAmps, GDouble* pcDevCalcAmp,
-                                        int nFact, int nPerm, int nEvents );
-
+                                        GDouble* pfDevAmps, GDouble* pcDevAmpFact,
+                                        int nFact, int nPerm, unsigned int nEvents );
 
 extern "C" void GPU_ExecNICalcKernel( dim3 dimGrid, dim3 dimBlock, unsigned int sharedSize,
                                       int nElements, double* pdDevNICalc,
                                       GDouble* pfDevAmps, GDouble* pfDevWeights,
-                                      int nEvents, int nTrueEvents, unsigned int maxSize );
+                                      unsigned int startEvent, unsigned int nEvents, 
+                                      unsigned int nTrueEvents, unsigned int maxSize );
 
 #endif //__GPU_KERNEL__H__