Skip to content

Commit

Permalink
[bug] fix CPU host fails to save trajectory, fix GPU traj duplicate
Browse files Browse the repository at this point in the history 
use omp atomic instead of critical in savedebugdata; otherwise speed drop by 25x
remove omp critical for photon sharing that Ruoyang added
  • Loading branch information
@fangq
fangq committed Mar 9, 2024
1 parent b0f9cfe commit dd9b4d5
Show file tree
Hide file tree
Showing 5 changed files with 41 additions and 28 deletions.
6 changes: 3 additions & 3 deletions src/mmc_cl_host.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -621,15 +621,15 @@ void mmc_run_cl(mcconfig* cfg, tetmesh* mesh, raytracer* tracer) {
reporter.jumpdebug += rep.jumpdebug;

if (cfg->debuglevel & dlTraj) {
uint debugrec = reporter.jumpdebug;
uint debugrec = rep.jumpdebug;

if (debugrec > 0) {
if (debugrec > cfg->maxdetphoton) {
if (debugrec > cfg->maxjumpdebug) {
MMC_FPRINTF(cfg->flog, S_RED "WARNING: the saved trajectory positions (%d) \
are more than what your have specified (%d), please use the --maxjumpdebug option to specify a greater number\n" S_RESET
, debugrec, cfg->maxjumpdebug);
} else {
MMC_FPRINTF(cfg->flog, "saved %u trajectory positions, total: %d\t", debugrec, cfg->maxjumpdebug + debugrec);
MMC_FPRINTF(cfg->flog, "saved %u trajectory positions, total: %d\t", debugrec, cfg->debugdatalen + debugrec);
}

debugrec = MIN(debugrec, cfg->maxjumpdebug);
Expand Down
4 changes: 0 additions & 4 deletions src/mmc_core.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1585,10 +1585,6 @@ __device__ void onephoton(unsigned int id, __local float* ppath, __constant MCXP
}
}

if (GPU_PARAM(gcfg, debuglevel) & dlTraj) {
savedebugdata(&r, id, reporter, gdebugdata, gcfg);
}

float mom = 0.f;
r.slen0 = mc_next_scatter(gmed[type[r.eid - 1]].g, &r.vec, ran, gcfg, &mom);
r.slen = r.slen0;
Expand Down
19 changes: 19 additions & 0 deletions src/mmc_host.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -407,13 +407,32 @@ int mmc_run_mp(mcconfig* cfg, tetmesh* mesh, raytracer* tracer) {
}
}

if ((cfg->debuglevel & dlTraj) && cfg->exportdebugdata) {
if (cfg->debugdatalen > cfg->maxjumpdebug) {
MMC_FPRINTF(cfg->flog, S_RED "WARNING: the saved trajectory positions (%d) \
are more than what your have specified (%d), please use the --maxjumpdebug option to specify a greater number\n" S_RESET
, cfg->debugdatalen, cfg->maxjumpdebug);
cfg->debugdatalen = cfg->maxjumpdebug;
} else {
MMC_FPRINTF(cfg->flog, "saved %u trajectory positions, total: %d\t", cfg->debugdatalen, cfg->debugdatalen);
}
}

#ifndef MCX_CONTAINER

if (cfg->issaveref) {
MMCDEBUG(cfg, dlTime, (cfg->flog, "saving surface diffuse reflectance ..."));
mesh_saveweight(mesh, cfg, 1);
}

if ((cfg->debuglevel & dlTraj) && cfg->parentid == mpStandalone && cfg->exportdebugdata) {
cfg->his.colcount = MCX_DEBUG_REC_LEN;
cfg->his.savedphoton = cfg->debugdatalen;
cfg->his.totalphoton = cfg->nphoton;
cfg->his.detected = 0; // this flag tells mcx_savedetphoton that the data is trajectory
mesh_savedetphoton(cfg->exportdebugdata, NULL, cfg->debugdatalen, 0, cfg);
}

#endif

MMCDEBUG(cfg, dlTime, (cfg->flog, "\tdone\t%d\n", GetTimeMillis() - t0));
Expand Down
1 change: 1 addition & 0 deletions src/mmc_mesh.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -309,6 +309,7 @@ void mesh_loadnode(tetmesh* mesh, mcconfig* cfg) {
if (cfg->method == rtBLBadouelGrid) {
mesh_createdualmesh(mesh, cfg);
}

return;
}

Expand Down
39 changes: 18 additions & 21 deletions src/mmc_raytrace.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1727,9 +1727,10 @@ void savedebugdata(ray* r, unsigned int id, mcconfig* cfg) {
unsigned int pos;
float* gdebugdata = cfg->exportdebugdata;

#pragma omp critical
#pragma omp atomic capture
{
pos = cfg->debugdatalen++;
pos = cfg->debugdatalen;
cfg->debugdatalen++;
}

if (pos < cfg->maxjumpdebug) {
Expand Down Expand Up @@ -1811,19 +1812,17 @@ void onephoton(size_t id, raytracer* tracer, tetmesh* mesh, mcconfig* cfg,
visit->launchweight[0] = kahant;
} else {
int psize = (int)cfg->srcparam1.w * (int)cfg->srcparam2.w;
#pragma omp critical
{
for (pidx = 0; pidx < cfg->srcnum; pidx++) {
if (cfg->seed == SEED_FROM_FILE && (cfg->outputtype == otWL || cfg->outputtype == otWP)) {
kahany = cfg->replayweight[r.photonid] - visit->kahanc0[pidx]; /* when replay mode, accumulate detected photon weight */
} else {
kahany = r.weight * cfg->srcpattern[pidx * psize + r.posidx] - visit->kahanc0[pidx];
}

kahant = visit->launchweight[pidx] + kahany;
visit->kahanc0[pidx] = (kahant - visit->launchweight[pidx]) - kahany;
visit->launchweight[pidx] = kahant;
for (pidx = 0; pidx < cfg->srcnum; pidx++) {
if (cfg->seed == SEED_FROM_FILE && (cfg->outputtype == otWL || cfg->outputtype == otWP)) {
kahany = cfg->replayweight[r.photonid] - visit->kahanc0[pidx]; /* when replay mode, accumulate detected photon weight */
} else {
kahany = r.weight * cfg->srcpattern[pidx * psize + r.posidx] - visit->kahanc0[pidx];
}

kahant = visit->launchweight[pidx] + kahany;
visit->kahanc0[pidx] = (kahant - visit->launchweight[pidx]) - kahany;
visit->launchweight[pidx] = kahant;
}
}

Expand Down Expand Up @@ -2085,14 +2084,12 @@ void onephoton(size_t id, raytracer* tracer, tetmesh* mesh, mcconfig* cfg,
visit->absorbweight[0] = kahant;
} else {
int psize = (int)cfg->srcparam1.w * (int)cfg->srcparam2.w;
#pragma omp critical
{
for (pidx = 0; pidx < cfg->srcnum; pidx++) {
kahany = r.Eabsorb * cfg->srcpattern[pidx * psize + r.posidx] - visit->kahanc1[pidx];
kahant = visit->absorbweight[pidx] + kahany;
visit->kahanc1[pidx] = (kahant - visit->absorbweight[pidx]) - kahany;
visit->absorbweight[pidx] = kahant;
}

for (pidx = 0; pidx < cfg->srcnum; pidx++) {
kahany = r.Eabsorb * cfg->srcpattern[pidx * psize + r.posidx] - visit->kahanc1[pidx];
kahant = visit->absorbweight[pidx] + kahany;
visit->kahanc1[pidx] = (kahant - visit->absorbweight[pidx]) - kahany;
visit->absorbweight[pidx] = kahant;
}
}
}
Expand Down

0 comments on commit dd9b4d5

Please sign in to comment.