Skip to content

Commit

Permalink
remove opencl jit warnings, adjust optlevel order
Browse files Browse the repository at this point in the history
  • Loading branch information
fangq committed Oct 7, 2023
1 parent 1a30e5c commit e930448
Show file tree
Hide file tree
Showing 3 changed files with 14 additions and 14 deletions.
12 changes: 6 additions & 6 deletions src/mcx_core.cl
Original file line number Diff line number Diff line change
Expand Up @@ -875,7 +875,7 @@ int launchnewphoton(float4* p, float4* v, float4* f, short4* flipdir, FLOAT4VEC*
#ifdef MCX_SAVE_DETECTORS

// let's handle detectors here
if ((isdet & DET_MASK) == DET_MASK && *mediaid == 0 && GPU_PARAM(gcfg, issaveref) < 2) {
if ((isdet & DET_MASK) == DET_MASK && *mediaid == 0 && (bool)(GPU_PARAM(gcfg, issaveref) < 2)) {
savedetphoton(n_det, dpnum, ppath, p, v, photonseed, gseeddata, gdetpos, gcfg, isdet);
}

Expand Down Expand Up @@ -1372,7 +1372,7 @@ __kernel void mcx_main_loop(__global const uint* media,

v.w += 1.f;

if (GPU_PARAM(gcfg, outputtype) == otWP || GPU_PARAM(gcfg, outputtype) == otDCS) {
if ((bool)(GPU_PARAM(gcfg, outputtype) == otWP) || (bool)(GPU_PARAM(gcfg, outputtype) == otDCS)) {
///< photontof[] and replayweight[] should be cached using local mem to avoid global read
int tshift = (idx * gcfg->threadphoton + min(idx, gcfg->oddphoton - 1) + (int)f.w);
tmp0 = (GPU_PARAM(gcfg, outputtype) == otDCS) ? (1.f - ctheta) : 1.f;
Expand Down Expand Up @@ -1472,7 +1472,7 @@ __kernel void mcx_main_loop(__global const uint* media,
/** calculate the quality to be accummulated */
if (GPU_PARAM(gcfg, outputtype) == otEnergy) {
weight = w0 - p.w;
} else if (GPU_PARAM(gcfg, outputtype) == otFluence || GPU_PARAM(gcfg, outputtype) == otFlux) {
} else if ((bool)(GPU_PARAM(gcfg, outputtype) == otFluence) || (bool)(GPU_PARAM(gcfg, outputtype) == otFlux)) {
weight = (prop.x * f.z < 0.001f) ? (w0 * f.z) : ((w0 - p.w) / (prop.x));
} else if (GPU_PARAM(gcfg, seed) == SEED_FROM_FILE) {
if (GPU_PARAM(gcfg, outputtype) == otJacobian) {
Expand Down Expand Up @@ -1523,8 +1523,8 @@ __kernel void mcx_main_loop(__global const uint* media,
}

/** launch new photon when exceed time window or moving from non-zero voxel to zero voxel without reflection */
if ((mediaid == 0 && ((!GPU_PARAM(gcfg, doreflect) || (GPU_PARAM(gcfg, doreflect) && n1 == gproperty[0].w)) || (((isdet & 0xF) == bcUnknown && !GPU_PARAM(gcfg, doreflect))
|| (isdet & 0xF) == bcAbsorb || (isdet & 0xF) == bcCyclic)) && (isdet & 0xF) != bcMirror && (isdet & 0xF) != bcReflect) || f.y > gcfg->twin1) {
if ((mediaid == 0 && (bool)(((bool)(!(GPU_PARAM(gcfg, doreflect))) || ((bool)(GPU_PARAM(gcfg, doreflect)) && n1 == gproperty[0].w)) || (((isdet & 0xF) == bcUnknown && (bool)(!(GPU_PARAM(gcfg, doreflect))))
|| (isdet & 0xF) == bcAbsorb || (isdet & 0xF) == bcCyclic)) && (isdet & 0xF) != bcMirror && (isdet & 0xF) != bcReflect) || f.y > gcfg->twin1) {
if (isdet == bcCyclic) {
if (flipdir.w == 0) {
p.x = mcx_nextafterf(convert_float_rte(((idx1d == OUTSIDE_VOLUME_MIN) ? gcfg->maxidx.x : 0)), (v.x > 0.f) - (v.x < 0.f));
Expand Down Expand Up @@ -1578,7 +1578,7 @@ __kernel void mcx_main_loop(__global const uint* media,
|| (mediaid == 0 && // or if out of bbx or enters 0-voxel
(((isdet & 0xF) == bcUnknown && GPU_PARAM(gcfg, doreflect)) // if cfg.bc is "_", check cfg.isreflect
|| (((isdet & 0xF) == bcReflect || (isdet & 0xF) == bcMirror))))) // or if cfg.bc is 'r' or 'm'
&& (((isdet & 0xF) == bcMirror) || n1 != ((GPU_PARAM(gcfg, mediaformat) < 100) ? (prop.w) : (gproperty[(mediaid > 0 && GPU_PARAM(gcfg, mediaformat) >= 100) ? 1 : mediaid].w)))) {
&& (((isdet & 0xF) == bcMirror) || n1 != ((GPU_PARAM(gcfg, mediaformat) < 100) ? (prop.w) : (gproperty[(mediaid > 0 && (bool)(GPU_PARAM(gcfg, mediaformat) >= 100)) ? 1 : mediaid].w)))) {
float Rtotal = 1.f;
float cphi, sphi, stheta, ctheta;

Expand Down
8 changes: 4 additions & 4 deletions src/mcx_host.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -757,15 +757,15 @@ void mcx_run_simulation(Config* cfg, float* fluence, float* totalenergy) {
}

if (cfg->optlevel >= 2) {
sprintf(opt + strlen(opt), "%s ", "-DUSE_MACRO_CONST");
sprintf(opt + strlen(opt), "%s ", "-DMCX_SIMPLIFY_BRANCH -DMCX_VECTOR_INDEX");
}

if (cfg->optlevel >= 3) {
sprintf(opt + strlen(opt), "%s ", "-DMCX_SIMPLIFY_BRANCH -DMCX_VECTOR_INDEX");
sprintf(opt + strlen(opt), "%s ", "-DGROUP_LOAD_BALANCE");
}

if (cfg->optlevel >= 4) {
sprintf(opt + strlen(opt), "%s ", "-DGROUP_LOAD_BALANCE");
if (cfg->optlevel >= 4 && cfg->seed != SEED_FROM_FILE) { // optlevel 2 - i.e. defining parameters as macros - fails replay test, need to debug
sprintf(opt + strlen(opt), "%s ", "-DUSE_MACRO_CONST");
}

for (i = 0; i < 3; i++)
Expand Down
8 changes: 4 additions & 4 deletions src/mcx_utils.c
Original file line number Diff line number Diff line change
Expand Up @@ -4488,9 +4488,9 @@ void mcx_printheader(Config* cfg) {
==============================================================================\n\
= Monte Carlo eXtreme (MCX) -- OpenCL =\n\
= Copyright (c) 2010-2023 Qianqian Fang <q.fang at neu.edu> =\n\
= https://mcx.space/ & https://neurojson.org/ =\n\
=" S_CYAN " https://mcx.space/ & https://neurojson.org/ " S_BLUE "=\n\
= =\n\
= Computational Optics&Translational Imaging (COTI) Lab - http://fanglab.org =\n\
= Computational Optics&Translational Imaging (COTI) Lab - " S_CYAN "http://fanglab.org " S_BLUE "=\n\
= Department of Bioengineering, Northeastern University, Boston, MA, USA =\n\
==============================================================================\n\
= The MCX Project is funded by the NIH/NIGMS under grant R01-GM114365 =\n\
Expand All @@ -4499,7 +4499,7 @@ void mcx_printheader(Config* cfg) {
= MCX proudly developed human-readable JSON-based data formats for easy reuse=\n\
= Please consider using JSON (https://neurojson.org/) for your research data =\n\
==============================================================================\n\
$Rev::4fdc45$ " MCX_VERSION " $Date::2018-03-29 00:35:53 -04$by $Author::Qianqian Fang$\n\
$Rev::4fdc45$ " S_CYAN MCX_VERSION S_BLUE " $Date::2018-03-29 00:35:53 -04$by $Author::Qianqian Fang$\n\
==============================================================================\n"S_RESET);
}

Expand Down Expand Up @@ -4564,7 +4564,7 @@ where possible parameters include (the first value in [*|*] is the default)\n\
-G '1101' (--gpu) using multiple devices (1 enable, 0 disable)\n\
-W '50,30,20' (--workload) workload for active devices; normalized by sum\n\
-I (--printgpu) print GPU information and run program\n\
-o [1|int] (--optlevel) optimization level 0-no opt;1,2,3 more optimized\n\
-o [1|int] (--optlevel) optimization level 0-no opt;1-4 more optimized\n\
-J '-DMACRO' (--compileropt) specify additional JIT compiler options\n\
A few built-in preprocessors include\n\
-DMCX_GPU_DEBUG - print step-by-step debug info\n\
Expand Down

0 comments on commit e930448

Please sign in to comment.