From e930448f34677071a0ec406b09002a589c51be3b Mon Sep 17 00:00:00 2001 From: Qianqian Fang Date: Sat, 7 Oct 2023 15:37:30 -0400 Subject: [PATCH] remove opencl jit warnings, adjust optlevel order --- src/mcx_core.cl | 12 ++++++------ src/mcx_host.cpp | 8 ++++---- src/mcx_utils.c | 8 ++++---- 3 files changed, 14 insertions(+), 14 deletions(-) diff --git a/src/mcx_core.cl b/src/mcx_core.cl index 503b859..f6ef2b8 100644 --- a/src/mcx_core.cl +++ b/src/mcx_core.cl @@ -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); } @@ -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; @@ -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) { @@ -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)); @@ -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; diff --git a/src/mcx_host.cpp b/src/mcx_host.cpp index 47bc1ef..1c2655d 100644 --- a/src/mcx_host.cpp +++ b/src/mcx_host.cpp @@ -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++) diff --git a/src/mcx_utils.c b/src/mcx_utils.c index 2ba3165..1872acb 100644 --- a/src/mcx_utils.c +++ b/src/mcx_utils.c @@ -4488,9 +4488,9 @@ void mcx_printheader(Config* cfg) { ==============================================================================\n\ = Monte Carlo eXtreme (MCX) -- OpenCL =\n\ = Copyright (c) 2010-2023 Qianqian Fang =\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\ @@ -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); } @@ -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\