Skip to content

Commit c7df8d9

Browse files
committed
Fix memory error by unifying handling of otRF and otRFmus
1 parent 7e7d88b commit c7df8d9

File tree

4 files changed

+21
-21
lines changed

4 files changed

+21
-21
lines changed

src/mcx_core.cu

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1648,7 +1648,7 @@ __device__ inline int launchnewphoton(MCXpos* p, MCXdir* v, Stokes* s, MCXtime*
16481648
ppath[2] = ((gcfg->srcnum > 1) ? ppath[2] : p->w); // store initial weight
16491649
v->nscat = EPS;
16501650

1651-
if (gcfg->outputtype == otRF || gcfg->outputtype == otRFmus) { // if run RF replay
1651+
if (gcfg->outputtype == otRF || gcfg->outputtype == otRFmus) { // if run RF replay
16521652
f->pathlen = photontof[(threadid * gcfg->threadphoton + min(threadid, gcfg->oddphotons - 1) + (int)f->ndone)];
16531653
sincosf(gcfg->omega * f->pathlen, ppath + 5 + gcfg->srcnum, ppath + 4 + gcfg->srcnum);
16541654
}
@@ -2007,7 +2007,7 @@ __global__ void mcx_main_loop(uint media[], OutputType field[], float genergy[],
20072007
#endif
20082008
}
20092009

2010-
if (gcfg->outputtype == otRFmus) {
2010+
if (gcfg->seed == SEED_FROM_FILE && gcfg->outputtype == otRFmus) {
20112011
OutputType w_N_scatt = replayweight[(idx * gcfg->threadphoton + min(idx, gcfg->oddphotons - 1) + (int)f.ndone)];
20122012
OutputType cos_omega_t = ppath[gcfg->w0offset + gcfg->srcnum];
20132013
OutputType sin_omega_t = ppath[gcfg->w0offset + gcfg->srcnum + 1];
@@ -2880,9 +2880,9 @@ void mcx_run_simulation(Config* cfg, GPUInfo* gpu) {
28802880
{
28812881
if (cfg->exportfield == NULL && cfg->issave2pt) {
28822882
if (cfg->seed == SEED_FROM_FILE && cfg->replaydet == -1) {
2883-
cfg->exportfield = (float*)calloc(sizeof(float) * dimxyz, gpu[gpuid].maxgate * (1 + (cfg->outputtype == otRF)) * cfg->detnum);
2883+
cfg->exportfield = (float*)calloc(sizeof(float) * dimxyz, gpu[gpuid].maxgate * (1 + (cfg->outputtype == otRF || cfg->outputtype == otRFmus)) * cfg->detnum);
28842884
} else {
2885-
cfg->exportfield = (float*)calloc(sizeof(float) * dimxyz, gpu[gpuid].maxgate * (1 + (cfg->outputtype == otRF)));
2885+
cfg->exportfield = (float*)calloc(sizeof(float) * dimxyz, gpu[gpuid].maxgate * (1 + (cfg->outputtype == otRF || cfg->outputtype == otRFmus)));
28862886
}
28872887
}
28882888

@@ -3302,7 +3302,7 @@ void mcx_run_simulation(Config* cfg, GPUInfo* gpu) {
33023302
*
33033303
* The calculation of the energy conservation will only reflect the last simulation.
33043304
*/
3305-
sharedbuf = (param.nphaselen + param.nanglelen) * sizeof(float) + gpu[gpuid].autoblock * (cfg->issaveseed * (RAND_BUF_LEN * sizeof(RandType)) + sizeof(float) * (param.w0offset + cfg->srcnum + 2 * (cfg->outputtype == otRF)));
3305+
sharedbuf = (param.nphaselen + param.nanglelen) * sizeof(float) + gpu[gpuid].autoblock * (cfg->issaveseed * (RAND_BUF_LEN * sizeof(RandType)) + sizeof(float) * (param.w0offset + cfg->srcnum + 2 * (cfg->outputtype == otRF || cfg->outputtype == otRFmus)));
33063306

33073307
MCX_FPRINTF(cfg->flog, "requesting %d bytes of shared memory\n", sharedbuf);
33083308

@@ -3656,14 +3656,14 @@ is more than what your have specified (%d), please use the -H option to specify
36563656
field[i] = rawfield[i];
36573657
#ifndef USE_DOUBLE
36583658

3659-
if (cfg->outputtype != otRF) {
3659+
if (cfg->outputtype != otRF && cfg->outputtype != otRFmus) {
36603660
field[i] += rawfield[i + fieldlen];
36613661
}
36623662

36633663
#endif
36643664
}
36653665

3666-
if (cfg->outputtype == otRF && cfg->omega > 0.f && SHADOWCOUNT == 2) {
3666+
if ((cfg->outputtype == otRF || cfg->outputtype == otRFmus) && cfg->omega > 0.f && SHADOWCOUNT == 2) {
36673667
rfimag = (OutputType*)malloc(fieldlen * sizeof(OutputType));
36683668
memcpy(rfimag, rawfield + fieldlen, fieldlen * sizeof(OutputType));
36693669
}
@@ -3720,7 +3720,7 @@ is more than what your have specified (%d), please use the -H option to specify
37203720
#pragma omp atomic
37213721
cfg->exportfield[i] += field[i];
37223722

3723-
if (cfg->outputtype == otRF && rfimag) {
3723+
if ((cfg->outputtype == otRF || cfg->outputtype == otRFmus) && rfimag) {
37243724
for (i = 0; i < fieldlen; i++)
37253725
#pragma omp atomic
37263726
cfg->exportfield[i + fieldlen] += rfimag[i];
@@ -3807,7 +3807,7 @@ is more than what your have specified (%d), please use the -H option to specify
38073807
}
38083808
} else if (cfg->outputtype == otEnergy || cfg->outputtype == otL) { /** If output is energy (joule), raw data is simply multiplied by 1/Nphoton */
38093809
scale[0] = 1.f / cfg->energytot;
3810-
} else if (cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF) {
3810+
} else if (cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF || cfg->outputtype == otRFmus) {
38113811
if (cfg->seed == SEED_FROM_FILE && cfg->replaydet == -1) {
38123812
int detid;
38133813

@@ -3831,7 +3831,7 @@ is more than what your have specified (%d), please use the -H option to specify
38313831
fflush(cfg->flog);
38323832
mcx_normalize(cfg->exportfield + (detid - 1)*dimxyz * gpu[gpuid].maxgate, scale[0], dimxyz * gpu[gpuid].maxgate, cfg->isnormalized, 0, 1);
38333833

3834-
if (cfg->outputtype == otRF) {
3834+
if (cfg->outputtype == otRF || cfg->outputtype == otRFmus) {
38353835
mcx_normalize(cfg->exportfield + fieldlen + (detid - 1)*dimxyz * gpu[gpuid].maxgate, scale[0], dimxyz * gpu[gpuid].maxgate, cfg->isnormalized, 0, 1);
38363836
}
38373837
}
@@ -3880,7 +3880,7 @@ is more than what your have specified (%d), please use the -H option to specify
38803880
for (i = 0; i < (int)cfg->srcnum; i++) {
38813881
MCX_FPRINTF(cfg->flog, "source %d, normalization factor alpha=%f\n", (i + 1), scale[i]);
38823882
fflush(cfg->flog);
3883-
mcx_normalize(cfg->exportfield, scale[i], fieldlen / cfg->srcnum * ((cfg->outputtype == otRF) + 1), cfg->isnormalized, i, cfg->srcnum);
3883+
mcx_normalize(cfg->exportfield, scale[i], fieldlen / cfg->srcnum * ((cfg->outputtype == otRF || cfg->outputtype == otRFmus) + 1), cfg->isnormalized, i, cfg->srcnum);
38843884
}
38853885
}
38863886

src/mcx_utils.c

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -893,7 +893,7 @@ void mcx_savedata(float* dat, size_t len, Config* cfg) {
893893
}
894894

895895
if (cfg->outputformat == ofNifti || cfg->outputformat == ofAnalyze) {
896-
mcx_savenii(dat, len * (1 + (cfg->outputtype == otRF)), name, NIFTI_TYPE_FLOAT32, cfg->outputformat, cfg);
896+
mcx_savenii(dat, len * (1 + (cfg->outputtype == otRF || cfg->outputtype == otRFmus)), name, NIFTI_TYPE_FLOAT32, cfg->outputformat, cfg);
897897
return;
898898
} else if (cfg->outputformat == ofJNifti || cfg->outputformat == ofBJNifti) {
899899
uint dims[6] = {cfg->dim.x, cfg->dim.y, cfg->dim.z, cfg->maxgate, cfg->srcnum, 1};
@@ -907,7 +907,7 @@ void mcx_savedata(float* dat, size_t len, Config* cfg) {
907907
dims[5] *= cfg->detnum;
908908
}
909909

910-
if (cfg->seed == SEED_FROM_FILE && cfg->outputtype == otRF) {
910+
if (cfg->seed == SEED_FROM_FILE && (cfg->outputtype == otRF || cfg->outputtype == otRFmus )) {
911911
dims[5] *= 2;
912912
}
913913

@@ -932,7 +932,7 @@ void mcx_savedata(float* dat, size_t len, Config* cfg) {
932932
fwrite(&(cfg->dim.x), sizeof(int), 3, fp);
933933
}
934934

935-
fwrite(dat, sizeof(float), len * (1 + (cfg->outputtype == otRF)), fp);
935+
fwrite(dat, sizeof(float), len * (1 + (cfg->outputtype == otRF || cfg->outputtype == otRFmus)), fp);
936936
fclose(fp);
937937
}
938938

@@ -1494,7 +1494,7 @@ void mcx_preprocess(Config* cfg) {
14941494
cfg->issavedet = 0;
14951495
}
14961496

1497-
if ((cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF)
1497+
if ((cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF || cfg->outputtype == otRFmus)
14981498
&& cfg->seed != SEED_FROM_FILE) {
14991499
MCX_ERROR(-6, "Jacobian output is only valid in the reply mode. Please define cfg.seed");
15001500
}
@@ -3956,7 +3956,7 @@ void mcx_loadseedfile(Config* cfg) {
39563956
cfg->seed = SEED_FROM_FILE;
39573957
cfg->nphoton = his.savedphoton;
39583958

3959-
if (cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF) { //cfg->replaydet>0
3959+
if (cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF || cfg->outputtype == otRFmus) { //cfg->replaydet>0
39603960
int i, j, hasdetid = 0, offset;
39613961
float plen, *ppath;
39623962
hasdetid = SAVE_DETID(his.savedetflag);
@@ -5210,7 +5210,7 @@ void mcx_parsecmd(int argc, char* argv[], Config* cfg) {
52105210
}
52115211
}
52125212

5213-
if ((cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == otRF) && cfg->seed != SEED_FROM_FILE) {
5213+
if ((cfg->outputtype == otJacobian || cfg->outputtype == otWP || cfg->outputtype == otDCS || cfg->outputtype == || cfg->outputtype == otRFmus) && cfg->seed != SEED_FROM_FILE) {
52145214
MCX_ERROR(-1, "Jacobian output is only valid in the reply mode. Please give an mch file after '-E'.");
52155215
}
52165216

src/mcxlab.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -267,7 +267,7 @@ void mexFunction(int nlhs, mxArray* plhs[], int nrhs, const mxArray* prhs[]) {
267267
fieldlen *= cfg.detnum;
268268
}
269269

270-
if (cfg.replay.seed != NULL && cfg.outputtype == otRF) {
270+
if (cfg.replay.seed != NULL && (cfg.outputtype == otRF || cfg.outputtype == otRFmus)) {
271271
fieldlen *= 2;
272272
}
273273

@@ -402,7 +402,7 @@ void mexFunction(int nlhs, mxArray* plhs[], int nrhs, const mxArray* prhs[]) {
402402
fielddim[4] = cfg.detnum;
403403
}
404404

405-
if (cfg.replay.seed != NULL && cfg.outputtype == otRF) {
405+
if (cfg.replay.seed != NULL && (cfg.outputtype == otRF || cfg.outputtype == otRFmus)) {
406406
fielddim[5] = 2;
407407
}
408408

src/pmcx.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1099,7 +1099,7 @@ py::dict pmcx_interface(const py::dict& user_cfg) {
10991099
field_len *= mcx_config.detnum;
11001100
}
11011101

1102-
if (mcx_config.replay.seed != nullptr && mcx_config.outputtype == otRF) {
1102+
if (mcx_config.replay.seed != nullptr && (mcx_config.outputtype == otRF || cfg->outputtype == otRFmus)) {
11031103
field_len *= 2;
11041104
}
11051105

@@ -1228,7 +1228,7 @@ py::dict pmcx_interface(const py::dict& user_cfg) {
12281228
field_dim[4] = mcx_config.detnum;
12291229
}
12301230

1231-
if (mcx_config.replay.seed != nullptr && mcx_config.outputtype == otRF) {
1231+
if (mcx_config.replay.seed != nullptr && (mcx_config.outputtype == otRF || cfg->outputtype == otRFmus)) {
12321232
field_dim[5] = 2;
12331233
}
12341234

0 commit comments

Comments
 (0)