Skip to content

Chunk memory debug #141

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
mashephe merged 11 commits into from
Aug 15, 2026
Merged

Chunk memory debug #141

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
dbaa1a0
change many unsigned int types used in array indexing to size_t
mashephe Aug 13, 2026
910bfba
change many unsigned int types used in array indexing to size_t
mashephe Aug 13, 2026
c6a30e0
reduce CPU allocation; remove unnecessary factor of 2
mashephe Aug 13, 2026
70f6dac
change many unsigned int types used in array indexing to size_t
mashephe Aug 13, 2026
b926a4f
rework how chunk size for generated MC is handled to allow chunk size…
mashephe Aug 14, 2026
9890da1
fix bugs in the chunked calculation of generated MC that caused some …
mashephe Aug 14, 2026
b974751
introduce event loader helper function and rework loadData and loadEv…
mashephe Aug 14, 2026
7c52c7d
tighten up memory allocations in places where they were unnecessarily…
mashephe Aug 14, 2026
de4552d
fix logic bug that caused the last chunk to get skipped
Aug 15, 2026
3643e1f
remove unnecessary optimization that is already built into logic to s…
Aug 15, 2026
6657278
remove unneeded and improperly written constructor
mashephe Aug 15, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions AmpTools/GPUManager/GPUAmpProductKernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,7 @@

__global__ void
amp_kernel( GDouble* pfDevAmps, GDouble* pfDevVVStar, GDouble* pfDevWeights,
int nAmps, unsigned int nEvents, double* pdDevRes )
int nAmps, size_t nEvents, double* pdDevRes )
{
int i = threadIdx.x + GPU_BLOCK_SIZE_X * threadIdx.y +
( blockIdx.x + blockIdx.y * gridDim.x ) * GPU_BLOCK_SIZE_SQ;
Expand Down Expand Up @@ -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, unsigned int nEvents, double* pdDevRes )
int nAmps, size_t nEvents, double* pdDevRes )
{
amp_kernel<<< dimGrid, dimBlock >>>( pfDevAmps, pfDevVVStar, pfDevWeights,
nAmps, nEvents, pdDevRes );
Expand Down
4 changes: 2 additions & 2 deletions AmpTools/GPUManager/GPUCustomTypes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,8 +52,8 @@ using namespace std;

// standard arguments to ampliutde kernel launches
#define GPU_AMP_PROTO GDouble* pfDevData, GDouble* pfDevUserVars, \
WCUComplex* pcDevAmp, int* piDevPerm, unsigned int iNParticles, \
unsigned int iNEvents, unsigned int startEvent
WCUComplex* pcDevAmp, int* piDevPerm, size_t iNParticles, \
size_t iNEvents, size_t startEvent
#define GPU_AMP_ARGS pfDevData, pfDevUserVars, pcDevAmp, piDevPerm, iNParticles, \
iNEvents, startEvent

Expand Down
8 changes: 4 additions & 4 deletions AmpTools/GPUManager/GPUFactPermKernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,22 +38,22 @@

__global__ void
fact_perm_kernel( GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact,
int nPerm, unsigned int nEvents )
int nPerm, size_t 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 ){

unsigned int offsetP = 2*nEvents*iPerm + 2*i;
size_t offsetP = 2*nEvents*iPerm + 2*i;

GDouble ampRe = pcDevAmpFact[offsetP];
GDouble ampIm = pcDevAmpFact[offsetP+1];

for( int iFactor = 1; iFactor < nFact; ++iFactor ){

unsigned int offsetF = 2*nEvents*nPerm*iFactor + offsetP;
size_t offsetF = 2*nEvents*nPerm*iFactor + offsetP;

GDouble re = ampRe;
GDouble im = ampIm;
Expand All @@ -71,7 +71,7 @@ fact_perm_kernel( GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact,
}

extern "C" void GPU_ExecFactPermKernel( dim3 dimGrid, dim3 dimBlock,
GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact, int nPerm, unsigned int nEvents )
GDouble* pfDevAmps, GDouble* pcDevAmpFact, int nFact, int nPerm, size_t nEvents )
{
fact_perm_kernel<<< dimGrid, dimBlock >>>( pfDevAmps, pcDevAmpFact,
nFact, nPerm, nEvents );
Expand Down
10 changes: 5 additions & 5 deletions AmpTools/GPUManager/GPUKernel.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,16 +41,16 @@

extern "C" void GPU_ExecAmpKernel(dim3 dimGrid,dim3 dimBlock,GDouble* pfDevAmps,
GDouble* pfDevVVStar, GDouble* pfDevWeights,
int nAmps, unsigned int nEvents, double* pfDevRes);
int nAmps, size_t nEvents, double* pfDevRes);

extern "C" void GPU_ExecFactPermKernel( dim3 dimGrid, dim3 dimBlock,
GDouble* pfDevAmps, GDouble* pcDevAmpFact,
int nFact, int nPerm, unsigned int nEvents );
int nFact, int nPerm, size_t nEvents );

extern "C" void GPU_ExecNICalcKernel( dim3 dimGrid, dim3 dimBlock, unsigned int sharedSize,
extern "C" void GPU_ExecNICalcKernel( dim3 dimGrid, dim3 dimBlock, size_t sharedSize,
int nElements, double* pdDevNICalc,
GDouble* pfDevAmps, GDouble* pfDevWeights,
unsigned int startEvent, unsigned int nEvents,
unsigned int nTrueEvents, unsigned int maxSize );
size_t startEvent, size_t nEvents,
size_t nTrueEvents, size_t maxSize );

#endif //__GPU_KERNEL__H__
Loading
Loading