int main() {
Kernel *ker; /* Kernel */
Module *world; /* World manager module */
/* Initialize kernel */
cout << "Begin execution" << endl;
ker = new Kernel();
ker->initialize();
/* Declare modules */
world = new WorldManager("WorldManager");
/* Register modules */
ker->register_module(world);
/* Initialize all modules */
ker->initialize_modules();
/* Load all modules */
ker->load_modules();
/* Configure simulation */
ker->configure_simulation(Kernel::FASTSIM, 20.0);
//SIMULATE
ker->simulate();
/* Unload modules */
ker->unload_modules();
cout << "Freeing last resources... ";
delete(world);
delete(ker);
cout << "DONE!" << endl;
return 0;
}
void run( int workgroup_size )
{
NDRange global,local;
Event event;
local = NDRange(workgroup_size);
global = NDRange(_nr_groups*workgroup_size);
posix_time::ptime t1 = posix_time::microsec_clock::local_time();
_kernel.setArg(0,_output);
_queue.enqueueNDRangeKernel( _kernel, global, local, &event );
_queue.flush();
_queue.waitForEvent(event);
posix_time::ptime t2 = posix_time::microsec_clock::local_time();
_exec_time = posix_time::time_period(t1,t2).length().total_microseconds();
}
double l1Distance(const Kernel &rhs, long nr = 1000)
{
double a,b,x;
a = min_ - 4.0*width_;
b = rhs.min_ - 4.0*rhs.width_;
double dm = (a < b ? a : b);
a = max_ + 4.0*width_;
b = rhs.max_ + 4.0*rhs.width_;
double dM = (a > b ? a : b);
double dD=(dM-dm)/(nr-1);
const function::Linear<double> lt(0, dm, nr-1, dM);
dM = 0.0;
for (long i=0; i<nr; ++i)
{
x = lt(i);
dM += fabs(density(x)-rhs.density(x))*dD;
}
return dM;
}
void setup( int dev, int workgroup_size )
{
std::vector<Device> devices = _context.getInfo<CAL_CONTEXT_DEVICES>();
// create program
std::string source = create_kernel_peekperf(workgroup_size);
//std::cout << source; // Uncomment to emit IL code
_program = Program( _context, source.c_str(), source.length() );
_program.build(devices);
//_program.disassemble(std::cout); // Uncomment to emit ISA code
// create kernel
_kernel = Kernel(_program,"main");
_kernel.setArgBind(0,"g[]");
_nr_groups = devices[dev].getInfo<CAL_DEVICE_NUMBEROFSIMD>();
_queue = CommandQueue(_context,devices[dev]);
// create output buffer
int width = 64*((workgroup_size + 63)/64);
_output = Image2D(_context, width, _nr_groups, CAL_FORMAT_FLOAT_4, CAL_RESALLOC_GLOBAL_BUFFER );
}
virtual bool main()
{
if(mainDisplay == 0)
return false;
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
glLoadIdentity();
glTranslatef(-1.5f, 0.0f, -6.0f);
glBegin(GL_TRIANGLES);
glVertex3f(0.0f, 1.0f, 0.0f);
glVertex3f(-1.0f, -1.0f, 0.0f);
glVertex3f(1.0f, -1.0f, 0.0f);
glEnd();
glTranslatef(3.0f, 0.0f, 0.0f);
glBegin(GL_QUADS);
glVertex3f(-1.0f, 1.0f, 0.0f);
glVertex3f(1.0f, 1.0f, 0.0f);
glVertex3f(1.0f, -1.0f, 0.0f);
glVertex3f(-1.0f, -1.0f, 0.0f);
glEnd();
kernel->refreshDisplay();
return true;
}
/*! \brief Apply the kernel expression to a particle
*
* \param vd vector with the particles positions
* \param cl Cell-list
* \param v_exp expression to execute for each particle
* \param key to which particle to apply the expression
* \param lker kernel function
*
* \return the result of apply the kernel on the particle key
*
*/
inline static typename std::remove_reference<rtype>::type apply(const vector & vd, NN_type & cl, const exp & v_exp, const vect_dist_key_dx & key, Kernel & lker)
{
// accumulator
typename std::remove_reference<rtype>::type pse = set_zero<typename std::remove_reference<rtype>::type>::create();
// position of particle p
Point<vector::dims,typename vector::stype> p = vd.getPos(key);
// property of the particle x
rtype prp_p = v_exp.value(key);
// Get the neighborhood of the particle
auto NN = cl.template getNNIterator<NO_CHECK>(cl.getCell(p));
while(NN.isNext())
{
auto nnp = NN.get();
// Calculate contribution given by the kernel value at position p,
// given by the Near particle, exclude itself
if (nnp != key.getKey())
{
// property of the particle x
rtype prp_q = v_exp.value(nnp);
// position of the particle q
Point<vector::dims,typename vector::stype> q = vd.getPos(nnp);
pse += lker.value(p,q,prp_p,prp_q);
}
// Next particle
++NN;
}
return pse;
}
int main(int argc, char* argv[])
{
// Allocate the kernel object and the application object
Kernel* kernel = Kernel::instance(); // Get the kernel
simpleApp* application = new simpleApp(); // Instantiate an instance of the app
#if ! defined(VRJ_USE_COCOA)
// IF not args passed to the program
// Display usage information and exit
if (argc <= 1)
{
std::cout << "\n\n";
std::cout << "Usage: " << argv[0] << " vjconfigfile[0] vjconfigfile[1] ... vjconfigfile[n]" << std::endl;
std::exit(EXIT_FAILURE);
}
#endif
kernel->init(argc, argv);
// Load any config files specified on the command line
for( int i = 1; i < argc; ++i )
{
kernel->loadConfigFile(argv[i]);
}
// Start the kernel running
kernel->start();
// Give the kernel an application to execute
kernel->setApplication(application);
// Keep thread alive and waiting
kernel->waitForKernelStop();
delete application;
return EXIT_SUCCESS;
}
bool searchModel_minimalSamples(const Kernel &kernel,
typename Kernel::Model* bestModel,
InliersVec *bestInliers = 0,
double *bestRMS = 0)
{
assert(kernel.NumSamples() == Kernel::MinimumSamples());
InliersVec isInlier(kernel.NumSamples());
int best_score = 0;
bool bestModelFound = false;
std::vector<typename Kernel::Model> possibleModels;
kernel.ComputeModelFromMinimumSamples(&possibleModels);
for (std::size_t i = 0; i < possibleModels.size(); ++i) {
double rms;
int model_score = kernel.ComputeInliersForModel(possibleModels[i], &isInlier, bestRMS ? &rms : 0);
if (model_score > best_score) {
if (bestRMS) {
*bestRMS = rms;
}
best_score = model_score;
*bestModel = possibleModels[i];
bestModelFound = true;
}
}
if (!bestModelFound) {
return false;
}
if (bestInliers) {
*bestInliers = isInlier;
}
if (bestRMS) {
*bestRMS = kernel.ScalarUnormalize(*bestRMS);
}
kernel.Unnormalize(bestModel);
return true;
}
/** Return Euler angles acording to a convention
* @param convention :: the convention used to calculate Euler Angles. The
* UniversalGoniometer is YZY, a triple axis goniometer at HFIR is YZX
*/
std::vector<double> Goniometer::getEulerAngles(std::string convention) {
return Quat(getR()).getEulerAngles(convention);
}
void two_level_test(const Kernel& K)
{
typedef Kernel kernel_type;
typedef typename kernel_type::point_type point_type;
typedef typename kernel_type::source_type source_type;
typedef typename kernel_type::target_type target_type;
typedef typename kernel_type::charge_type charge_type;
typedef typename kernel_type::result_type result_type;
typedef typename kernel_type::multipole_type multipole_type;
typedef typename kernel_type::local_type local_type;
// init source
std::vector<source_type> s(1);
s[0] = source_type(0,0,0);
// init charge
std::vector<charge_type> c(1);
c[0][0] = 1.;
c[0][1] = 2.;
c[0][2] = 3.;
c[0][3] = 1.;
c[0][4] = 0.;
c[0][5] = 1.;
// init target
std::vector<target_type> t(1);
t[0] = target_type(0.98,0.98,0.98);
// init results vectors for exact, FMM
std::vector<result_type> rexact(1);
result_type rm2p;
result_type rfmm;
// test direct
K.P2P(s.begin(),s.end(),c.begin(),t.begin(),t.end(),rexact.begin());
// setup intial multipole expansion
multipole_type M;
const point_type M_center(0.05, 0.05, 0.05);
INITM::eval(K, M, M_center, 2u);
K.P2M(s[0], c[0], M_center, M);
// perform M2M
multipole_type M2;
const point_type M2_center(0.1, 0.1, 0.1);
INITM::eval(K, M2, M2_center, 1u);
K.M2M(M, M2, M2_center - M_center);
//K.P2M(s,c,M2_center,M2);
// test M2P
K.M2P(M2, M2_center, t[0], rm2p);
// test M2L, L2P
#ifndef TREECODE_ONLY
local_type L2;
point_type L2_center(0.9, 0.9, 0.9);
INITL::eval(K, L2, L2_center, 1u);
K.M2L(M2, L2, L2_center - M2_center);
// test L2L
local_type L;
point_type L_center(0.95, 0.95, 0.95);
INITL::eval(K, L, L_center, 2u);
K.L2L(L2, L, L_center - L2_center);
// test L2P
K.L2P(L2, L2_center, t[0], rfmm);
#endif
// check errors
std::cout << "rexact = " << rexact[0] << std::endl;
std::cout << "rm2p = " << 1./6*rm2p << std::endl;
std::cout << "rfmm = " << 1./6*rfmm << std::endl;
/*
std::cout << "M2P L1 rel error: "
<< std::scientific << l1_rel_error(rm2p, rexact) << std::endl;
std::cout << "M2P L2 error: "
<< std::scientific << l2_error(rm2p, rexact) << std::endl;
std::cout << "M2P L2 rel error: "
<< std::scientific << l2_rel_error(rm2p, rexact) << std::endl;
#ifndef TREECODE_ONLY
std::cout << "FMM L1 rel error: "
<< std::scientific << l1_rel_error(rfmm, rexact) << std::endl;
std::cout << "FMM L2 error: "
<< std::scientific << l2_error(rfmm, rexact) << std::endl;
std::cout << "FMM L2 rel error: "
<< std::scientific << l2_rel_error(rfmm, rexact) << std::endl;
#endif
*/
}
void init() {
// Default pins to low status
for (int i = 0; i < 5; i++){
leds[i].output();
leds[i]= 0;
}
#ifdef STEPTICKER_DEBUG_PIN
stepticker_debug_pin.output();
stepticker_debug_pin= 0;
#endif
Kernel* kernel = new Kernel();
kernel->streams->printf("Smoothie Running @%ldMHz\r\n", SystemCoreClock / 1000000);
Version version;
kernel->streams->printf(" Build version %s, Build date %s\r\n", version.get_build(), version.get_build_date());
//some boards don't have leds.. TOO BAD!
kernel->use_leds= !kernel->config->value( disable_leds_checksum )->by_default(false)->as_bool();
bool sdok= (sd.disk_initialize() == 0);
if(!sdok) kernel->streams->printf("SDCard is disabled\r\n");
#ifdef DISABLEMSD
// attempt to be able to disable msd in config
if(sdok && !kernel->config->value( disable_msd_checksum )->by_default(false)->as_bool()){
// HACK to zero the memory USBMSD uses as it and its objects seem to not initialize properly in the ctor
size_t n= sizeof(USBMSD);
void *v = AHB0.alloc(n);
memset(v, 0, n); // clear the allocated memory
msc= new(v) USBMSD(&u, &sd); // allocate object using zeroed memory
}else{
msc= NULL;
kernel->streams->printf("MSD is disabled\r\n");
}
#endif
// Create and add main modules
kernel->add_module( new SimpleShell() );
kernel->add_module( new Configurator() );
kernel->add_module( new CurrentControl() );
kernel->add_module( new PauseButton() );
kernel->add_module( new PlayLed() );
kernel->add_module( new Endstops() );
kernel->add_module( new Player() );
// these modules can be completely disabled in the Makefile by adding to EXCLUDE_MODULES
#ifndef NO_TOOLS_SWITCH
SwitchPool *sp= new SwitchPool();
sp->load_tools();
delete sp;
#endif
#ifndef NO_TOOLS_EXTRUDER
// NOTE this must be done first before Temperature control so ToolManager can handle Tn before temperaturecontrol module does
ExtruderMaker *em= new ExtruderMaker();
em->load_tools();
delete em;
#endif
#ifndef NO_TOOLS_TEMPERATURECONTROL
// Note order is important here must be after extruder so Tn as a parameter will get executed first
TemperatureControlPool *tp= new TemperatureControlPool();
tp->load_tools();
kernel->temperature_control_pool= tp;
#else
kernel->temperature_control_pool= new TemperatureControlPool(); // so we can get just an empty temperature control array
#endif
#ifndef NO_TOOLS_LASER
kernel->add_module( new Laser() );
#endif
#ifndef NO_TOOLS_SPINDLE
kernel->add_module( new Spindle() );
#endif
#ifndef NO_UTILS_PANEL
kernel->add_module( new Panel() );
#endif
#ifndef NO_TOOLS_TOUCHPROBE
kernel->add_module( new Touchprobe() );
#endif
#ifndef NO_TOOLS_ZPROBE
kernel->add_module( new ZProbe() );
#endif
#ifndef NO_TOOLS_SCARACAL
kernel->add_module( new SCARAcal() );
#endif
#ifndef NONETWORK
kernel->add_module( new Network() );
#endif
#ifndef NO_TOOLS_TEMPERATURESWITCH
// Must be loaded after TemperatureControlPool
kernel->add_module( new TemperatureSwitch() );
#endif
#ifndef NO_TOOLS_DRILLINGCYCLES
kernel->add_module( new Drillingcycles() );
#endif
#ifndef NO_TOOLS_FILAMENTDETECTOR
kernel->add_module( new FilamentDetector() );
#endif
// Create and initialize USB stuff
u.init();
#ifdef DISABLEMSD
if(sdok && msc != NULL){
kernel->add_module( msc );
}
#else
kernel->add_module( &msc );
#endif
kernel->add_module( &usbserial );
if( kernel->config->value( second_usb_serial_enable_checksum )->by_default(false)->as_bool() ){
kernel->add_module( new(AHB0) USBSerial(&u) );
}
if( kernel->config->value( dfu_enable_checksum )->by_default(false)->as_bool() ){
kernel->add_module( new(AHB0) DFU(&u));
}
kernel->add_module( &u );
// clear up the config cache to save some memory
kernel->config->config_cache_clear();
if(kernel->use_leds) {
// set some leds to indicate status... led0 init doe, led1 mainloop running, led2 idle loop running, led3 sdcard ok
leds[0]= 1; // indicate we are done with init
leds[3]= sdok?1:0; // 4th led inidicates sdcard is available (TODO maye should indicate config was found)
}
if(sdok) {
// load config override file if present
// NOTE only Mxxx commands that set values should be put in this file. The file is generated by M500
FILE *fp= fopen(kernel->config_override_filename(), "r");
if(fp != NULL) {
char buf[132];
kernel->streams->printf("Loading config override file: %s...\n", kernel->config_override_filename());
while(fgets(buf, sizeof buf, fp) != NULL) {
kernel->streams->printf(" %s", buf);
if(buf[0] == ';') continue; // skip the comments
struct SerialMessage message= {&(StreamOutput::NullStream), buf};
kernel->call_event(ON_CONSOLE_LINE_RECEIVED, &message);
}
kernel->streams->printf("config override file executed\n");
fclose(fp);
}
}
THEKERNEL->step_ticker->start();
}
//----------------------------------------------------------------------------------------------
///< Render Object or ObjComponent
void BitmapGeometryHandler::Render() {
// std::cout << "BitmapGeometryHandler::Render() called\n";
V3D pos;
// Wait for no error
while (glGetError() != GL_NO_ERROR)
;
// Because texture colours are combined with the geometry colour
// make sure the current colour is white
glColor3f(1.0f, 1.0f, 1.0f);
// Nearest-neighbor scaling
GLint texParam = GL_NEAREST;
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, texParam);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, texParam);
glEnable(GL_TEXTURE_2D); // enable texture mapping
int texx, texy;
m_rectDet->getTextureSize(texx, texy);
double tex_frac_x = (1.0 * m_rectDet->xpixels()) / (texx);
double tex_frac_y = (1.0 * m_rectDet->ypixels()) / (texy);
// Point to the ID of the texture that was created before - in
// RectangularDetectorActor.
// int texture_id = m_rectDet->getTextureID();
// glBindTexture (GL_TEXTURE_2D, texture_id);
// if (glGetError()>0) std::cout << "OpenGL error in glBindTexture \n";
glBegin(GL_QUADS);
glTexCoord2f(0.0, 0.0);
pos = m_rectDet->getRelativePosAtXY(0, 0);
pos += V3D(m_rectDet->xstep() * (-0.5), m_rectDet->ystep() * (-0.5),
0.0); // Adjust to account for the size of a pixel
glVertex3f((GLfloat)pos.X(), (GLfloat)pos.Y(), (GLfloat)pos.Z());
glTexCoord2f((GLfloat)tex_frac_x, 0.0);
pos = m_rectDet->getRelativePosAtXY(m_rectDet->xpixels() - 1, 0);
pos += V3D(m_rectDet->xstep() * (+0.5), m_rectDet->ystep() * (-0.5),
0.0); // Adjust to account for the size of a pixel
glVertex3f((GLfloat)pos.X(), (GLfloat)pos.Y(), (GLfloat)pos.Z());
glTexCoord2f((GLfloat)tex_frac_x, (GLfloat)tex_frac_y);
pos = m_rectDet->getRelativePosAtXY(m_rectDet->xpixels() - 1,
m_rectDet->ypixels() - 1);
pos += V3D(m_rectDet->xstep() * (+0.5), m_rectDet->ystep() * (+0.5),
0.0); // Adjust to account for the size of a pixel
glVertex3f((GLfloat)pos.X(), (GLfloat)pos.Y(), (GLfloat)pos.Z());
glTexCoord2f(0.0, (GLfloat)tex_frac_y);
pos = m_rectDet->getRelativePosAtXY(0, m_rectDet->ypixels() - 1);
pos += V3D(m_rectDet->xstep() * (-0.5), m_rectDet->ystep() * (+0.5),
0.0); // Adjust to account for the size of a pixel
glVertex3f((GLfloat)pos.X(), (GLfloat)pos.Y(), (GLfloat)pos.Z());
glEnd();
if (glGetError() > 0)
std::cout << "OpenGL error in BitmapGeometryHandler::Render \n";
glDisable(
GL_TEXTURE_2D); // stop texture mapping - not sure if this is necessary.
}
int main(void)
{
// Setup scene details
scene.fogDensity = 0.003f;
scene.ambientColour = vec3(0.10f, 0.16f, 0.20f);
scene.lightColour = vec3(0.63f, 0.55f, 0.5f);
scene.lightDirection = normalize(vec3(1.0f, 0.25f, 0.5f));
// Setup raytracer
raytracer = new Raytracer(WIDTH, HEIGHT, scene, voxels, camera);
// Setup SDL
SDL_Init(SDL_INIT_EVERYTHING);
SDL_SetVideoMode(WIDTH, HEIGHT, 32, SDL_OPENGL | (FULLSCREEN ? SDL_FULLSCREEN : 0));
TTF_Init();
// Setup screen
setupOpenGL();
gui = new GUI(ivec2(WIDTH, HEIGHT));
// Get platform and device information
cl_platform_id platform = selectPlatform();
cl_device_id device = selectDevice(platform);
// Create an OpenCL context
cl_int ret;
cl_context_properties properties[] =
{
CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(),
CL_GLX_DISPLAY_KHR, (cl_context_properties)glXGetCurrentDisplay(),
CL_CONTEXT_PLATFORM, (cl_context_properties)platform,
0
};
cl_context context = clCreateContext(properties, 1, &device, openCLLog, NULL, &ret);
printCLErrorCode("createContext", ret);
// Create a command queue
cl_command_queue command_queue = clCreateCommandQueue(context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &ret);
// Create texture
outputTexture = createTexture();
screenTexture[0] = createTexture();
screenTexture[1] = createTexture();
finalTexture = createTexture();
// Load textures
grassTexture = new Texture("textures/grass.png");
// Create pixel array
int bytes = WIDTH * HEIGHT * sizeof(vec4);
vec4 *pixels = (vec4*)malloc(bytes);
glBindTexture(GL_TEXTURE_2D, outputTexture);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F, WIDTH, HEIGHT, 0, GL_RGBA, GL_FLOAT, NULL);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glBindTexture(GL_TEXTURE_2D, 0);
glBindTexture(GL_TEXTURE_2D, screenTexture[0]);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F, WIDTH, HEIGHT, 0, GL_RGBA, GL_FLOAT, NULL);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glBindTexture(GL_TEXTURE_2D, 0);
glBindTexture(GL_TEXTURE_2D, screenTexture[1]);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F, WIDTH, HEIGHT, 0, GL_RGBA, GL_FLOAT, NULL);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glBindTexture(GL_TEXTURE_2D, 0);
glBindTexture(GL_TEXTURE_2D, finalTexture);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F, WIDTH, HEIGHT, 0, GL_RGBA, GL_FLOAT, NULL);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glBindTexture(GL_TEXTURE_2D, 0);
// Create memory buffers on the device for each vector
cl_mem cameraPosMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(vec3), NULL, &ret);
cl_mem cameraViewMatrixMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(mat4), NULL, &ret);
cl_mem lightDirectionMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(vec3), NULL, &ret);
cl_mem voxelSizeMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * 3, NULL, &ret);
printCLErrorCode("clCreateBuffer", ret);
if (ret != 0) return 1;
cl_mem outputMemory = clCreateFromGLTexture2D(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, outputTexture, &ret);
cl_mem screenMemory[2] = {
clCreateFromGLTexture2D(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, screenTexture[0], &ret),
clCreateFromGLTexture2D(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, screenTexture[1], &ret)
};
cl_mem finalMemory = clCreateFromGLTexture2D(context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, finalTexture, &ret);
// Create the OpenCL kernel
Kernel* kernelFast = new Kernel("Fast", "raytracer_fast.cl", "raytrace", context, device);
Kernel* kernelNormal = new Kernel("AO", "raytracer_normal.cl", "raytrace", context, device);
Kernel* kernelSlow = new Kernel("GI", "raytracer.cl", "raytrace", context, device);
Kernel* kernelPath = new Kernel("Path", "raytracer_path.cl", "raytrace", context, device);
vector<Kernel*> raytracerKernels = {
kernelFast,
kernelNormal,
kernelSlow,
kernelPath
};
for (Kernel* kernel : raytracerKernels) kernel->build();
Kernel postprocess("postprocess", "postprocess.cl", "main", context, device);
postprocess.build();
Kernel tonemapper("tonemapper", "tonemap.cl", "main", context, device);
tonemapper.build();
bool reloadKernel = true;
int kernelType = 0;
Kernel* kernel = kernelFast;
// Create voxels
const int AIR = 0, GRASS = 1, GROUND = 2, STONE = 3, WATER = 4, WOOD = 5, LEAVES = 6;
#pragma omp parallel for schedule(dynamic)
for (int z = 0; z < voxelSize.z; z++)
for (int y = 0; y < voxelSize.y; y++)
for (int x = 0; x < voxelSize.x; x++)
{
float h = (perlinNoise(0.6f, 6, x / 128.0f, y / 128.0f, z / 128.0f, 0) - (0.5f - (float)y / voxelSize.y)) * 3.0f;
int v = AIR;
if (h >= 1.2f)
v = STONE;
else if (h >= 1.0f)
v = GROUND;
if (y > 0 && v == GROUND && voxels.get(x, y - 1, z) == AIR)
v = GRASS;
if (v == 0 && y > voxelSize.y - seaLevel)
v = WATER;
voxels.get(x, y, z) = v;
}
int trees = voxelSize.x*voxelSize.z*0.01f;
int plantedTrees = 0;
for (int i = 0; i < trees*trees && plantedTrees < trees; i++)
{
int x = noise(i, 10) * (voxelSize.x - 8) + 4;
int y = 0;
int z = noise(i, 20) * (voxelSize.z - 8) + 4;
bool valid = true;
for (y = 0; y < voxelSize.y; y++)
{
int type = voxels.get(x, y, z);
valid = (type == GRASS);
if (type != AIR) break;
}
if (!valid || y < 10) continue;
plantedTrees++;
for (int yt = y - 6; yt < y; yt++)
voxels.get(x, yt, z) = WOOD;
for (int yt = y - 8; yt < y - 2; yt++)
{
for (int xt = -4; xt <= 4; xt++)
for (int zt = -4; zt <= 4; zt++)
{
int yft = y - yt - 6;
int dist = zt*zt + xt*xt + yft*yft + perlinNoise(0.5, 3, (x+xt)/2.0f, yt, (z+zt)/2.0f, 13) * 8.0f;
if (dist <= 12)
voxels.get(x+xt, yt, z+zt) = LEAVES;
}
}
}
voxels.generateOctree();
voxelsDirty = false;
//voxels.remake();
cout << "Node: " << sizeof(Node) << endl;
// Create OpenCL data buffers
int voxelsSpace = voxels.space * 1.2;
int voxelsPtrSpace = voxels.ptrSpace * 1.5;
size_t gpuMemoryEstimate = sizeof(uint8_t) * voxelCount + sizeof(Node) * voxelsSpace + sizeof(uint32_t) * voxelsPtrSpace;
cout << "GPU Memory Estimate: " << ((gpuMemoryEstimate / 1024.0) / 1024) << "MB" << endl;
cl_mem voxelMemory, octreeMemory, ptrTableMemory;
voxelMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(uint8_t) * voxelCount, NULL, &ret);
octreeMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(Node) * voxelsSpace, NULL, &ret);
ptrTableMemory = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(uint32_t) * voxelsPtrSpace, NULL, &ret);
// Main loop
SDL_Event event;
float delta = -1.0f;
long last = SDL_GetTicks();
long lastTimePhysics = last;
long nextReport = last;
bool running = true;
bool postprocessEnabled = true;
int framesLeftToRender = -1;
int forceMixOff = 2;
int succesiveFrames = 1;
while (running)
{
bool onceBomb = false;
while (SDL_PollEvent(&event))
switch (event.type)
{
case SDL_QUIT:
running = false;
break;
case SDL_KEYUP:
if (!grab) gui->keyPress(event.key.keysym.sym, false);
key[event.key.keysym.sym] = false;
if (event.key.keysym.sym == SDLK_ESCAPE)
{
running = false;
}
else if (event.key.keysym.sym == SDLK_t)
{
kernelType = (kernelType + 1) % raytracerKernels.size();
kernel = raytracerKernels[kernelType];
cout << "Kernel Type: " << kernelType << endl;
cout << "Kernel: " << kernel->getName() << endl;
reloadKernel = true;
}
else if (event.key.keysym.sym == SDLK_y)
{
if (framesLeftToRender == -1)
{
framesLeftToRender = 3;
kernel = kernelPath;
kernelType = 3;
}
else
{
framesLeftToRender = -1;
kernel = kernelFast;
kernelType = 0;
}
reloadKernel = true;
delta = 0; // Reset delta smoothing
}
else if (event.key.keysym.sym == SDLK_l)
{
opencl = !opencl;
delta = 0; // Reset delta smoothing
}
else if (event.key.keysym.sym == SDLK_g)
{
grab = !grab;
if (grab)
{
SDL_ShowCursor(0);
SDL_WM_GrabInput(SDL_GRAB_ON);
}
else
{
SDL_ShowCursor(1);
SDL_WM_GrabInput(SDL_GRAB_OFF);
}
}
else if (event.key.keysym.sym == SDLK_i)
{
cout << camera.position.x << ", " << camera.position.y << ", " << camera.position.z << endl;
}
else if (event.key.keysym.sym == SDLK_o)
{
ofstream out("camera.txt");
out << camera.position.x << " " << camera.position.y << " " << camera.position.z;
out << " " << camera.getYaw() << " " << camera.getPitch() << endl;
out.close();
cout << "Camera state saved" << endl;
}
else if (event.key.keysym.sym == SDLK_p)
{
ifstream in("camera.txt");
float yaw, pitch;
in >> camera.position.x >> camera.position.y >> camera.position.z;
in >> yaw >> pitch;
in.close();
camera.setYaw(yaw);
camera.setPitch(pitch);
cout << "Camera state restored" << endl;
}
//if (event.key.keysym.sym == SDLK_ESCAPE) running = false;
break;
case SDL_KEYDOWN:
if (!grab) gui->keyPress(event.key.keysym.sym, true);
if (event.key.keysym.sym == SDLK_n)
onceBomb = true;
key[event.key.keysym.sym] = true;
break;
case SDL_MOUSEBUTTONDOWN:
if (!grab)
{
gui->mouseButton(ivec2(event.button.x, event.button.y), true);
gui->mouseMove(ivec2(event.button.x, event.button.y));
}
else if (event.button.button == 1)
deleteBlock();
break;
case SDL_MOUSEBUTTONUP:
if (!grab)
{
gui->mouseButton(ivec2(event.button.x, event.button.y), false);
gui->mouseMove(ivec2(event.button.x, event.button.y));
}
break;
case SDL_MOUSEMOTION:
if (grab)
{
if (event.motion.xrel != 0 || event.motion.yrel != 0)
forceMixOff = 1;
camera.changeYaw(0.012f * event.motion.xrel);
camera.changePitch(-0.015f * event.motion.yrel);
}
else
gui->mouseMove(ivec2(event.motion.x, event.motion.y));
break;
}
if (!running) break;
if (framesLeftToRender > 0)
{
if (--framesLeftToRender == 0)
cout << "Paused" << endl;
}
else if (framesLeftToRender == 0)
continue;
if (gui->getOption(1, 0).choice != kernelType)
{
kernelType = gui->getOption(1, 0).choice;
kernel = raytracerKernels[kernelType];
cout << "Kernel Type: " << kernelType << endl;
cout << "Kernel: " << kernel->getName() << endl;
reloadKernel = true;
}
{
int newState = gui->getOption(2, 1).choice == 1;
if (postprocessEnabled != newState)
{
postprocessEnabled = newState;
forceMixOff = 1;
}
}
{
timeOfDay = gui->getOption(0, 0).value;
//timeOfDay = (timeOfDay + delta * 0.025f);
if (timeOfDay > 24.0f) timeOfDay -= 24.0f;
float sunAngle = timeOfDay * M_PI / 12.0f;
scene.lightDirection = normalize(vec3(sin(sunAngle), -cos(sunAngle), 0.25f));
}
{
exposure = gui->getOption(2, 0).value;
exposure *= exposure;
}
{
float speed = 8.0f;
//acceleration.y = 9.8f * 4.0f;
velocity.x *= delta * 0.1f;
velocity.y *= delta * 0.1f;
velocity.z *= delta * 0.1f;
vec3 up = vec3(0.0f, 1.0f, 0.0f);
vec3 forward = camera.getWalkDirection();
vec3 right = cross(camera.getWalkDirection(), up);
if (key[SDLK_a])
{
velocity += right * speed;
forceMixOff = 1;
}
if (key[SDLK_d])
{
velocity -= right * speed;
forceMixOff = 1;
}
if (key[SDLK_w])
{
velocity += forward * speed;
forceMixOff = 1;
}
if (key[SDLK_s])
{
velocity -= forward * speed;
forceMixOff = 1;
}
if (key[SDLK_q])
{
velocity -= up * speed;
forceMixOff = 1;
}
if (key[SDLK_z])
{
velocity += up * speed;
forceMixOff = 1;
}
if (key[SDLK_b] || onceBomb)
{
Ray actionRay = Ray(camera.position, vec3(camera.getViewMatrix() * vec4(0.0f, 0.0f, 1.0f, 0.0f)));
VoxelGrid::Result result = voxels.intersect(actionRay, 100.0f);
if (result.hit)
{
voxelsDirty = true;
for (int x = -4; x <= 4; x++)
for (int y = -4; y <= 4; y++)
for (int z = -4; z <= 4; z++)
if (x*x + y*y + z*z < 16)
{
int xp = x + result.pos.x, yp = y + result.pos.y, zp = z + result.pos.z;
if (xp < 0 || xp >= voxelSize.x || yp < 0 || yp >= voxelSize.y || zp < 0 || zp >= voxelSize.z) continue;
voxels.get(xp, yp, zp) = 0;
}
}
}
else if (key[SDLK_DELETE])
{
deleteBlock();
}
else if (key[SDLK_RETURN])
{
Ray actionRay(camera.position, vec3(camera.getViewMatrix() * vec4(0.0f, 0.0f, 1.0f, 0.0f)));
VoxelGrid::Result result = voxels.intersect(actionRay, 100.0f);
if (result.hit)
voxels.get(result.pos.x, result.pos.y, result.pos.z) = 3;
}
long now = SDL_GetTicks();
bool firstPhysics = true;
for (;lastTimePhysics < now; lastTimePhysics += 10)
{
float pdelta = 10.0f / 1000.0f;
// Integrate position
camera.position += velocity * pdelta + acceleration * 0.5f * pdelta * pdelta;
velocity += acceleration * pdelta;
// Move out of terrain
Ray groundRay(camera.position, vec3(0.0f, 1.0f, 0.0f));
VoxelGrid::Result groundResult = voxels.intersect(groundRay, 100.0f);
if (groundResult.hit && groundResult.t < 2.0f)
{
camera.position.y -= 2.0f - groundResult.t;
if (velocity.y > 0.0f) velocity.y = 0.0f;
if (key[SDLK_SPACE] && velocity.y > -0.1f && firstPhysics) velocity.y = -20.0f;
}
firstPhysics = false;
}
}
if (opencl)
{
static int bufferIndex = 1;
bufferIndex = 1 - bufferIndex;
if (reloadKernel)
{
reloadKernel = false;
ret |= clSetKernelArg(kernel->kernel, 1, sizeof(cl_mem), (void *)&outputMemory);
ret |= clSetKernelArg(kernel->kernel, 2, sizeof(cl_mem), (void *)&cameraPosMemory);
ret |= clSetKernelArg(kernel->kernel, 3, sizeof(cl_mem), (void *)&cameraViewMatrixMemory);
ret |= clSetKernelArg(kernel->kernel, 4, sizeof(cl_mem), (void *)&lightDirectionMemory);
ret |= clSetKernelArg(kernel->kernel, 6, sizeof(cl_mem), (void *)&voxelSizeMemory);
ret |= clEnqueueWriteBuffer(command_queue, voxelMemory, CL_TRUE, 0, sizeof(uint8_t) * voxelCount, voxels.getVoxels(), 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(command_queue, octreeMemory, CL_TRUE, 0, sizeof(Node) * voxels.space, voxels.nodes, 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(command_queue, ptrTableMemory, CL_TRUE, 0, sizeof(uint32_t) * voxels.ptrSpace, voxels.ptrTable, 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(command_queue, voxelSizeMemory, CL_TRUE, 0, sizeof(int) * 3, &voxelSize, 0, NULL, NULL);
forceMixOff = 2;
}
if (forceMixOff > 0)
{
forceMixOff--;
succesiveFrames = 1;
}
else if (succesiveFrames < 128)
succesiveFrames++;
float mixFactor = 1.0f - 1.0f / succesiveFrames;
static int k = 0;
k++;
vec3 cameraPos = camera.position;
mat4 cameraViewMatrix = transpose(camera.getViewMatrix());
ret = clEnqueueWriteBuffer(command_queue, cameraPosMemory, CL_FALSE, 0, sizeof(vec3), &cameraPos,
0, NULL, NULL);
ret = clEnqueueWriteBuffer(command_queue, cameraViewMatrixMemory, CL_FALSE, 0, sizeof(mat4), &cameraViewMatrix,
0, NULL, NULL);
ret = clEnqueueWriteBuffer(command_queue, lightDirectionMemory, CL_FALSE, 0, sizeof(vec3), &scene.lightDirection,
0, NULL, NULL);
if (voxelsDirty)
{
voxelsDirty = false;
voxels.generateOctree();
ret |= clEnqueueWriteBuffer(command_queue, voxelMemory, CL_FALSE, 0, sizeof(uint8_t) * voxelCount, voxels.getVoxels(), 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(command_queue, octreeMemory, CL_FALSE, 0, sizeof(Node) * voxels.space, voxels.nodes, 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(command_queue, ptrTableMemory, CL_FALSE, 0, sizeof(uint32_t) * voxels.ptrSpace, voxels.ptrTable, 0, NULL, NULL);
}
// Execute the OpenCL kernel on the list
ret = clSetKernelArg(kernel->kernel, 0, sizeof(int), (void *)&k);
ret = clSetKernelArg(kernel->kernel, 5, sizeof(cl_mem), (void *)&voxelMemory);
ret = clSetKernelArg(kernel->kernel, 7, sizeof(cl_mem), (void *)&octreeMemory);
ret = clSetKernelArg(kernel->kernel, 8, sizeof(cl_mem), (void *)&ptrTableMemory);
glFinish();
ret = clEnqueueAcquireGLObjects(command_queue, 1, &outputMemory, 0, NULL, NULL);
ret = clEnqueueAcquireGLObjects(command_queue, 1, &screenMemory[0], 0, NULL, NULL);
ret = clEnqueueAcquireGLObjects(command_queue, 1, &screenMemory[1], 0, NULL, NULL);
ret = clEnqueueAcquireGLObjects(command_queue, 1, &finalMemory, 0, NULL, NULL);
size_t global_item_size[2];
size_t local_item_size[2] = {8, 8};
global_item_size[0] = WIDTH;//((WIDTH + 7) >> 3) << 3;
global_item_size[1] = HEIGHT;//((HEIGHT + 15) >> 4) << 4;
clEnqueueNDRangeKernel(command_queue, kernel->kernel, 2, NULL, global_item_size, local_item_size, 0, NULL, NULL);
if (postprocessEnabled)
{
clSetKernelArg(postprocess.kernel, 0, sizeof(int), (void *)&k);
clSetKernelArg(postprocess.kernel, 1, sizeof(float), (void *)&mixFactor);
clSetKernelArg(postprocess.kernel, 2, sizeof(cl_mem), (void *)&outputMemory);
clSetKernelArg(postprocess.kernel, 3, sizeof(cl_mem), (void *)&screenMemory[1-bufferIndex]);
clSetKernelArg(postprocess.kernel, 4, sizeof(cl_mem), (void *)&screenMemory[bufferIndex]);
clEnqueueNDRangeKernel(command_queue, postprocess.kernel, 2, NULL, global_item_size, local_item_size, 0, NULL, NULL);
}
if (postprocessEnabled)
clSetKernelArg(tonemapper.kernel, 0, sizeof(cl_mem), (void *)&screenMemory[bufferIndex]);
else
clSetKernelArg(tonemapper.kernel, 0, sizeof(cl_mem), (void *)&outputMemory);
clSetKernelArg(tonemapper.kernel, 1, sizeof(cl_mem), (void *)&finalMemory);
clSetKernelArg(tonemapper.kernel, 2, sizeof(float), (void *)&exposure);
clEnqueueNDRangeKernel(command_queue, tonemapper.kernel, 2, NULL, global_item_size, local_item_size, 0, NULL, NULL);
clEnqueueReleaseGLObjects(command_queue, 1, &outputMemory, 0, NULL, NULL);
clEnqueueReleaseGLObjects(command_queue, 1, &screenMemory[0], 0, NULL, NULL);
clEnqueueReleaseGLObjects(command_queue, 1, &screenMemory[1], 0, NULL, NULL);
clEnqueueReleaseGLObjects(command_queue, 1, &finalMemory, 0, NULL, NULL);
clFinish(command_queue);
drawToScreen();
}
else
{
static int k = 0;
#pragma omp parallel for schedule(dynamic) lastprivate(k)
for (int y = 0; y < HEIGHT; y++)
for (int x = 0; x < WIDTH; x++)
pixels[x + y * WIDTH] = raytracer->trace(x, y, x + y * WIDTH + k++);
glBindTexture(GL_TEXTURE_2D, screenTexture[0]);
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, WIDTH, HEIGHT, GL_RGBA, GL_FLOAT, (GLvoid*)pixels);
glBindTexture(GL_TEXTURE_2D, 0);
drawToScreen();
}
long now = SDL_GetTicks();
//if (now - start > 10000) break;
float cdelta = (now - last) * 0.001f;
last = now;
delta = (delta > 0.0f) ? (delta * 2.0f + cdelta) * 0.333f : cdelta;
if (nextReport < now)
{
cout << fixed << setprecision(2) << (delta * 1000) << "ms - FPS: " << (1.0f / delta) << endl;
nextReport = now + 500;
}
gui->update(delta, !grab);
}
static int crash(lua_State *L) {
Kernel *pKernel = Kernel::getInstance();
assert(pKernel);
pKernel->crash();
return 0;
}
int main() {
// Default pins to low status
for (int i = 0; i < 5; i++){
leds[i].output();
leds[i] = (i & 1) ^ 1;
}
//bool sdok= (sd.disk_initialize() == 0);
sd.disk_initialize();
Kernel* kernel = new Kernel();
kernel->streams->printf("Smoothie ( grbl port ) version 0.8.0-rc1 with new accel @%ldMHz\r\n", SystemCoreClock / 1000000);
Version version;
kernel->streams->printf(" Build version %s, Build date %s\r\n", version.get_build(), version.get_build_date());
// Create and add main modules
kernel->add_module( new Laser() );
kernel->add_module( new Extruder() );
kernel->add_module( new SimpleShell() );
kernel->add_module( new Configurator() );
kernel->add_module( new CurrentControl() );
kernel->add_module( new TemperatureControlPool() );
kernel->add_module( new SwitchPool() );
kernel->add_module( new PauseButton() );
kernel->add_module( new PlayLed() );
kernel->add_module( new Endstops() );
kernel->add_module( new Player() );
kernel->add_module( new Panel() );
kernel->add_module( new Touchprobe() );
// Create and initialize USB stuff
u.init();
//if(sdok) { // only do this if there is an sd disk
// msc= new USBMSD(&u, &sd);
// kernel->add_module( msc );
//}
kernel->add_module( &msc );
kernel->add_module( &usbserial );
if( kernel->config->value( second_usb_serial_enable_checksum )->by_default(false)->as_bool() ){
kernel->add_module( new USBSerial(&u) );
}
kernel->add_module( &dfu );
kernel->add_module( &u );
// clear up the config cache to save some memory
kernel->config->config_cache_clear();
// Main loop
while(1){
kernel->call_event(ON_MAIN_LOOP);
kernel->call_event(ON_IDLE);
}
}
void build_kernel(Call<B, I, J> e, Kernel<T>& k)
{
vsip::index_type y = k.origin(0) + offset(e.i());
vsip::index_type x = k.origin(1) + offset(e.j());
k(y, x) += T(1);
};
static int sleep(lua_State *L) {
Kernel *pKernel = Kernel::getInstance();
assert(pKernel);
pKernel->sleep(static_cast<uint>(luaL_checknumber(L, 1) * 1000));
return 0;
}
double LeastMedianOfSquares(const Kernel &kernel,
typename Kernel::Model * model = NULL,
double* outlierThreshold = NULL,
double outlierRatio=0.5,
double minProba=0.99)
{
const size_t min_samples = Kernel::MINIMUM_SAMPLES;
const size_t total_samples = kernel.NumSamples();
std::vector<double> residuals(total_samples); // Array for storing residuals
std::vector<size_t> vec_sample(min_samples);
double dBestMedian = std::numeric_limits<double>::max();
// Required number of iterations is evaluated from outliers ratio
const size_t N = (min_samples<total_samples)?
getNumSamples(minProba, outlierRatio, min_samples): 0;
for (size_t i=0; i < N; i++) {
// Get Samples indexes
UniformSample(min_samples, total_samples, &vec_sample);
// Estimate parameters: the solutions are stored in a vector
std::vector<typename Kernel::Model> models;
kernel.Fit(vec_sample, &models);
// Now test the solutions on the whole data
for (size_t k = 0; k < models.size(); ++k) {
//Compute Residuals :
for (size_t l = 0; l < total_samples; ++l) {
double error = kernel.Error(l, models[k]);
residuals[l] = error;
}
// Compute median
std::vector<double>::iterator itMedian = residuals.begin() +
std::size_t( total_samples*(1.-outlierRatio) );
std::nth_element(residuals.begin(), itMedian, residuals.end());
double median = *itMedian;
// Store best solution
if(median < dBestMedian) {
dBestMedian = median;
if (model) (*model) = models[k];
}
}
}
// This array of precomputed values corresponds to the inverse
// cumulative function for a normal distribution. For more information
// consult the litterature (Robust Regression for Outlier Detection,
// rouseeuw-leroy). The values are computed for each 5%
static const double ICDF[21] = {
1.4e16, 15.94723940, 7.957896558, 5.287692054,
3.947153876, 3.138344200, 2.595242369, 2.203797543,
1.906939402, 1.672911853, 1.482602218, 1.323775627,
1.188182950, 1.069988721, 0.9648473415, 0.8693011162,
0.7803041458, 0.6946704675, 0.6079568319,0.5102134568,
0.3236002672
};
// Evaluate the outlier threshold
if(outlierThreshold) {
double sigma = ICDF[int((1.-outlierRatio)*20.)] *
(1. + 5. / double(total_samples - min_samples));
*outlierThreshold = (double)(sigma * sigma * dBestMedian * 4.);
if (N==0) *outlierThreshold = std::numeric_limits<double>::max();
}
return dBestMedian;
}
/** Set the outline of the assembly. Creates an Object and sets m_shape point to
* it.
* All child components must be detectors and positioned along a straight line
* and have the same shape.
* The shape can be either a box or a cylinder.
* @return The shape of the outline: "cylinder", "box", ...
*/
boost::shared_ptr<Object> ObjCompAssembly::createOutline() {
if (group.empty()) {
throw Kernel::Exception::InstrumentDefinitionError("Empty ObjCompAssembly");
}
if (nelements() < 2) {
g_log.warning("Creating outline with fewer than 2 elements. The outline displayed may be inaccurate.");
}
// Get information about the shape and size of a detector
std::string type;
int otype;
std::vector<Kernel::V3D> vectors;
double radius, height;
boost::shared_ptr<const Object> obj = group.front()->shape();
if (!obj) {
throw Kernel::Exception::InstrumentDefinitionError(
"Found ObjComponent without shape");
}
obj->GetObjectGeom(otype, vectors, radius, height);
if (otype == 1) {
type = "box";
} else if (otype == 3) {
type = "cylinder";
} else {
throw std::runtime_error(
"IDF \"outline\" option is only allowed for assemblies containing "
"components of types \"box\" or \"cylinder\".");
}
// Calculate the dimensions of the outline object
// find the 'moments of inertia' of the assembly
double Ixx = 0, Iyy = 0, Izz = 0, Ixy = 0, Ixz = 0, Iyz = 0;
V3D Cmass; // 'center of mass' of the assembly
for (const_comp_it it = group.begin(); it != group.end(); it++) {
V3D p = (**it).getRelativePos();
Cmass += p;
}
Cmass /= nelements();
for (const_comp_it it = group.begin(); it != group.end(); it++) {
V3D p = (**it).getRelativePos();
double x = p.X() - Cmass.X(), x2 = x * x;
double y = p.Y() - Cmass.Y(), y2 = y * y;
double z = p.Z() - Cmass.Z(), z2 = z * z;
Ixx += y2 + z2;
Iyy += x2 + z2;
Izz += y2 + x2;
Ixy -= x * y;
Ixz -= x * z;
Iyz -= y * z;
}
// principal axes of the outline shape
// vz defines the line through all pixel centres
V3D vx, vy, vz;
if (Ixx == 0) // pixels along x axis
{
vx = V3D(0, 1, 0);
vy = V3D(0, 0, 1);
vz = V3D(1, 0, 0);
} else if (Iyy == 0) // pixels along y axis
{
vx = V3D(0, 0, 1);
vy = V3D(1, 0, 0);
vz = V3D(0, 1, 0);
} else if (Izz == 0) // pixels along z axis
{
vx = V3D(1, 0, 0);
vy = V3D(0, 1, 0);
vz = V3D(0, 0, 1);
} else {
// Either the detectors are not perfectrly aligned or
// vz is parallel to neither of the 3 axis x,y,z
// This code is unfinished
DblMatrix II(3, 3), Vec(3, 3), D(3, 3);
II[0][0] = Ixx;
II[0][1] = Ixy;
II[0][2] = Ixz;
II[1][0] = Ixy;
II[1][1] = Iyy;
II[1][2] = Iyz;
II[2][0] = Ixz;
II[2][1] = Iyz;
II[2][2] = Izz;
std::cerr << "Inertia matrix: " << II << II.determinant() << '\n';
II.Diagonalise(Vec, D);
std::cerr << "Vectors: " << Vec << '\n';
std::cerr << "Moments: " << D << '\n';
vx = V3D(Vec[0][0], Vec[1][0], Vec[2][0]);
vy = V3D(Vec[0][1], Vec[1][1], Vec[2][1]);
vz = V3D(Vec[0][2], Vec[1][2], Vec[2][2]);
}
// find assembly sizes along the principal axes
double hx, hy, hz; // sizes along x,y, and z axis
// maximum displacements from the mass centre along axis vx,vy, and vz
// in positive (p) and negative (n) directions.
// positive displacements are positive numbers and negative ones are negative
double hxn = 0, hyn = 0, hzn = 0;
double hxp = 0, hyp = 0, hzp = 0;
for (const_comp_it it = group.begin(); it != group.end(); it++) {
// displacement vector of a detector
V3D p = (**it).getRelativePos() - Cmass;
// its projection on the vx axis
double h = p.scalar_prod(vx);
if (h > hxp)
hxp = h;
if (h < hxn)
hxn = h;
// its projection on the vy axis
h = p.scalar_prod(vy);
if (h > hyp)
hyp = h;
if (h < hyn)
hyn = h;
// its projection on the vz axis
h = p.scalar_prod(vz);
if (h > hzp)
hzp = h;
if (h < hzn)
hzn = h;
}
// calc the assembly sizes
hx = hxp - hxn;
hy = hyp - hyn;
hz = hzp - hzn;
// hx and hy must be practically zero
if (hx > 1e-3 || hy > 1e-3) // arbitrary numbers
{
throw Kernel::Exception::InstrumentDefinitionError(
"Detectors of a ObjCompAssembly do not lie on a staraight line");
}
// determine the order of the detectors to make sure that the texture
// coordinates are correct
bool firstAtBottom; // first detector is at the bottom of the outline shape
// the bottom end is the one with the negative displacement from the centre
firstAtBottom =
((**group.begin()).getRelativePos() - Cmass).scalar_prod(vz) < 0;
// form the input string for the ShapeFactory
std::ostringstream obj_str;
if (type == "box") {
if (hz == 0)
hz = 0.1;
hx = hy = 0;
height = 0;
V3D p0 = vectors[0];
for (size_t i = 1; i < vectors.size(); ++i) {
V3D p = vectors[i] - p0;
double h = fabs(p.scalar_prod(vx));
if (h > hx)
hx = h;
h = fabs(p.scalar_prod(vy));
if (h > hy)
hy = h;
height = fabs(p.scalar_prod(vz));
if (h > height)
height = h;
}
vx *= hx / 2;
vy *= hy / 2;
vz *= hzp + height / 2;
if (!firstAtBottom) {
vz = vz * (-1);
}
// define the outline shape as cuboid
V3D p_lfb = Cmass - vx - vy - vz;
V3D p_lft = Cmass - vx - vy + vz;
V3D p_lbb = Cmass - vx + vy - vz;
V3D p_rfb = Cmass + vx - vy - vz;
obj_str << "<cuboid id=\"shape\">";
obj_str << "<left-front-bottom-point ";
obj_str << "x=\"" << p_lfb.X();
obj_str << "\" y=\"" << p_lfb.Y();
obj_str << "\" z=\"" << p_lfb.Z();
obj_str << "\" />";
obj_str << "<left-front-top-point ";
obj_str << "x=\"" << p_lft.X();
obj_str << "\" y=\"" << p_lft.Y();
obj_str << "\" z=\"" << p_lft.Z();
obj_str << "\" />";
obj_str << "<left-back-bottom-point ";
obj_str << "x=\"" << p_lbb.X();
obj_str << "\" y=\"" << p_lbb.Y();
obj_str << "\" z=\"" << p_lbb.Z();
obj_str << "\" />";
obj_str << "<right-front-bottom-point ";
obj_str << "x=\"" << p_rfb.X();
obj_str << "\" y=\"" << p_rfb.Y();
obj_str << "\" z=\"" << p_rfb.Z();
obj_str << "\" />";
obj_str << "</cuboid>";
} else if (type == "cylinder") {
// the outline is one detector height short
hz += height;
// shift Cmass to the end of the cylinder where the first detector is
if (firstAtBottom) {
Cmass += vz * hzn;
} else {
hzp += height;
Cmass += vz * hzp;
// inverse the vz axis
vz = vz * (-1);
}
obj_str << "<segmented-cylinder id=\"stick\">";
obj_str << "<centre-of-bottom-base ";
obj_str << "x=\"" << Cmass.X();
obj_str << "\" y=\"" << Cmass.Y();
obj_str << "\" z=\"" << Cmass.Z();
obj_str << "\" />";
obj_str << "<axis x=\"" << vz.X() << "\" y=\"" << vz.Y() << "\" z=\""
<< vz.Z() << "\" /> ";
obj_str << "<radius val=\"" << radius << "\" />";
obj_str << "<height val=\"" << hz << "\" />";
obj_str << "</segmented-cylinder>";
}
if (!obj_str.str().empty()) {
boost::shared_ptr<Object> s = ShapeFactory().createShape(obj_str.str());
setOutline(s);
return s;
}
return boost::shared_ptr<Object>();
}
void LibSvmUtils::setKernelParams(const Kernel& kernel, struct svm_parameter *params) const {
LibSvmKernelParamSetter paramSetter(params);
kernel.accept(paramSetter);
}
void kernelStartPoint(multiboot_info_t* mbd, unsigned int magic)
{
Kernel kernel;
kernel.main(mbd, magic);
}
typename Kernel::Model RANSAC(
const Kernel &kernel,
const Scorer &scorer,
std::vector<size_t> *best_inliers = nullptr ,
size_t *best_score = nullptr , // Found number of inliers
double outliers_probability = 1e-2)
{
assert(outliers_probability < 1.0);
assert(outliers_probability > 0.0);
size_t iteration = 0;
const size_t min_samples = Kernel::MINIMUM_SAMPLES;
const size_t total_samples = kernel.NumSamples();
size_t max_iterations = 100;
const size_t really_max_iterations = 4096;
size_t best_num_inliers = 0;
double best_inlier_ratio = 0.0;
typename Kernel::Model best_model;
// Test if we have sufficient points for the kernel.
if (total_samples < min_samples) {
if (best_inliers) {
best_inliers->resize(0);
}
return best_model;
}
// In this robust estimator, the scorer always works on all the data points
// at once. So precompute the list ahead of time [0,..,total_samples].
std::vector<size_t> all_samples(total_samples);
std::iota(all_samples.begin(), all_samples.end(), 0);
std::vector<size_t> sample;
for (iteration = 0;
iteration < max_iterations &&
iteration < really_max_iterations; ++iteration) {
UniformSample(min_samples, &all_samples, &sample);
std::vector<typename Kernel::Model> models;
kernel.Fit(sample, &models);
// Compute the inlier list for each fit.
for (size_t i = 0; i < models.size(); ++i) {
std::vector<size_t> inliers;
scorer.Score(kernel, models[i], all_samples, &inliers);
if (best_num_inliers < inliers.size()) {
best_num_inliers = inliers.size();
best_inlier_ratio = inliers.size() / double(total_samples);
best_model = models[i];
if (best_inliers) {
best_inliers->swap(inliers);
}
}
if (best_inlier_ratio) {
max_iterations = IterationsRequired(min_samples,
outliers_probability,
best_inlier_ratio);
}
}
}
if (best_score)
*best_score = best_num_inliers;
return best_model;
}
int main (int argc, char **argv)
{
//Initialize GLUT
initGlut(argc, argv);
Kernel kernel;
kernel.enableVerticalSync( false );
renderer = kernel.getRenderer();
printf( "Kernel loaded\n" );
//Load shape model
LoaderObj ldr;
File f( "zekko.obj" );
ldr.setUVMeshClass( Class( ETexMesh ));
int start = GE::Time::GetTicks();
ldr.loadFile( f.getPathName() );
int end = GE::Time::GetTicks();
printf( "Time: %d\n", end - start );
PolyMesh *dmesh;
uvmesh = (ETexMesh*) ldr.getFirstResource( Class(TexMesh) );
dmesh = (PolyMesh*) ldr.getFirstResource( Class(PolyMesh) );
zekko = (PolyMeshActor*) ldr.getFirstObject( Class(PolyMeshActor) );
if (dmesh == NULL) return EXIT_FAILURE;
if (uvmesh == NULL) return EXIT_FAILURE;
if (zekko == NULL) return EXIT_FAILURE;
//Check if UV mesh type correct
printf( "uvmesh = %s\n", StringOf( ClassOf( zekko->getTexMesh() )));
printf( "uvmesh %s ETexMesh\n",
SafeCast( ETexMesh, zekko->getTexMesh() ) ?
"IS" : "is NOT");
printf( "uvmesh %s TexMesh\n",
SafeCast( TexMesh, zekko->getTexMesh() ) ?
"IS" : "is NOT");
printf( "uvmesh %s Resource\n",
SafeCast( Resource, zekko->getTexMesh() ) ?
"IS" : "is NOT");
printf( "uvmesh %s PolyMesh\n",
SafeCast( PolyMesh, zekko->getTexMesh() ) ?
"IS" : "is NOT");
//Setup camera
cam3D.translate( 0,35,-100 );
cam3D.setCenter( center );
//cam3D.setNearClipPlane(100.0f);
//cam3D.setFarClipPlane(100000.0f);
//Find model center
findCenter();
//Set half bunny green
for (PolyMesh::FaceIter f(dmesh); !f.end(); ++f) {
//f->smoothGroups = 0x1;
if (f->firstHedge()->dstVertex()->point.y > center.y)
dmesh->setMaterialID( *f, 1 ); }
//Convert to static mesh
//dmesh->updateNormals();
smesh.fromPoly( dmesh, uvmesh );
//Load texture image
imgDiff.readFile( "texture.jpg", "" );
//Create texture from image
texDiff = new Texture;
texDiff->fromImage( &imgDiff );
//Load specularity image
Image imgSpec;
imgSpec.readFile( "specularity.jpg", "" );
//Create specularity texture
texSpec = new Texture();
texSpec->fromImage(&imgSpec);
//Create material using texture
PhongMaterial mat;
mat.setDiffuseTexture( texDiff );
//mat.setSpecularityTexture( texSpec );
mat.setDiffuseColor( Vector3( 1,0.0f,0.0f ));
mat.setAmbientColor( Vector3( 0.2f,0.2f,0.2f ));
mat.setSpecularity( 1.0f );
mat.setGlossiness( 0.2f );
PhongMaterial mat2;
mat2.setDiffuseTexture( texDiff );
//mat2.setSpecularityTexture( texSpec );
mat2.setDiffuseColor( Vector3( 0.0f, 1.0f, 0.0f ));
mat2.setAmbientColor( Vector3( 0.2f, 0.2f, 0.2f ));
mat2.setSpecularity( 1.0f );
mat2.setGlossiness( 0.2f );
MultiMaterial mm;
mm.setNumSubMaterials( 2 );
mm.setSubMaterial( 0, &mat );
mm.setSubMaterial( 1, &mat2 );
zekko->setMaterial( &mm );
scene = new Scene;
scene->addChild( zekko );
//Light *light = new SpotLight( Vector3(-200,200,-200), Vector3(1,-1,1), 60, 0 );
Light *light = new HeadLight;
scene->addChild( light );
lblFps.setLocation( Vector2( 0.0f,(Float)resY ));
lblFps.setColor( Vector3( 1.0f,1.0f,1.0f ));
//Run application
atexit( cleanup );
glutMainLoop();
cleanup();
return EXIT_SUCCESS;
}
/**
* @param x X coordinate
* @param y Y coordinate
*/
void ConvexPolygon::insert(double x, double y) { this->insert(V2D(x, y)); }
/* ========================================================================== */
int sci_gpuApplyFunction(char *fname)
{
SciErr sciErr;
int block_w = 0, block_h = 0, grid_w = 0, grid_h = 0;
int row = 0, col = 0;
int argnum = 0;
int retnumbvalue = 0;
double* MM = NULL;
int *ptr = NULL;
int *lstptr = NULL;
int *dptr[4];
int inputType_1 = 0;
int inputType_2 = 0;
CheckRhs(6, 6);
CheckLhs(1, 1);
try
{
void *fptr = NULL;
if (!isGpuInit())
{
throw "gpu is not initialised. Please launch gpuInit() before use this function.";
}
sciErr = getVarAddressFromPosition(pvApiCtx, 1, &ptr);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getVarType(pvApiCtx, ptr, &inputType_1);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getVarAddressFromPosition(pvApiCtx, 2, &lstptr);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getVarType(pvApiCtx, lstptr, &inputType_2);
if (sciErr.iErr)
{
throw sciErr;
}
if (inputType_1 != sci_pointer)
{
throw "gpuApplyFuntion : Bad type for input argument #1: A string expected.";
}
else if (inputType_2 != sci_list)
{
throw "gpuApplyFuntion : Bad type for input argument #2: A list expected.";
}
else
{
for (int i = 3; i < 7; i++)
{
int inputType = 0;
sciErr = getVarAddressFromPosition(pvApiCtx, i, &dptr[i - 3]);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getVarType(pvApiCtx, dptr[i - 3], &inputType);
if (sciErr.iErr)
{
throw sciErr;
}
if (inputType != sci_matrix)
{
char string[64];
sprintf(string, "gpuApplyFunction : Bad type for input argument #%d : A matrix expected.", i);
throw string;
}
}
sciErr = getPointer(pvApiCtx, ptr, (void**)&fptr);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getListItemNumber(pvApiCtx, lstptr, &argnum);
if (sciErr.iErr)
{
throw sciErr;
}
sciErr = getMatrixOfDouble(pvApiCtx, dptr[0], &row, &col, &MM);
if (sciErr.iErr)
{
throw sciErr;
}
if (row * col != 1)
{
throw "gpuApplyFunction : Bad size for input argument #3: A scalar expected.";
}
block_h = (int)MM[0];
sciErr = getMatrixOfDouble(pvApiCtx, dptr[1], &row, &col, &MM);
if (sciErr.iErr)
{
throw sciErr;
}
if (row * col != 1)
{
throw "gpuApplyFunction : Bad size for input argument #4: A scalar expected.";
}
block_w = (int)MM[0];
sciErr = getMatrixOfDouble(pvApiCtx, dptr[2], &row, &col, &MM);
if (sciErr.iErr)
{
throw sciErr;
}
if (row * col != 1)
{
throw "gpuApplyFunction : Bad size for input argument #5: A scalar expected.";
}
grid_h = (int)MM[0];
sciErr = getMatrixOfDouble(pvApiCtx, dptr[3], &row, &col, &MM);
if (sciErr.iErr)
{
throw sciErr;
}
if (row * col != 1)
{
throw "gpuApplyFunction : Bad size for input argument #6: A scalar expected.";
}
grid_w = (int)MM[0];
#ifdef WITH_CUDA
if (useCuda())
{
Kernel<ModeDefinition<CUDA> >* fptrCuda = (Kernel<ModeDefinition<CUDA> >*)fptr;
sci_CUDA_getArgs(fptrCuda, lstptr, argnum, fname);
fptrCuda->launch(getCudaQueue() , block_w, block_h, grid_w, grid_h);
}
#endif
#ifdef WITH_OPENCL
if (!useCuda())
{
Kernel<ModeDefinition<OpenCL> >* fptrOpenCL = (Kernel<ModeDefinition<OpenCL> >*)fptr;
sci_OpenCL_getArgs(fptrOpenCL, lstptr, argnum, fname);
fptrOpenCL->launch(getOpenClQueue(), block_w, block_h, grid_w * block_w, grid_h * block_h);
}
#endif
PutLhsVar();
}
}
catch (const char* str)
{
Scierror(999, "%s: %s\n", fname, str);
}
catch (SciErr E)
{
printError(&E, 0);
}
return 0;
}
int LoadPlugins (Kernel & k)
{
k.RegisterNode (new Node0_5) ;
return 1 ;
}
static void runPingTest(TestRunner& tr)
{
tr.test("Ping");
Config cfg = tr.getApp()->getConfig()["pong"];
// Stats and control
PingPong pingPong(cfg);
// create kernel
Kernel k;
// set thread stack size in engine (128k)
k.getEngine()->getThreadPool()->setThreadStackSize(
cfg["threadStackSize"]->getUInt32());
// optional for testing --
// limit threads to 2: one for accepting, 1 for handling
//k.getEngine()->getThreadPool()->setPoolSize(2);
k.getEngine()->getThreadPool()->setPoolSize(cfg["threads"]->getUInt32());
// start engine
k.getEngine()->start();
// create server
Server server;
server.setMaxConnectionCount(cfg["maxConnections"]->getInt32());
InternetAddress address("0.0.0.0", cfg["port"]->getUInt32());
// create SSL/generic http connection servicer
HttpConnectionServicer hcs;
// SslContext context;
// SslSocketDataPresenter presenter1(&context);
// NullSocketDataPresenter presenter2;
// SocketDataPresenterList list(false);
// list.add(&presenter1);
// list.add(&presenter2);
//server.addConnectionService(&address, &hcs);//, &list);
server.addConnectionService(
&address, &hcs, NULL, "pong",
cfg["maxConnections"]->getInt32(),
cfg["backlog"]->getInt32());
// create test http request servicer
NoContentServicer noContentSrv(&pingPong, "/");
hcs.addRequestServicer(&noContentSrv, false);
PingServicer ping(&pingPong, "/pong");
hcs.addRequestServicer(&ping, false);
const int bufsize = 4096;
char buf[bufsize];
memset(buf, '.', bufsize);
DataServicer data(&pingPong, "/data", buf, bufsize);
hcs.addRequestServicer(&data, false);
StatsServicer stats(&pingPong, "/stats");
hcs.addRequestServicer(&stats, false);
ResetServicer reset(&pingPong, "/reset");
hcs.addRequestServicer(&reset, false);
QuitServicer quit(&pingPong, "/quit");
hcs.addRequestServicer(&quit, false);
if(server.start(&k))
{
uint64_t num = cfg["num"]->getUInt64();
MO_INFO("Server started.");
if(num == 0)
{
MO_INFO("Servicing forever. CTRL-C to quit.");
}
{
MO_INFO("Servicing approximately %" PRIu64 " connections.", num);
}
}
else if(Exception::get() != NULL)
{
MO_ERROR("Server started with errors=%s\n",
Exception::get()->getMessage());
}
// start timing
pingPong.reset();
// either serve for limited time, or wait for lock
uint32_t time = cfg["time"]->getUInt32();
if(time != 0)
{
Thread::sleep(time);
}
else
{
pingPong.getLock().wait();
}
server.stop();
MO_INFO("Server stopped.");
// stop kernel engine
k.getEngine()->stop();
tr.passIfNoException();
}
/**
Run solver iterations.
*/
int SweepSolver (Grid_Data *grid_data, bool block_jacobi)
{
Kernel *kernel = grid_data->kernel;
int mpi_rank = 0;
#ifdef KRIPKE_USE_MPI
MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);
#endif
BLOCK_TIMER(grid_data->timing, Solve);
// Loop over iterations
double part_last = 0.0;
for(int iter = 0;iter < grid_data->niter;++ iter){
/*
* Compute the RHS: rhs = LPlus*S*L*psi + Q
*/
// Discrete to Moments transformation (phi = L*psi)
{
BLOCK_TIMER(grid_data->timing, LTimes);
kernel->LTimes(grid_data);
}
// Compute Scattering Source Term (psi_out = S*phi)
{
BLOCK_TIMER(grid_data->timing, Scattering);
kernel->scattering(grid_data);
}
// Compute External Source Term (psi_out = psi_out + Q)
{
BLOCK_TIMER(grid_data->timing, Source);
kernel->source(grid_data);
}
// Moments to Discrete transformation (rhs = LPlus*psi_out)
{
BLOCK_TIMER(grid_data->timing, LPlusTimes);
kernel->LPlusTimes(grid_data);
}
/*
* Sweep (psi = Hinv*rhs)
*/
{
BLOCK_TIMER(grid_data->timing, Sweep);
if(true){
// Create a list of all groups
std::vector<int> sdom_list(grid_data->subdomains.size());
for(int i = 0;i < grid_data->subdomains.size();++ i){
sdom_list[i] = i;
}
// Sweep everything
SweepSubdomains(sdom_list, grid_data, block_jacobi);
}
// This is the ARDRA version, doing each groupset sweep independently
else{
for(int group_set = 0;group_set < grid_data->num_group_sets;++ group_set){
std::vector<int> sdom_list;
// Add all subdomains for this groupset
for(int s = 0;s < grid_data->subdomains.size();++ s){
if(grid_data->subdomains[s].idx_group_set == group_set){
sdom_list.push_back(s);
}
}
// Sweep the groupset
SweepSubdomains(sdom_list, grid_data, block_jacobi);
}
}
}
{
BLOCK_TIMER(grid_data->timing, ParticleEdit);
double part = grid_data->particleEdit();
if(mpi_rank==0){
printf("iter %d: particle count=%e, change=%e\n", iter, part, (part-part_last)/part);
}
part_last = part;
}
}
return(0);
}
int main(int argc, char **argv)
{
srand((unsigned)time(NULL));
Kernel kernel;
CommandQueue queue;
Context context;
{
std::vector<Platform> platformList;
Platform::get(&platformList);
clog << "Platform number is: " << platformList.size() << endl;
std::string platformVendor;
platformList[0].getInfo((cl_platform_info)CL_PLATFORM_VENDOR, &platformVendor);
clog << "Platform is by: " << platformVendor << "\n";
cl_context_properties cprops[] = {
CL_CONTEXT_PLATFORM, (cl_context_properties) platformList[0](),
0
};
context = Context(GET_TARGET_PLATFORM, cprops);
std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
queue = CommandQueue(context, devices[0]);
std::string sourceCode = "#include \"es.cl\"\n";
Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1));
Program program = Program(context, source);
try
{
program.build(devices, "-I.");
}
catch (Error &)
{
std::string errors;
program.getBuildInfo(devices[0], CL_PROGRAM_BUILD_LOG, &errors);
std::cerr << "Build log: " << endl << errors << endl;
return 1;
}
kernel = Kernel(program, "es");
}
individual *individuals = new individual[LAMBDA];
for (int i = 0; i < LAMBDA; i++)
{
for (int j = 0; j < DIM; ++j)
{
individuals[i].x[j] = (rand()/((float)RAND_MAX)) * (XMAX-XMIN) + XMIN;
individuals[i].s[j] = (XMAX-XMIN) / 6.f;
}
for (int j = 0; j < DIM_A; ++j)
{
individuals[i].a[j] = (rand()/((float)RAND_MAX)) * (2*PI) - PI;
}
individuals[i].fitness = 0;
}
float gbest = std::numeric_limits<float>::infinity(), xbest[DIM];
Buffer esBuffer = Buffer(context, 0, INDIVIDUALS_SIZE);
Event ev;
queue.enqueueMapBuffer(esBuffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, INDIVIDUALS_SIZE);
for (int i = 0; i < 1000; i++)
{
queue.enqueueWriteBuffer(esBuffer, CL_TRUE, 0, INDIVIDUALS_SIZE, individuals);
kernel.setArg(1, (cl_ulong)rand());
kernel.setArg(0, esBuffer);
queue.enqueueNDRangeKernel(kernel, NullRange, NDRange(LAMBDA), NDRange(1), NULL, &ev);
ev.wait();
queue.enqueueReadBuffer(esBuffer, CL_TRUE, 0, INDIVIDUALS_SIZE, individuals);
std::sort(individuals, individuals + LAMBDA, individual_comp);
individual mean = get_mean(individuals);
for (int j = 0; j < LAMBDA; ++j)
{
individuals[j] = mean;
}
}
gbest = individuals[0].fitness;
for (int i = 0; i < DIM; ++i) xbest[i] = individuals[0].x[i];
clog << "Best value " << gbest << " found at (";
for (int i = 0; i < DIM; ++i) clog << xbest[i] << (i == DIM-1 ? ")" : ", ");
clog << "\n";
clog << "Our computation estemates it: f(" << xbest[0] << ", ..., " << xbest[DIM-1] << ") = " << es_f(xbest) << endl;
delete[] individuals;
return 0;
}
int LoadPlugins (Kernel & k)
{
//FETS printf ("dht11::LoadPlugins %d\n", __LINE__);
k.RegisterNode (new NodeDhtXX) ;
return 1 ;
}
