From 6baa839c136b583e3699b2214579a5b0cfbfcfe7 Mon Sep 17 00:00:00 2001 From: Qianqian Fang Date: Sun, 3 Nov 2024 18:22:37 -0500 Subject: [PATCH] [bug] fix incorrect shared memory initialization, fix #101, fix #103 --- mmclab/example/demo_example_replay.m | 1 + src/mmc_core.cl | 7 ++++--- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/mmclab/example/demo_example_replay.m b/mmclab/example/demo_example_replay.m index e7ef49d..1909290 100644 --- a/mmclab/example/demo_example_replay.m +++ b/mmclab/example/demo_example_replay.m @@ -36,6 +36,7 @@ cfg.issaveexit = 1; % save detected photon exit position and angles cfg.issaveseed = 1; % save detected photon seeds to replay later % cfg.method=0; +cfg.compute='cuda'; newcfg = mmclab(cfg, 'prep'); % preprocessing of the mesh to get the missing fields diff --git a/src/mmc_core.cl b/src/mmc_core.cl index 764c8ca..7dbf0df 100644 --- a/src/mmc_core.cl +++ b/src/mmc_core.cl @@ -1455,6 +1455,10 @@ __device__ void onephoton(unsigned int id, __local float* ppath, __constant MCXP RandType initseed[RAND_BUF_LEN]; #endif + clearpath(ppath, (GPU_PARAM(gcfg, reclen) + (GPU_PARAM(gcfg, srcnum) > 1) * GPU_PARAM(gcfg, srcnum))); + clearpath(energytot, GPU_PARAM(gcfg, srcnum)); + clearpath(energyesc, GPU_PARAM(gcfg, srcnum)); + r.photonid = id; #ifdef MCX_SAVE_SEED @@ -1734,9 +1738,6 @@ __kernel void mmc_main_loop(const int nphoton, const int ophoton, gpu_rng_init(t, n_seed, idx); } - clearpath(sharedmem, get_local_size(0) * ((GPU_PARAM(gcfg, srcnum) << 1) + - (GPU_PARAM(gcfg, reclen) + (GPU_PARAM(gcfg, srcnum) > 1) * GPU_PARAM(gcfg, srcnum)))); - /*launch photons*/ for (int i = 0; i < nphoton + (idx < ophoton); i++) { if (GPU_PARAM(gcfg, seed) == SEED_FROM_FILE)