Skip to content
Merged
Show file tree
Hide file tree
Changes from 105 commits
Commits
Show all changes
116 commits
Select commit Hold shift + click to select a range
e2393c7
clearRegisteredPieces called before tree build
spencerw Sep 25, 2024
7dec4c5
CUDA streams assigned to TreePieces after tree build
spencerw Sep 25, 2024
ba603d6
Move clearRegisteredPieces inside buildTree
spencerw Sep 25, 2024
8bbeccb
Remove unused code
spencerw Sep 25, 2024
049e4fa
Move assignCUDAStreams() call inside of buildTree()
spencerw Sep 30, 2024
8ddb491
Remove DataTransfer calls from TreePieceCellListDataTransferLocal whe…
spencerw Oct 22, 2024
a5e500b
Hack to launch gpuLocalTreeWalk from DataManager
spencerw Oct 22, 2024
83961cf
Hack to run remote gravity without bookkeeping
spencerw Oct 23, 2024
f24d350
Restore bookkeeping around GPU local tree walk
spencerw Nov 12, 2024
bb0d5bb
Use bare callback for local walk
spencerw Nov 13, 2024
5eb603a
First attempt at consolidated Ewald GPU kernel
spencerw Nov 13, 2024
85191fe
Merge remote-tracking branch 'remotes/origin/dm_tp' into kernelfix
spencerw Nov 14, 2024
37ba423
More fixes to consolidated Ewald
spencerw Nov 18, 2024
336b959
Ensure device memory pointers are passed to TreePieces even though co…
spencerw Nov 19, 2024
ee3d927
Call finishBucket after local tree walk, ensure smooth happens before…
spencerw Nov 19, 2024
aa919f8
Pass nReplicas and fPeriod to DataManagerLocalTreeWalk
spencerw Nov 20, 2024
40de4f5
Removed unused Ewald GPU code
spencerw Nov 20, 2024
6b012e5
Ewald GPU data to pinned host memory, remove markers
spencerw Nov 20, 2024
e4cb17f
Comments and code cleanup
spencerw Nov 20, 2024
2866fdb
Fix for bucket bookkeeping
spencerw Nov 21, 2024
bce757c
Temporarily move initiatePrefetch to commenceCalculateGravityLocal
spencerw Dec 9, 2024
e2a9686
Interaction list kernels launch once per PE
spencerw Apr 15, 2025
87a10dd
Bug fixes for Ewald and multistepping
spencerw Apr 17, 2025
0847e56
GPU local walk and Ewald callbacks use treeProxy
spencerw Apr 28, 2025
ae0c589
Remove small phase code
spencerw Apr 30, 2025
c886eed
Remove CUDA_2D_FLAT option
spencerw Apr 30, 2025
de673ee
Use HAPI callbacks for PEList functions
spencerw Apr 30, 2025
08686a6
dMProxy mem log writer
mrcawood Apr 30, 2025
7dc78bd
fixed headers
mrcawood May 1, 2025
08cde7a
CUDA streams handled by PE groups
spencerw May 1, 2025
01ae6d5
fixed init, added nodeID
mrcawood May 1, 2025
274ff25
PEList vectors use pinned host memory
spencerw May 2, 2025
5112ab0
merging into kernelfix
mrcawood May 5, 2025
1372e29
ifdefs around GPU code
spencerw May 5, 2025
3db6c66
added origin function tag
mrcawood May 5, 2025
dfb999d
Fix DataManager creating CUDA stream too early
spencerw May 11, 2025
11c0574
Callbacks for PEList functions managed by class
spencerw May 28, 2025
a6ac836
Fix memory leaks
spencerw Jun 3, 2025
631a2f0
Various code cleanup
spencerw Jun 10, 2025
bc97624
updates for memory logger
mrcawood Jul 7, 2025
58d9c60
Merge branch 'kernelfix' into kernelfix
spencerw Jul 8, 2025
97479f4
Fixes to Ewald on GPU
spencerw Jul 9, 2025
40af695
Merge pull request #3 from mrcawood/kernelfix
spencerw Jul 14, 2025
0c8da91
Cleanup for GPU memory logger
spencerw Jul 14, 2025
c1e5eb0
Remove HAPI_MEMPOOL config option
spencerw Jul 16, 2025
8cb9796
Fixes to CUDA build options
spencerw Jul 18, 2025
1344955
Comments and docstrings
spencerw Jul 21, 2025
f644e62
Fix hang when a PE has no TreePieces
spencerw Jul 23, 2025
dd205c4
Performance tweaks
spencerw Aug 8, 2025
20aa6ef
Move initiatePrefetch back to startGravity
spencerw Aug 15, 2025
9510da7
Fix PINNED_HOST_MEMORY option
spencerw Aug 27, 2025
281fe39
Free remote buffers in resumeRemoteChunk
spencerw Sep 5, 2025
dd79d45
Fixes according to Tom's PR comments
spencerw Sep 4, 2025
d40b963
Move TreePiece::EwaldGPU code to DataManager::startEwald
spencerw Sep 9, 2025
95c3e15
Ensure Ewald still triggers if GPU_LOCAL_TREE_WALK is disabled
spencerw Sep 9, 2025
12b5f21
Various minor fixes
spencerw Sep 9, 2025
89fae31
Add back in PinnedHostAllocator for PEList
spencerw Oct 1, 2025
de3f9d2
integrated memory pools \w warmup back into upstream kernelfix
mrcawood Oct 1, 2025
3f25c73
increased host pool warmup buckets
mrcawood Oct 1, 2025
433bd8e
added logging to CPU memory pool with +cpumemlog
mrcawood Oct 2, 2025
a77533e
removed large long lived mallocs from host mem pool
mrcawood Oct 2, 2025
2283f08
Many improvements for host memory pool, include trimming and refill. …
mrcawood Oct 14, 2025
f7f918d
removed large mallocs from pool, further tuned warmup and trim params
mrcawood Oct 21, 2025
9c5b451
Allow remote walk before GPU data transfer
spencerw Oct 20, 2025
e0f384e
resumeRemoteChunk iterates over PEList proxies correctly
spencerw Oct 24, 2025
9d887ee
fillGPUBuffer given priority over SPH
spencerw Oct 29, 2025
10ed833
removed stale code
mrcawood Nov 6, 2025
4b4fc4c
Merge remote-tracking branch 'upstream/kernelfix' into kernelfix
mrcawood Nov 25, 2025
ff764d5
Fix thread safety: synchronize shared data access and prevent deadloc…
mrcawood Nov 25, 2025
b688f14
* Fix hang when an entire PE has no particles
spencerw Nov 25, 2025
dc56105
Merge remote-tracking branch 'upstream/master' into kernelfix
spencerw Dec 2, 2025
967e408
Fix GPU memory pool race condition causing non-determinism
mrcawood Dec 3, 2025
19758e5
Merge remote-tracking branch 'matt/kernelfix' into kernelfix
spencerw Dec 3, 2025
789cce3
added more null pointer guardrails
mrcawood Dec 8, 2025
8426073
Rework PEList callbacks
spencerw Dec 9, 2025
6de2c13
Remove pinned host allocator
spencerw Dec 9, 2025
c169dba
Fix minor memory leaks
spencerw Dec 11, 2025
3c48a26
Remove stream init from DataManager constructor
spencerw Dec 11, 2025
04724af
Merge pull request #5 from mrcawood/kernelfix
spencerw Dec 12, 2025
8314b2a
TreePiece::finishWalk(): make HAPI_TRACE less verbose.
trquinn Jan 19, 2026
87a10b8
PEList and DataManager: add booleans to track data transfer to the GPU.
trquinn Jan 20, 2026
2b486c8
Synchronously call EwaldInit() to avoid use-before-init problem.
trquinn Jan 25, 2026
92b7598
PEList::launch_kernel(): remove asserts that are not valid.
trquinn Feb 23, 2026
42c9a1d
Host pool init race fix + MemoryPool review cleanups
mrcawood Feb 24, 2026
8ec8b4a
Fix SMP memory visibility in remote gravity data transfer gating
mrcawood Feb 23, 2026
470e5af
Fix startEwaldGPU crash when no TreePiece has valid Ewald data (#13)
mrcawood Mar 6, 2026
cbf84f0
Allocate particleInterRemote before finishWalk to avoid NULL derefere…
mrcawood Mar 6, 2026
292715f
Inactive TreePiece optimization (#18)
trquinn Mar 18, 2026
c985482
Fix PEList launch readiness when local transfer lags remote transfer …
mrcawood Mar 18, 2026
c1dfed7
TreePiece::updateParticles(): skip inactive TreePieces.
trquinn May 4, 2026
3eb45b2
class PEList: remove unused attribute.
trquinn Mar 24, 2026
82c6f53
TreePiece constructors: initialize nEwhLoop.
trquinn May 9, 2026
5d7c5e0
White space cleanup.
trquinn May 9, 2026
ad49143
TreePiece::startGravity(): empty TPs don't checkin with PEList::finis…
trquinn Jul 21, 2026
0b9139e
PEList: change TP checkin counter from vector to an int.
trquinn Jul 21, 2026
96f3800
ListCompute::stateReady(): deleted unused variables.
trquinn Aug 22, 2026
318b06c
class TreePiece: consolidate numActiveParticles and myNumActiveParticles
trquinn Aug 22, 2026
cfe9bf9
DataManager::transformLocalTreeRecursive(): don't process NonLocal no…
trquinn Aug 30, 2026
8829b1b
Add a TreePiece entry method to wait for GPU buffer to be filled.
trquinn Sep 6, 2026
fd48e77
class PEList: remove unused attributes.
trquinn Sep 10, 2026
e427a8e
Normalized arguments to CUDA Ewald call.
trquinn Sep 19, 2026
41121ff
Make EwaldInit() a DataManager method (#23)
trquinn Sep 19, 2026
4e22f13
DataManager: set cuda device and stream on restart from checkpoint.
trquinn Sep 22, 2026
5338c19
Main::setupICs(): move initMemLog() out of if(bStarForm) block.
trquinn Sep 22, 2026
6100e6c
Fixed small comment change.
trquinn Sep 22, 2026
437dd36
Host pool diagnostics: replace sprintf with make_formatted_string().
trquinn Sep 23, 2026
f0983b2
Merge remote-tracking branch 'upstream_public/master' into kernelfix
trquinn Sep 25, 2026
a04413f
CUDA: allocate ewtGPU buffer once for the entire run.
trquinn Sep 26, 2026
1af6bd4
Merge branch 'trq/one_ewt_alloc' into kernelfix
trquinn Sep 26, 2026
4c9769f
CUDA: initialize memory logging on restart.
trquinn Sep 27, 2026
778fa59
TreePiece class with CUDA: eliminate unused attributes.
trquinn Sep 28, 2026
cac97c7
CUDA: remove transfer of acceleration vector to the GPU.
trquinn Sep 28, 2026
de32d30
CUDA: add check for Ewald completion before transferParticlesBack().
trquinn Sep 30, 2026
0f217eb
calculateEwaldUsingCkLoop(): delete unused parameter.
trquinn Sep 30, 2026
0a61303
HostCUDA.cu: remove debugging cudaPeekError()s
trquinn Oct 8, 2026
d6bf8bd
Merge branch 'master' of github.com:N-BodyShop/changa into kernelfix
trquinn Oct 8, 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
429 changes: 77 additions & 352 deletions Compute.cpp

Large diffs are not rendered by default.

3 changes: 0 additions & 3 deletions Compute.h
Original file line number Diff line number Diff line change
Expand Up @@ -185,9 +185,6 @@ class ListCompute : public Compute{
void enableCpu() {bUseCpu = 1;}

#ifdef CUDA
#ifdef GPU_LOCAL_TREE_WALK
void sendLocalTreeWalkTriggerToGpu(State *state, TreePiece *tp, int activeRung, int startBucket, int endBucket);
#endif //GPU_LOCAL_TREE_WALK
void sendNodeInteractionsToGpu(DoubleWalkState *state, TreePiece *tp);
void sendPartInteractionsToGpu(DoubleWalkState *state, TreePiece *tp);
#endif
Expand Down
17 changes: 2 additions & 15 deletions CudaFunctions.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,8 +4,8 @@
#ifdef CUDA
#include "hapi.h"
#include "HostCUDA.h"
void TreePieceDataTransferBasic(CudaRequest *data, CudaDevPtr *ptr);
void TreePieceDataTransferBasicCleanup(CudaDevPtr *ptr);
void DataTransferBasic(CudaRequest *data, CudaDevPtr *ptr, const char* functionTag);
void DataTransferBasicCleanup(CudaDevPtr *ptr, cudaStream_t stream, const char* functionTag);

#ifdef GPU_LOCAL_TREE_WALK
__global__ void gpuLocalTreeWalk(
Expand Down Expand Up @@ -36,7 +36,6 @@ __global__ void nodeGravityComputation(

__global__ void ZeroVars(VariablePartData *particleVars, int nVars);

#ifdef CUDA_2D_TB_KERNEL
__global__ void particleGravityComputation(
CompactPartData *targetCores,
VariablePartData *targetVars,
Expand All @@ -46,18 +45,6 @@ __global__ void particleGravityComputation(
int *bucketStarts,
int *bucketSizes,
cudatype fperiod);
#else
__global__ void particleGravityComputation(
CompactPartData *targetCores,
VariablePartData *targetVars,
CompactPartData *sourceCores,
ILPart *ils,
int *ilmarks,
int *bucketStarts,
int *bucketSizes,
cudatype fperiod);
#endif

#endif

#endif
Loading
Loading