From d431eba01db09ed403d2e5a41d7a5987d2ec9554 Mon Sep 17 00:00:00 2001 From: lungd Date: Sun, 27 Jan 2019 22:14:34 +0100 Subject: [PATCH 1/5] Add skip_display params --- src/main.cpp | 12 ++ src/owWorldSimulation.cpp | 277 ++++++++++++++++++++------------------ 2 files changed, 160 insertions(+), 129 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index c5f7071b..0618fdce 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -39,6 +39,9 @@ bool load_from_file = false; bool load_to = false; +bool skip_display_particles = false; +bool skip_display_membranes = false; +bool skip_display_connections = false; std::string version = "0.0.6"; int usage() { @@ -128,6 +131,15 @@ int main(int argc, char **argv) { graph = true; load_from_file = true; } + if (std::string("-skip_display_particles").compare(argv[i]) == 0) { + skip_display_particles = true; + } + if (std::string("-skip_display_membranes").compare(argv[i]) == 0) { + skip_display_membranes = true; + } + if (std::string("-skip_display_connections").compare(argv[i]) == 0) { + skip_display_connections = true; + } if (std::string("-test").compare(argv[i]) == 0) { // run tests run_tests = true; } diff --git a/src/owWorldSimulation.cpp b/src/owWorldSimulation.cpp index d236ce58..3d0ae650 100644 --- a/src/owWorldSimulation.cpp +++ b/src/owWorldSimulation.cpp @@ -40,6 +40,9 @@ extern bool load_from_file; extern bool load_to; +extern bool skip_display_particles; +extern bool skip_display_membranes; +extern bool skip_display_connections; int old_x = 0, old_y = 0; // Used for mouse event float camera_trans[] = {0, 0, -8.0}; @@ -203,69 +206,71 @@ void display(void) { glBegin(GL_POINTS); float dc, rho; // Display all particles - for (i = 0; i < localConfig->getParticleCount(); ++i) { - if (!load_from_file) { - rho = d_cpp[p_indexb[i * 2 + 0]]; - if (rho < 0) - rho = 0; - if (rho > 2 * localConfig->getConst("rho0")) - rho = 2 * localConfig->getConst("rho0"); - dc = 100.f * (rho - localConfig->getConst("rho0")) / - localConfig->getConst("rho0"); - if (dc > 1.f) - dc = 1.f; - // R G B - glColor4f(0, 0, 1, 1.0f); // blue - if (!load_from_file) { - if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.00f) / - localConfig->getConst("rho0")) > 0) - glColor4f(0, dc, 1, 1.0f); // cyan - if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.01f) / - localConfig->getConst("rho0")) > 0) - glColor4f(0, 1, 1 - dc, 1.0f); // green - if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.02f) / - localConfig->getConst("rho0")) > 0) - glColor4f(dc, 1, 0, 1.0f); // yellow - if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.03f) / - localConfig->getConst("rho0")) > 0) - glColor4f(1, 1 - dc, 0, 1.0f); // red - if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.04f) / - localConfig->getConst("rho0")) > 0) - glColor4f(1, 0, 0, 1.0f); - } - } else - glColor4f(0, 0, 1, 1.0f); // blue - if ((int)p_cpp[i * 4 + 3] != - BOUNDARY_PARTICLE /*&& (int)p_cpp[i*4 + 3] != ELASTIC_PARTICLE*/) { - glBegin(GL_POINTS); - if ((int)p_cpp[i * 4 + 3] == 2) { - glColor4f(0, 0, 0, 1.0f); // color of elastic particles - glPointSize(1.3f * sqrt(sc / 0.025f)); - } - glVertex3f((p_cpp[i * 4] - localConfig->xmax / 2) * sc, - (p_cpp[i * 4 + 1] - localConfig->ymax / 2) * sc, - (p_cpp[i * 4 + 2] - localConfig->zmax / 2) * sc); - glPointSize(1.3f * sqrt(sc / 0.025f)); - glEnd(); + if (!skip_display_particles) { + for (i = 0; i < localConfig->getParticleCount(); ++i) { + if (!load_from_file) { + rho = d_cpp[p_indexb[i * 2 + 0]]; + if (rho < 0) + rho = 0; + if (rho > 2 * localConfig->getConst("rho0")) + rho = 2 * localConfig->getConst("rho0"); + dc = 100.f * (rho - localConfig->getConst("rho0")) / + localConfig->getConst("rho0"); + if (dc > 1.f) + dc = 1.f; + // R G B + glColor4f(0, 0, 1, 1.0f); // blue + if (!load_from_file) { + if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.00f) / + localConfig->getConst("rho0")) > 0) + glColor4f(0, dc, 1, 1.0f); // cyan + if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.01f) / + localConfig->getConst("rho0")) > 0) + glColor4f(0, 1, 1 - dc, 1.0f); // green + if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.02f) / + localConfig->getConst("rho0")) > 0) + glColor4f(dc, 1, 0, 1.0f); // yellow + if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.03f) / + localConfig->getConst("rho0")) > 0) + glColor4f(1, 1 - dc, 0, 1.0f); // red + if ((dc = 100 * (rho - localConfig->getConst("rho0") * 1.04f) / + localConfig->getConst("rho0")) > 0) + glColor4f(1, 0, 0, 1.0f); + } + } else + glColor4f(0, 0, 1, 1.0f); // blue + if ((int)p_cpp[i * 4 + 3] != + BOUNDARY_PARTICLE /*&& (int)p_cpp[i*4 + 3] != ELASTIC_PARTICLE*/) { + glBegin(GL_POINTS); + if ((int)p_cpp[i * 4 + 3] == 2) { + glColor4f(0, 0, 0, 1.0f); // color of elastic particles + glPointSize(1.3f * sqrt(sc / 0.025f)); + } + glVertex3f((p_cpp[i * 4] - localConfig->xmax / 2) * sc, + (p_cpp[i * 4 + 1] - localConfig->ymax / 2) * sc, + (p_cpp[i * 4 + 2] - localConfig->zmax / 2) * sc); + glPointSize(1.3f * sqrt(sc / 0.025f)); + glEnd(); - if (!((p_cpp[i * 4] >= 0) && (p_cpp[i * 4] <= localConfig->xmax) && - (p_cpp[i * 4 + 1] >= 0) && - (p_cpp[i * 4 + 1] <= localConfig->ymax) && - (p_cpp[i * 4 + 2] >= 0) && - (p_cpp[i * 4 + 2] <= localConfig->zmax))) { - /*char label[50]; - beginWinCoords(); - glRasterPos2f (0.01F, 0.05F); - if(err_coord_cnt<50){ - sprintf(label,"%d: %f , %f , %f",i,p_cpp[i*4 - ],p_cpp[i*4+1],p_cpp[i*4+2]); - glPrint( 0.f, (float)(50+err_coord_cnt*11), label, m_font);} - if(err_coord_cnt==50) { - glPrint( 0, (float)(50+err_coord_cnt*11), "............", m_font);} - err_coord_cnt++; - endWinCoords();*/ + if (!((p_cpp[i * 4] >= 0) && (p_cpp[i * 4] <= localConfig->xmax) && + (p_cpp[i * 4 + 1] >= 0) && + (p_cpp[i * 4 + 1] <= localConfig->ymax) && + (p_cpp[i * 4 + 2] >= 0) && + (p_cpp[i * 4 + 2] <= localConfig->zmax))) { + /*char label[50]; + beginWinCoords(); + glRasterPos2f (0.01F, 0.05F); + if(err_coord_cnt<50){ + sprintf(label,"%d: %f , %f , %f",i,p_cpp[i*4 + ],p_cpp[i*4+1],p_cpp[i*4+2]); + glPrint( 0.f, (float)(50+err_coord_cnt*11), label, m_font);} + if(err_coord_cnt==50) { + glPrint( 0, (float)(50+err_coord_cnt*11), "............", m_font);} + err_coord_cnt++; + endWinCoords();*/ + } + } } - } } glLineWidth((GLfloat)0.1); // Display elastic connections @@ -277,6 +282,16 @@ void display(void) { // (generateInitialConfiguration!=1)*numOfBoundaryP; if (i < j) { glColor4b(150 / 2, 125 / 2, 0, 100 / 2 /*alpha*/); + + if (p_cpp[j * 4 + 3] > 2.25) { + //glColor4b(250 / 2, 250 / 2, 200 / 2, 255 / 2); // agar + continue; + } else if ((p_cpp[i * 4 + 3] > 2.31) && (p_cpp[i * 4 + 3] < 2.33)) { + // agar particles which contacted the worm + continue; + } + + if (ec_cpp[4 * i_ec + 2] > 1.f) // muscles { glLineWidth((GLfloat)1.0); @@ -368,8 +383,10 @@ void display(void) { glColor4b(150 / 2, 125 / 2, 0, 100 / 2); if (p_cpp[i * 4 + 3] > 2.15) glColor4b(50 / 2, 125 / 2, 0, 255 / 2); - if (p_cpp[j * 4 + 3] > 2.25) + if (p_cpp[j * 4 + 3] > 2.25) { glColor4b(250 / 2, 250 / 2, 200 / 2, 255 / 2); // agar + continue; + } if ((p_cpp[i * 4 + 3] > 2.31) && (p_cpp[i * 4 + 3] < 2.33)) glColor4b(200 / 2, 250 / 2, 000 / 2, 255 / 2); // agar particles which contacted the worm @@ -393,76 +410,78 @@ void display(void) { } } // Draw membranes - glColor4b(0, 200 / 2, 150 / 2, 255 / 2 /*alpha*/); - for (unsigned int i_m = 0; i_m < localConfig->numOfMembranes; ++i_m) { - i = md_cpp[i_m * 3 + 0]; - j = md_cpp[i_m * 3 + 1]; - k = md_cpp[i_m * 3 + 2]; + if (!skip_display_membranes) { + glColor4b(0, 200 / 2, 150 / 2, 255 / 2 /*alpha*/); + for (unsigned int i_m = 0; i_m < localConfig->numOfMembranes; ++i_m) { + i = md_cpp[i_m * 3 + 0]; + j = md_cpp[i_m * 3 + 1]; + k = md_cpp[i_m * 3 + 2]; - glBegin(GL_LINES); - glVertex3f( - ((p_cpp[i * 4] + p_cpp[j * 4] + 4 * p_cpp[k * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[i * 4 + 1] + p_cpp[j * 4 + 1] + 4 * p_cpp[k * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[i * 4 + 2] + p_cpp[j * 4 + 2] + 4 * p_cpp[k * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); - glVertex3f( - ((p_cpp[i * 4] + p_cpp[k * 4] + 4 * p_cpp[j * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[i * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[j * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[i * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[j * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); + glBegin(GL_LINES); + glVertex3f( + ((p_cpp[i * 4] + p_cpp[j * 4] + 4 * p_cpp[k * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[i * 4 + 1] + p_cpp[j * 4 + 1] + 4 * p_cpp[k * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[i * 4 + 2] + p_cpp[j * 4 + 2] + 4 * p_cpp[k * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); + glVertex3f( + ((p_cpp[i * 4] + p_cpp[k * 4] + 4 * p_cpp[j * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[i * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[j * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[i * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[j * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); - glVertex3f( - ((p_cpp[i * 4] + p_cpp[k * 4] + 4 * p_cpp[j * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[i * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[j * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[i * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[j * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); - glVertex3f( - ((p_cpp[j * 4] + p_cpp[k * 4] + 4 * p_cpp[i * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[j * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[i * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[j * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[i * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); + glVertex3f( + ((p_cpp[i * 4] + p_cpp[k * 4] + 4 * p_cpp[j * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[i * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[j * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[i * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[j * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); + glVertex3f( + ((p_cpp[j * 4] + p_cpp[k * 4] + 4 * p_cpp[i * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[j * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[i * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[j * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[i * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); - glVertex3f( - ((p_cpp[j * 4] + p_cpp[k * 4] + 4 * p_cpp[i * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[j * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[i * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[j * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[i * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); - glVertex3f( - ((p_cpp[i * 4] + p_cpp[j * 4] + 4 * p_cpp[k * 4]) / 6 - - localConfig->xmax / 2) * - sc, - ((p_cpp[i * 4 + 1] + p_cpp[j * 4 + 1] + 4 * p_cpp[k * 4 + 1]) / 6 - - localConfig->ymax / 2) * - sc, - ((p_cpp[i * 4 + 2] + p_cpp[j * 4 + 2] + 4 * p_cpp[k * 4 + 2]) / 6 - - localConfig->zmax / 2) * - sc); - glEnd(); + glVertex3f( + ((p_cpp[j * 4] + p_cpp[k * 4] + 4 * p_cpp[i * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[j * 4 + 1] + p_cpp[k * 4 + 1] + 4 * p_cpp[i * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[j * 4 + 2] + p_cpp[k * 4 + 2] + 4 * p_cpp[i * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); + glVertex3f( + ((p_cpp[i * 4] + p_cpp[j * 4] + 4 * p_cpp[k * 4]) / 6 - + localConfig->xmax / 2) * + sc, + ((p_cpp[i * 4 + 1] + p_cpp[j * 4 + 1] + 4 * p_cpp[k * 4 + 1]) / 6 - + localConfig->ymax / 2) * + sc, + ((p_cpp[i * 4 + 2] + p_cpp[j * 4 + 2] + 4 * p_cpp[k * 4 + 2]) / 6 - + localConfig->zmax / 2) * + sc); + glEnd(); + } } glLineWidth((GLfloat)1.0); glutSwapBuffers(); From 62d4cac4b589bc7c715046f465cfb7b2f5aa2843 Mon Sep 17 00:00:00 2001 From: lungd Date: Mon, 28 Jan 2019 15:42:56 +0100 Subject: [PATCH 2/5] Fix skip_display_connections --- src/owWorldSimulation.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/src/owWorldSimulation.cpp b/src/owWorldSimulation.cpp index 3d0ae650..676cee2c 100644 --- a/src/owWorldSimulation.cpp +++ b/src/owWorldSimulation.cpp @@ -283,10 +283,10 @@ void display(void) { if (i < j) { glColor4b(150 / 2, 125 / 2, 0, 100 / 2 /*alpha*/); - if (p_cpp[j * 4 + 3] > 2.25) { + if (skip_display_connections && (p_cpp[j * 4 + 3] > 2.25)) { //glColor4b(250 / 2, 250 / 2, 200 / 2, 255 / 2); // agar continue; - } else if ((p_cpp[i * 4 + 3] > 2.31) && (p_cpp[i * 4 + 3] < 2.33)) { + } else if (skip_display_connections && (p_cpp[i * 4 + 3] > 2.31) && (p_cpp[i * 4 + 3] < 2.33)) { // agar particles which contacted the worm continue; } @@ -385,7 +385,6 @@ void display(void) { glColor4b(50 / 2, 125 / 2, 0, 255 / 2); if (p_cpp[j * 4 + 3] > 2.25) { glColor4b(250 / 2, 250 / 2, 200 / 2, 255 / 2); // agar - continue; } if ((p_cpp[i * 4 + 3] > 2.31) && (p_cpp[i * 4 + 3] < 2.33)) glColor4b(200 / 2, 250 / 2, 000 / 2, From 193793433990eed447a9bf89ecfc5f324e8a6e63 Mon Sep 17 00:00:00 2001 From: lungd Date: Mon, 28 Jan 2019 15:47:45 +0100 Subject: [PATCH 3/5] Pass dt to C302NRNSimulation instance after initialization --- inc/owSignalSimulator.h | 2 +- main_sim.py | 23 ++++++++++++++++++----- src/owConfigProperty.cpp | 4 ++-- src/owSignalSimulator.cpp | 11 ++++++++++- 4 files changed, 31 insertions(+), 9 deletions(-) diff --git a/inc/owSignalSimulator.h b/inc/owSignalSimulator.h index f341eba7..b702014a 100644 --- a/inc/owSignalSimulator.h +++ b/inc/owSignalSimulator.h @@ -53,7 +53,7 @@ class SignalSimulator: public owINeuronSimulator{ public: - SignalSimulator(const std::string & simFileName = "main_sim", const std::string & simClassName = "MuscleSimulation"); + SignalSimulator(const std::string & simFileName = "main_sim", const std::string & simClassName = "MuscleSimulation", float timeStep=0.005); std::vector run(); ~SignalSimulator(); }; diff --git a/main_sim.py b/main_sim.py index 64a79b03..1bcde0fd 100644 --- a/main_sim.py +++ b/main_sim.py @@ -120,6 +120,9 @@ def __init__(self,increment=1.0): self.increment = increment self.step = 0 + def set_timestep(self, dt): + pass + def run(self, skip_to_time=0, do_plot = True): self.contraction_array = parallel_waves(step = self.step) self.step += self.increment @@ -147,15 +150,25 @@ class C302NRNSimulation(): def __init__(self, tstop=100, dt=0.005, activity_file=None, verbose=True): #from LEMS_c302_C1_Full_nrn import NeuronSimulation - from LEMS_c302_nrn import NeuronSimulation + #from LEMS_c302_nrn import NeuronSimulation - import neuron - self.h = neuron.h + #import neuron + #self.h = neuron.h + self.tstop = tstop self.verbose = verbose - self.ns = NeuronSimulation(tstop, dt) - print_("Initialised C302NRNSimulation of length %s ms and dt = %s ms..."%(tstop,dt)) + #self.ns = NeuronSimulation(tstop, dt) + #print_("Initialised C302NRNSimulation of length %s ms and dt = %s ms..."%(tstop,dt)) + + def set_timestep(self, dt): + dt = float('{:0.1e}'.format(dt)) * 1000.0 # memory issue fix + from LEMS_c302_nrn import NeuronSimulation + import neuron + self.h = neuron.h + + self.ns = NeuronSimulation(self.tstop, dt) + print_("Initialised C302NRNSimulation of length %s ms and dt = %s ms..."%(self.tstop,dt)) def save_results(self): diff --git a/src/owConfigProperty.cpp b/src/owConfigProperty.cpp index 41283fb1..d2d9e3bf 100755 --- a/src/owConfigProperty.cpp +++ b/src/owConfigProperty.cpp @@ -147,9 +147,9 @@ owConfigProperty::owConfigProperty(int argc, char **argv) } if (simName.compare("") == 0) - simulation = new SignalSimulator("main_sim", pythonClass); + simulation = new SignalSimulator("main_sim", pythonClass, this->timeStep); else - simulation = new SignalSimulator(simName, pythonClass); + simulation = new SignalSimulator(simName, pythonClass, this->timeStep); } else { simulation = new owNeuronSimulator(1, this->timeStep, nrnSimulationFileName); diff --git a/src/owSignalSimulator.cpp b/src/owSignalSimulator.cpp index e0014841..d76529b0 100644 --- a/src/owSignalSimulator.cpp +++ b/src/owSignalSimulator.cpp @@ -46,7 +46,8 @@ #include "owSignalSimulator.h" SignalSimulator::SignalSimulator(const std::string &simFileName, - const std::string &simClassName) { + const std::string &simClassName, + float timeStep) { // char pyClass[] = "SiberneticNEURONWrapper"; // char pyClass[] = "C302Simulation"; @@ -81,6 +82,14 @@ SignalSimulator::SignalSimulator(const std::string &simFileName, pInstance = PyObject_CallObject(pClass, nullptr); if (PyErr_Occurred()) PyErr_Print(); + PyObject *dt = Py_BuildValue("f", timeStep); // Create tuple of arguments for initialization + PyObject *pFuncName = Py_BuildValue("s", "set_timestep"); + //pInstance = PyObject_CallMethod(pInstance, "set_timestep", "(f)", timeStep); + PyObject_CallMethodObjArgs(pInstance, pFuncName, dt, nullptr); + if (PyErr_Occurred()) + PyErr_Print(); + Py_DECREF(dt); + Py_DECREF(pFuncName); std::cout << "Python muscle signal generator class: " << simClassName << " loaded!" << std::endl; } else { From c3152486460bf7376493dc88af22a177a6c48520 Mon Sep 17 00:00:00 2001 From: lungd Date: Mon, 11 Feb 2019 12:55:15 +0100 Subject: [PATCH 4/5] Set local_NDRange_size to 32 (MAX_NEIGHBOR_COUNT). Better performance with nvidia 940m. unchanged performance with intel cpu/gpu --- inc/owOpenCLConstant.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/inc/owOpenCLConstant.h b/inc/owOpenCLConstant.h index 3cf2bb89..ba0bf4f6 100644 --- a/inc/owOpenCLConstant.h +++ b/inc/owOpenCLConstant.h @@ -50,7 +50,7 @@ #define INTEL_OPENCL_DEBUG 0 -const int local_NDRange_size = 256; +const int local_NDRange_size = MAX_NEIGHBOR_COUNT; enum DEVICE { CPU = 0, GPU = 1, ALL = 2}; enum INTEGRATOR { EULER = 0, LEAPFROG = 1 }; From 0ada241505fe72ceced2a099aa8ac6de67a148e3 Mon Sep 17 00:00:00 2001 From: lungd Date: Mon, 11 Feb 2019 12:56:44 +0100 Subject: [PATCH 5/5] Test --- src/owOpenCLSolver.cpp | 5 +++-- src/sphFluid.cl | 46 ++++++++++++++++++++++++++++-------------- 2 files changed, 34 insertions(+), 17 deletions(-) diff --git a/src/owOpenCLSolver.cpp b/src/owOpenCLSolver.cpp index 4a468ce9..b599f284 100644 --- a/src/owOpenCLSolver.cpp +++ b/src/owOpenCLSolver.cpp @@ -37,6 +37,7 @@ #include #include "owOpenCLSolver.h" +#include "owOpenCLConstant.h" int comparator(const void *v1, const void *v2); @@ -731,11 +732,11 @@ owOpenCLSolver::_run_pcisph_computeDensity(owConfigProperty *config) { pcisph_computeDensity.setArg(5, config->getParticleCount()); int err = queue.enqueueNDRangeKernel( pcisph_computeDensity, cl::NullRange, - cl::NDRange((int)(config->getParticleCount_RoundUp())), + cl::NDRange((int)(config->getParticleCount_RoundUp()*local_NDRange_size/2)), #if defined(__APPLE__) cl::NullRange, nullptr, nullptr); #else - cl::NDRange((int)(local_NDRange_size)), nullptr, nullptr); + cl::NDRange((int)(local_NDRange_size/2)), nullptr, nullptr); #endif #if QUEUE_EACH_KERNEL queue.finish(); diff --git a/src/sphFluid.cl b/src/sphFluid.cl index a7b4f942..75d72db4 100644 --- a/src/sphFluid.cl +++ b/src/sphFluid.cl @@ -478,7 +478,7 @@ __kernel void pcisph_computeDensity( __global uint * particleIndexBack, uint PARTICLE_COUNT ) { - int id = get_global_id( 0 ); + int id = get_group_id( 0 ); if( id >= PARTICLE_COUNT ) return; id = particleIndexBack[id]; //track selected particle (indices are not shuffled anymore) int idx = id * MAX_NEIGHBOR_COUNT; @@ -488,25 +488,41 @@ __kernel void pcisph_computeDensity( float hScaled6 = hScaled2*hScaled2*hScaled2; int real_nc = 0; - do // gather density contribution from all neighbors (if they exist) - { - if( NEIGHBOR_MAP_ID( neighborMap[ idx + nc ] ) != NO_PARTICLE_ID ) - { + int lid = get_local_id(0); + __local float l_density[MAX_NEIGHBOR_COUNT/2]; + + l_density[lid] = 0.0; + barrier(CLK_LOCAL_MEM_FENCE); + + nc = lid; + int offset = MAX_NEIGHBOR_COUNT/2; + + if( NEIGHBOR_MAP_ID( neighborMap[ idx + nc ] ) != NO_PARTICLE_ID ) { r_ij2= NEIGHBOR_MAP_DISTANCE( neighborMap[ idx + nc ] ); // distance is already scaled here r_ij2 *= r_ij2; - if(r_ij2hScaled2) printf("=Error: r_ij/h = %f\n", NEIGHBOR_MAP_DISTANCE( neighborMap[ idx + nc ] ) / hScaled); - real_nc++; + if(r_ij2