// compute and draw image
void ComputeImage() {
const double t0 = WallClockTime();
// copy scene from host to device
cl::CommandQueue &oclQueue = deviceQueues[0];
oclQueue.enqueueWriteBuffer(*sceneBuff,
CL_FALSE,
0,
sceneBuff->getInfo<CL_MEM_SIZE>(),
scene);
size_t globalThreads = windowWidth * windowHeight;
if (globalThreads % kernelsWorkGroupSize != 0)
globalThreads = (globalThreads / kernelsWorkGroupSize + 1) * kernelsWorkGroupSize;
// Enqueue a kernel run
oclQueue.enqueueNDRangeKernel(kernelsJugCLer, cl::NullRange,
cl::NDRange(globalThreads), cl::NDRange(kernelsWorkGroupSize));
// Read back the result
oclQueue.enqueueReadBuffer(
*(pixelsBuff),
CL_FALSE,
0,
pixelsBuff->getInfo<CL_MEM_SIZE>(),
bitmap->pixels);
oclQueue.finish();
const double t1 = WallClockTime();
// A simple trick to smooth sample/sec value
const double k = 0.01;
frameSec = (1.0 - k) * frameSec + k * (1.0 / (t1 -t0));
}
void RenderEngine::UpdateFilm() {
boost::unique_lock<boost::mutex> lock(engineMutex);
if (started) {
elapsedTime = WallClockTime() - startTime;
UpdateFilmLockLess();
UpdateCounters();
const float haltthreshold = renderConfig->cfg.GetFloat("batch.haltthreshold", -1.f);
if (haltthreshold >= 0.f) {
// Check if it is time to run the convergence test again
const u_int imgWidth = film->GetWidth();
const u_int imgHeight = film->GetHeight();
const u_int pixelCount = imgWidth * imgHeight;
const double now = WallClockTime();
// Do not run the test if we don't have at least 16 new samples per pixel
if ((samplesCount - lastConvergenceTestSamplesCount > pixelCount * 16) &&
((now - lastConvergenceTestTime) * 1000.0 >= renderConfig->GetScreenRefreshInterval())) {
film->UpdateScreenBuffer(); // Required in order to have a valid convergence test
convergence = 1.f - film->RunConvergenceTest() / (float)pixelCount;
lastConvergenceTestTime = now;
lastConvergenceTestSamplesCount = samplesCount;
}
}
}
}
void UpdateRenderingCPU(void) {
double startTime = WallClockTime();
const float invWidth = 1.f / width;
const float invHeight = 1.f / height;
int x, y;
for (y = 0; y < height; y++) { /* Loop over image rows */
for (x = 0; x < width; x++) { /* Loop cols */
const int i = (height - y - 1) * width + x;
const int i2 = 2 * i;
const float r1 = GetRandom(&seeds[i2], &seeds[i2 + 1]) - .5f;
const float r2 = GetRandom(&seeds[i2], &seeds[i2 + 1]) - .5f;
const float kcx = (x + r1) * invWidth - .5f;
const float kcy = (y + r2) * invHeight - .5f;
Vec rdir;
vinit(rdir,
camera.x.x * kcx + camera.y.x * kcy + camera.dir.x,
camera.x.y * kcx + camera.y.y * kcy + camera.dir.y,
camera.x.z * kcx + camera.y.z * kcy + camera.dir.z);
Vec rorig;
vsmul(rorig, 0.1f, rdir);
vadd(rorig, rorig, camera.orig)
vnorm(rdir);
const Ray ray = {rorig, rdir};
Vec r;
RadiancePathTracing(spheres, sphereCount, &ray,
&seeds[i2], &seeds[i2 + 1], &r);
if (currentSample == 0)
colors[i] = r;
else {
const float k1 = currentSample;
const float k2 = 1.f / (k1 + 1.f);
colors[i].x = (colors[i].x * k1 + r.x) * k2;
colors[i].y = (colors[i].y * k1 + r.y) * k2;
colors[i].z = (colors[i].z * k1 + r.z) * k2;
}
pixels[y * width + x] = toInt(colors[i].x) |
(toInt(colors[i].y) << 8) |
(toInt(colors[i].z) << 16);
}
}
const float elapsedTime = WallClockTime() - startTime;
const float sampleSec = height * width / elapsedTime;
sprintf(captionBuffer, "Rendering time %.3f sec (pass %d) Sample/sec %.1fK\n",
elapsedTime, currentSample, sampleSec / 1000.f);
currentSample++;
}
size_t OCLRenderer::DrawFrame() {
const GameConfig &gameConfig(*(gameLevel->gameConfig));
//--------------------------------------------------------------------------
// Recompile the scene
//--------------------------------------------------------------------------
{
//const double t1 = WallClockTime();
boost::unique_lock<boost::mutex> lock(gameLevel->levelMutex);
compiledScene->Recompile(gameLevel->editActionList);
gameLevel->editActionList.Reset();
const float ghostTimeLength = gameConfig.GetRendererGhostFactorTime();
float k;
if (gameLevel->camera->IsChangedSinceLastUpdate()) {
timeSinceLastCameraEdit = WallClockTime();
const double dt = Min<double>(WallClockTime() - timeSinceLastNoCameraEdit, ghostTimeLength);
k = 1.f - dt / ghostTimeLength;
} else {
timeSinceLastNoCameraEdit = WallClockTime();
const double dt = Min<double>(WallClockTime() - timeSinceLastCameraEdit, ghostTimeLength);
k = dt / ghostTimeLength;
}
blendFactor = (1.f - k) * gameConfig.GetRendererGhostFactorCameraEdit() +
k * gameConfig.GetRendererGhostFactorNoCameraEdit();
//SFERA_LOG("Mutex time: " << ((WallClockTime() - t1) * 1000.0));
}
//--------------------------------------------------------------------------
// Render
//--------------------------------------------------------------------------
barrier->wait();
// Other threads do the rendering
barrier->wait();
//--------------------------------------------------------------------------
// Blend frames, tone mapping and copy the OpenCL frame buffer to OpenGL one
//--------------------------------------------------------------------------
renderThread[0]->DrawFrame();
return totSamplePerPass *
gameConfig.GetScreenWidth() *
gameConfig.GetScreenHeight();
}
void Renderer::waitAndSave(const int image_number) {
const double startTime = WallClockTime();
for (;;) {
boost::this_thread::sleep(boost::posix_time::millisec(1000));
const double elapsedTime = WallClockTime() - startTime;
// Print some information about the rendering progress
stats();
if (elapsedTime > renderTime) {
// Time to stop the rendering
break;
}
}
session.GetFilm().SaveOutputs();
renameFile(image_number);
}
void updateCamera(Scene* scene) {
double time = WallClockTime() - sceneTimeOffset;
double rtime = (time / 21.35634); // camera rotation time
rtime = rtime - floor(rtime);
double btime = (time / 59.8752); // camera bobbing time
btime = btime - floor(btime);
double ltime = (time / 98.7654); // light rotation time
ltime = ltime - floor(ltime);
double lbtime = (time / 92.764); // light bobbing time
lbtime = lbtime - floor(lbtime);
camInit(scene->cam);
cl_float3 pos;
cl_float3 center;
vecInit(center, 14.9f, 9.7f, -14.85f);
vecInit(pos, 33.0f * sin(rtime * M_PI * 2.0),
8.0f * sin(btime * M_PI * 2.0 + 0.3),
33.0f * (cos(rtime * M_PI * 2.0) - 0.11));
vecAdd(pos, center);
camMove(scene->cam, pos);
camLookAt(scene->cam, center);
vecInit(pos, 5.0f * sin(-ltime * M_PI * 2.0),
8.0f + 6.0f * sin(lbtime * M_PI * 2.0),
5.0f * cos(ltime * M_PI * 2.0));
vecNormalize(pos);
scene->lightDir = pos;
}
void PathIntegrator::ReInit() {
for (size_t i = 0; i < paths.size(); ++i)
paths[i]->Init(renderEngine, sampler);
firstPath = 0;
sampleBuffer->Reset();
statsRenderingStart = WallClockTime();
statsTotalSampleCount = 0;
}
void UpdateMandel() {
const double startTime = WallClockTime();
const int maxSize = width > height ? width : height;
int s;
for (s = 0; s < width * height / 4; ++s) {
int t;
for (t = 0; t < 4; ++t) {
const int tid = s * 4 + t;
const int i = tid % width;
const int j = tid / width;
const float x0 = ((i * scale) - ((scale / 2) * width)) / maxSize + offsetX;
const float y0 = ((j * scale) - ((scale / 2) * height)) / maxSize + offsetY;
float x = x0;
float y = y0;
float x2 = x * x;
float y2 = y * y;
uint iter = 0;
for (iter = 0; (x2 + y2 <= 4.f) && (iter < maxIterations); ++iter) {
y = 2 * x * y + y0;
x = x2 - y2 + x0;
x2 = x * x;
y2 = y * y;
}
unsigned char *p = (unsigned char *)pixels;
if (iter == maxIterations)
p[tid] = 0;
else
p[tid] = (unsigned char)(255.f * (iter / (float)maxIterations) + 0.5f);
}
}
const double elapsedTime = WallClockTime() - startTime;
const double sampleSec = height * width / elapsedTime;
sprintf(captionBuffer, "Rendering time: %.3f secs (Sample/sec %.1fK Max. Iterations %d)",
elapsedTime, sampleSec / 1000.f, maxIterations);
}
bool RenderSession::NeedPeriodicSave() {
if (periodicSaveEnabled) {
const double now = WallClockTime();
if (now - lastPeriodicSave > periodiceSaveTime) {
lastPeriodicSave = now;
return true;
} else
return false;
} else
return false;
}
void CPU_Worker::Intersect(RayBuffer *rayBuffer) {
// Trace rays
const Ray *rb = rayBuffer->GetRayBuffer();
RayHit *hb = rayBuffer->GetHitBuffer();
double start = WallClockTime();
#ifndef __DEBUG
omp_set_num_threads(config->max_threads);
#pragma omp parallel for schedule(guided)
#endif
for (unsigned int i = 0; i < rayBuffer->GetRayCount(); ++i) {
hb[i].SetMiss();
IntersectRay(&rb[i], &hb[i]);
}
profiler->addRayTracingTime(WallClockTime() - start);
profiler->addRaysTraced(rayBuffer->GetSize());
}
void CUDA_Worker::IntersectGPU(RayBuffer *rayBuffer) {
//const double t1 = WallClockTime();
cudaMemset(hraysBuff, 0, sizeof(RayHit) * rayBuffer->GetRayCount());
cudaMemset(raysBuff, 0, sizeof(Ray) * rayBuffer->GetRayCount());
double start = WallClockTime();
cudaMemcpy(raysBuff, rayBuffer->GetRayBuffer(), sizeof(Ray) * rayBuffer->GetRayCount(),
cudaMemcpyHostToDevice);
intersect_wrapper(raysBuff, hraysBuff, (POINTERFREESCENE::QBVHNode*) d_qbvhBuff,
(POINTERFREESCENE::QuadTriangle*) d_qbvhTrisBuff, rayBuffer->GetRayCount());
cudaMemcpy(rayBuffer->GetHitBuffer(), hraysBuff, sizeof(RayHit) * rayBuffer->GetRayCount(),
cudaMemcpyDeviceToHost);
profiler->addRayTracingTime(WallClockTime() - start);
profiler->addRaysTraced(rayBuffer->GetRayCount());
}
void RenderEngine::EndEdit(const EditActionList &editActions) {
boost::unique_lock<boost::mutex> lock(engineMutex);
assert (started);
assert (editMode);
bool dataSetUpdated;
if (editActions.Has(GEOMETRY_EDIT) ||
((renderConfig->scene->dataSet->GetAcceleratorType() != ACCEL_MQBVH) &&
editActions.Has(INSTANCE_TRANS_EDIT))) {
// To avoid reference to the DataSet de-allocated inside UpdateDataSet()
ctx->SetDataSet(NULL);
// For all other accelerator, I have to rebuild the DataSet
renderConfig->scene->UpdateDataSet(ctx);
// Set the LuxRays SataSet
ctx->SetDataSet(renderConfig->scene->dataSet);
dataSetUpdated = true;
} else
dataSetUpdated = false;
// Restart all intersection devices
ctx->Start();
if (!dataSetUpdated &&
(renderConfig->scene->dataSet->GetAcceleratorType() == ACCEL_MQBVH) &&
editActions.Has(INSTANCE_TRANS_EDIT)) {
// Update the DataSet
ctx->UpdateDataSet();
}
samplesCount = 0;
elapsedTime = 0.0f;
startTime = WallClockTime();
film->ResetConvergenceTest();
convergence = 0.f;
lastConvergenceTestTime = startTime;
lastConvergenceTestSamplesCount = 0;
editMode = false;
EndEditLockLess(editActions);
}
PathRenderEngine::PathRenderEngine(SLGScene *scn, Film *flm, boost::mutex *filmMutex,
vector<IntersectionDevice *> intersectionDev, const bool onlySpecular,
const Properties &cfg) : RenderEngine(scn, flm, filmMutex) {
intersectionDevices = intersectionDev;
samplePerPixel = max(1, cfg.GetInt("path.sampler.spp", cfg.GetInt("sampler.spp", 4)));
lowLatency = (cfg.GetInt("opencl.latency.mode", 1) != 0);
maxPathDepth = cfg.GetInt("path.maxdepth", 3);
int strat = cfg.GetInt("path.lightstrategy", 0);
if (strat == 0)
lightStrategy = ONE_UNIFORM;
else
lightStrategy = ALL_UNIFORM;
shadowRayCount = cfg.GetInt("path.shadowrays", 1);
onlySampleSpecular = onlySpecular;
// Russian Roulette parameters
int rrStrat = cfg.GetInt("path.russianroulette.strategy", 1);
if (rrStrat == 0)
rrStrategy = PROBABILITY;
else
rrStrategy = IMPORTANCE;
rrDepth = cfg.GetInt("path.russianroulette.depth", 2);
rrProb = cfg.GetFloat("path.russianroulette.prob", 0.5f);
rrImportanceCap = cfg.GetFloat("path.russianroulette.cap", 0.125f);
const unsigned long seedBase = (unsigned long)(WallClockTime() / 1000.0);
// Create and start render threads
const size_t renderThreadCount = intersectionDevices.size();
cerr << "Starting "<< renderThreadCount << " Path render threads" << endl;
for (size_t i = 0; i < renderThreadCount; ++i) {
if (intersectionDevices[i]->GetType() == DEVICE_TYPE_NATIVE_THREAD) {
PathNativeRenderThread *t = new PathNativeRenderThread(i, seedBase, i / (float)renderThreadCount,
samplePerPixel, (NativeThreadIntersectionDevice *)intersectionDevices[i], this);
renderThreads.push_back(t);
} else {
PathDeviceRenderThread *t = new PathDeviceRenderThread(i, seedBase, i / (float)renderThreadCount,
samplePerPixel, intersectionDevices[i], this);
renderThreads.push_back(t);
}
}
}
static void PrintCaptions() {
glEnable(GL_BLEND);
glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
glColor4f(0.f, 0.f, 0.f, 0.8f);
glRecti(0, engine->height - 15, engine->width - 1, engine->height - 1);
glRecti(0, 0, engine->width - 1, 18);
glDisable(GL_BLEND);
// Title
glColor3f(1.f, 1.f, 1.f);
glRasterPos2i(4, engine->height - 10);
PrintString(GLUT_BITMAP_8_BY_13, engine->SPPMG_LABEL);
// Stats
glRasterPos2i(4, 5);
char captionBuffer[512];
const double elapsedTime = WallClockTime() - engine->startTime;
const unsigned int kPhotonsSec = engine->getPhotonTracedTotal() / (elapsedTime * 1000.f);
sprintf(captionBuffer, "[Photons %.2fM][Avg. photons/sec % 4dK][Elapsed time %dsecs]",
float(engine->getPhotonTracedTotal() / 1000000.0), kPhotonsSec, int(elapsedTime));
PrintString(GLUT_BITMAP_8_BY_13, captionBuffer);
}
void RenderEngine::Start() {
boost::unique_lock<boost::mutex> lock(engineMutex);
assert (!started);
started = true;
const float epsilonMin = renderConfig->cfg.GetFloat("scene.epsilon.min", DEFAULT_EPSILON_MIN);
MachineEpsilon::SetMin(epsilonMin);
const float epsilonMax = renderConfig->cfg.GetFloat("scene.epsilon.max", DEFAULT_EPSILON_MAX);
MachineEpsilon::SetMax(epsilonMax);
ctx->Start();
StartLockLess();
samplesCount = 0;
elapsedTime = 0.0f;
startTime = WallClockTime();
film->ResetConvergenceTest();
convergence = 0.f;
lastConvergenceTestTime = startTime;
lastConvergenceTestSamplesCount = 0;
}
void PathOCLRenderThread::InitKernels() {
//--------------------------------------------------------------------------
// Compile kernels
//--------------------------------------------------------------------------
CompiledScene *cscene = renderEngine->compiledScene;
cl::Context &oclContext = intersectionDevice->GetOpenCLContext();
cl::Device &oclDevice = intersectionDevice->GetOpenCLDevice();
// Set #define symbols
stringstream ss;
ss.precision(6);
ss << scientific <<
" -D PARAM_TASK_COUNT=" << renderEngine->taskCount <<
" -D PARAM_IMAGE_WIDTH=" << renderEngine->film->GetWidth() <<
" -D PARAM_IMAGE_HEIGHT=" << renderEngine->film->GetHeight() <<
" -D PARAM_RAY_EPSILON=" << renderEngine->epsilon << "f" <<
" -D PARAM_SEED=" << seed <<
" -D PARAM_MAX_PATH_DEPTH=" << renderEngine->maxPathDepth <<
" -D PARAM_MAX_DIFFUSE_PATH_VERTEX_COUNT=" << renderEngine->maxDiffusePathVertexCount <<
" -D PARAM_RR_DEPTH=" << renderEngine->rrDepth <<
" -D PARAM_RR_CAP=" << renderEngine->rrImportanceCap << "f"
;
switch (renderEngine->renderConfig->scene->dataSet->GetAcceleratorType()) {
case ACCEL_BVH:
ss << " -D PARAM_ACCEL_BVH";
break;
case ACCEL_QBVH:
ss << " -D PARAM_ACCEL_QBVH";
break;
case ACCEL_MQBVH:
ss << " -D PARAM_ACCEL_MQBVH";
break;
default:
assert (false);
}
if (cscene->enable_MAT_MATTE)
ss << " -D PARAM_ENABLE_MAT_MATTE";
if (cscene->enable_MAT_AREALIGHT)
ss << " -D PARAM_ENABLE_MAT_AREALIGHT";
if (cscene->enable_MAT_MIRROR)
ss << " -D PARAM_ENABLE_MAT_MIRROR";
if (cscene->enable_MAT_GLASS)
ss << " -D PARAM_ENABLE_MAT_GLASS";
if (cscene->enable_MAT_MATTEMIRROR)
ss << " -D PARAM_ENABLE_MAT_MATTEMIRROR";
if (cscene->enable_MAT_METAL)
ss << " -D PARAM_ENABLE_MAT_METAL";
if (cscene->enable_MAT_MATTEMETAL)
ss << " -D PARAM_ENABLE_MAT_MATTEMETAL";
if (cscene->enable_MAT_ALLOY)
ss << " -D PARAM_ENABLE_MAT_ALLOY";
if (cscene->enable_MAT_ARCHGLASS)
ss << " -D PARAM_ENABLE_MAT_ARCHGLASS";
if (cscene->camera.lensRadius > 0.f)
ss << " -D PARAM_CAMERA_HAS_DOF";
if (infiniteLightBuff)
ss << " -D PARAM_HAS_INFINITELIGHT";
if (skyLightBuff)
ss << " -D PARAM_HAS_SKYLIGHT";
if (sunLightBuff) {
ss << " -D PARAM_HAS_SUNLIGHT";
if (!areaLightsBuff) {
ss <<
" -D PARAM_DIRECT_LIGHT_SAMPLING" <<
" -D PARAM_DL_LIGHT_COUNT=0"
;
}
}
if (areaLightsBuff) {
ss <<
" -D PARAM_DIRECT_LIGHT_SAMPLING" <<
" -D PARAM_DL_LIGHT_COUNT=" << renderEngine->compiledScene->areaLights.size()
;
}
if (texMapRGBBuff || texMapAlphaBuff)
ss << " -D PARAM_HAS_TEXTUREMAPS";
if (texMapAlphaBuff)
ss << " -D PARAM_HAS_ALPHA_TEXTUREMAPS";
if (meshBumpsBuff)
ss << " -D PARAM_HAS_BUMPMAPS";
if (meshNormalMapsBuff)
ss << " -D PARAM_HAS_NORMALMAPS";
const PathOCL::Filter *filter = renderEngine->filter;
switch (filter->type) {
case PathOCL::NONE:
ss << " -D PARAM_IMAGE_FILTER_TYPE=0";
break;
case PathOCL::BOX:
ss << " -D PARAM_IMAGE_FILTER_TYPE=1" <<
" -D PARAM_IMAGE_FILTER_WIDTH_X=" << filter->widthX << "f" <<
" -D PARAM_IMAGE_FILTER_WIDTH_Y=" << filter->widthY << "f";
break;
case PathOCL::GAUSSIAN:
ss << " -D PARAM_IMAGE_FILTER_TYPE=2" <<
" -D PARAM_IMAGE_FILTER_WIDTH_X=" << filter->widthX << "f" <<
" -D PARAM_IMAGE_FILTER_WIDTH_Y=" << filter->widthY << "f" <<
" -D PARAM_IMAGE_FILTER_GAUSSIAN_ALPHA=" << ((PathOCL::GaussianFilter *)filter)->alpha << "f";
break;
case PathOCL::MITCHELL:
ss << " -D PARAM_IMAGE_FILTER_TYPE=3" <<
" -D PARAM_IMAGE_FILTER_WIDTH_X=" << filter->widthX << "f" <<
" -D PARAM_IMAGE_FILTER_WIDTH_Y=" << filter->widthY << "f" <<
" -D PARAM_IMAGE_FILTER_MITCHELL_B=" << ((PathOCL::MitchellFilter *)filter)->B << "f" <<
" -D PARAM_IMAGE_FILTER_MITCHELL_C=" << ((PathOCL::MitchellFilter *)filter)->C << "f";
break;
default:
assert (false);
}
if (renderEngine->usePixelAtomics)
ss << " -D PARAM_USE_PIXEL_ATOMICS";
const PathOCL::Sampler *sampler = renderEngine->sampler;
switch (sampler->type) {
case PathOCL::INLINED_RANDOM:
ss << " -D PARAM_SAMPLER_TYPE=0";
break;
case PathOCL::RANDOM:
ss << " -D PARAM_SAMPLER_TYPE=1";
break;
case PathOCL::METROPOLIS:
ss << " -D PARAM_SAMPLER_TYPE=2" <<
" -D PARAM_SAMPLER_METROPOLIS_LARGE_STEP_RATE=" << ((PathOCL::MetropolisSampler *)sampler)->largeStepRate << "f" <<
" -D PARAM_SAMPLER_METROPOLIS_MAX_CONSECUTIVE_REJECT=" << ((PathOCL::MetropolisSampler *)sampler)->maxConsecutiveReject <<
" -D PARAM_SAMPLER_METROPOLIS_IMAGE_MUTATION_RANGE=" << ((PathOCL::MetropolisSampler *)sampler)->imageMutationRate << "f";
break;
case PathOCL::STRATIFIED:
ss << " -D PARAM_SAMPLER_TYPE=3" <<
" -D PARAM_SAMPLER_STRATIFIED_X_SAMPLES=" << ((PathOCL::StratifiedSampler *)sampler)->xSamples <<
" -D PARAM_SAMPLER_STRATIFIED_Y_SAMPLES=" << ((PathOCL::StratifiedSampler *)sampler)->ySamples;
break;
default:
assert (false);
}
// Check the OpenCL vendor and use some specific compiler options
#if defined(__APPLE__) // OSX version detection
{
struct utsname retval;
uname(&retval);
if(retval.release[0] == '1' && retval.release[1] < '1') // result < darwin 11
ss << " -D __APPLE_FIX__";
}
#endif
//--------------------------------------------------------------------------
const double tStart = WallClockTime();
// Check if I have to recompile the kernels
string newKernelParameters = ss.str();
if (kernelsParameters != newKernelParameters) {
kernelsParameters = newKernelParameters;
// Compile sources
stringstream ssKernel;
ssKernel <<
_LUXRAYS_UV_OCLDEFINE
_LUXRAYS_SPECTRUM_OCLDEFINE
_LUXRAYS_POINT_OCLDEFINE
_LUXRAYS_VECTOR_OCLDEFINE
_LUXRAYS_TRIANGLE_OCLDEFINE
_LUXRAYS_RAY_OCLDEFINE
_LUXRAYS_RAYHIT_OCLDEFINE <<
KernelSource_PathOCL_kernel_datatypes <<
KernelSource_PathOCL_kernel_core <<
KernelSource_PathOCL_kernel_filters <<
KernelSource_PathOCL_kernel_scene <<
KernelSource_PathOCL_kernel_samplers <<
KernelSource_PathOCL_kernels;
string kernelSource = ssKernel.str();
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Defined symbols: " << kernelsParameters);
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Compiling kernels ");
bool cached;
cl::STRING_CLASS error;
cl::Program *program = kernelCache->Compile(oclContext, oclDevice,
kernelsParameters, kernelSource,
&cached, &error);
if (!program) {
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] PathOCL kernel compilation error" << std::endl << error);
throw std::runtime_error("PathOCL kernel compilation error");
}
if (cached) {
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Kernels cached");
} else {
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Kernels not cached");
}
//----------------------------------------------------------------------
// Init kernel
//----------------------------------------------------------------------
delete initKernel;
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Compiling Init Kernel");
initKernel = new cl::Kernel(*program, "Init");
initKernel->getWorkGroupInfo<size_t>(oclDevice, CL_KERNEL_WORK_GROUP_SIZE, &initWorkGroupSize);
if (intersectionDevice->GetForceWorkGroupSize() > 0)
initWorkGroupSize = intersectionDevice->GetForceWorkGroupSize();
else if (renderEngine->sampler->type == PathOCL::STRATIFIED) {
// Resize the workgroup to have enough local memory
size_t localMem = oclDevice.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>();
while ((initWorkGroupSize > 64) && (stratifiedDataSize * initWorkGroupSize > localMem))
initWorkGroupSize /= 2;
if (stratifiedDataSize * initWorkGroupSize > localMem)
throw std::runtime_error("Not enough local memory to run, try to reduce path.sampler.xsamples and path.sampler.xsamples values");
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Cap work group size to: " << initWorkGroupSize);
}
//--------------------------------------------------------------------------
// InitFB kernel
//--------------------------------------------------------------------------
delete initFBKernel;
initFBKernel = new cl::Kernel(*program, "InitFrameBuffer");
initFBKernel->getWorkGroupInfo<size_t>(oclDevice, CL_KERNEL_WORK_GROUP_SIZE, &initFBWorkGroupSize);
if (intersectionDevice->GetForceWorkGroupSize() > 0)
initFBWorkGroupSize = intersectionDevice->GetForceWorkGroupSize();
//----------------------------------------------------------------------
// Sampler kernel
//----------------------------------------------------------------------
delete samplerKernel;
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Compiling Sampler Kernel");
samplerKernel = new cl::Kernel(*program, "Sampler");
samplerKernel->getWorkGroupInfo<size_t>(oclDevice, CL_KERNEL_WORK_GROUP_SIZE, &samplerWorkGroupSize);
if (intersectionDevice->GetForceWorkGroupSize() > 0)
samplerWorkGroupSize = intersectionDevice->GetForceWorkGroupSize();
else if (renderEngine->sampler->type == PathOCL::STRATIFIED) {
// Resize the workgroup to have enough local memory
size_t localMem = oclDevice.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>();
while ((samplerWorkGroupSize > 64) && (stratifiedDataSize * samplerWorkGroupSize > localMem))
samplerWorkGroupSize /= 2;
if (stratifiedDataSize * samplerWorkGroupSize > localMem)
throw std::runtime_error("Not enough local memory to run, try to reduce path.sampler.xsamples and path.sampler.xsamples values");
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Cap work group size to: " << samplerWorkGroupSize);
}
//----------------------------------------------------------------------
// AdvancePaths kernel
//----------------------------------------------------------------------
delete advancePathsKernel;
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Compiling AdvancePaths Kernel");
advancePathsKernel = new cl::Kernel(*program, "AdvancePaths");
advancePathsKernel->getWorkGroupInfo<size_t>(oclDevice, CL_KERNEL_WORK_GROUP_SIZE, &advancePathsWorkGroupSize);
if (intersectionDevice->GetForceWorkGroupSize() > 0)
advancePathsWorkGroupSize = intersectionDevice->GetForceWorkGroupSize();
//----------------------------------------------------------------------
const double tEnd = WallClockTime();
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Kernels compilation time: " << int((tEnd - tStart) * 1000.0) << "ms");
delete program;
} else
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Using cached kernels");
}
void PathOCLRenderThread::RenderThreadImpl(PathOCLRenderThread *renderThread) {
//LM_LOG_ENGINE("[PathOCLRenderThread::" << renderThread->threadIndex << "] Rendering thread started");
cl::CommandQueue &oclQueue = renderThread->intersectionDevice->GetOpenCLQueue();
const unsigned int taskCount = renderThread->renderEngine->taskCount;
oclQueue.finish();
// Wait for the signal to start the rendering
renderThread->renderEngine->renderStartBarrier->wait();
try {
double startTime = WallClockTime();
while (!boost::this_thread::interruption_requested()) {
/*if(renderThread->threadIndex == 0)
cerr<< "[DEBUG] =================================");*/
// Async. transfer of the frame buffer
oclQueue.enqueueReadBuffer(
*(renderThread->frameBufferBuff),
CL_FALSE,
0,
renderThread->frameBufferBuff->getInfo<CL_MEM_SIZE>(),
renderThread->frameBuffer);
// Async. transfer of GPU task statistics
oclQueue.enqueueReadBuffer(
*(renderThread->taskStatsBuff),
CL_FALSE,
0,
sizeof(PathOCL::GPUTaskStats) * taskCount,
renderThread->gpuTaskStats);
for (;;) {
cl::Event event;
// Decide how many kernels to enqueue
const unsigned int screenRefreshInterval = renderThread->renderEngine->renderConfig->GetScreenRefreshInterval();
unsigned int iterations;
if (screenRefreshInterval <= 100)
iterations = 1;
else if (screenRefreshInterval <= 500)
iterations = 2;
else if (screenRefreshInterval <= 1000)
iterations = 4;
else
iterations = 8;
for (unsigned int i = 0; i < iterations; ++i) {
// Generate the samples and paths
if (i == 0)
oclQueue.enqueueNDRangeKernel(*(renderThread->samplerKernel), cl::NullRange,
cl::NDRange(taskCount), cl::NDRange(renderThread->samplerWorkGroupSize),
NULL, &event);
else
oclQueue.enqueueNDRangeKernel(*(renderThread->samplerKernel), cl::NullRange,
cl::NDRange(taskCount), cl::NDRange(renderThread->samplerWorkGroupSize));
// Trace rays
renderThread->intersectionDevice->EnqueueTraceRayBuffer(*(renderThread->raysBuff),
*(renderThread->hitsBuff), taskCount, NULL, NULL);
// Advance to next path state
oclQueue.enqueueNDRangeKernel(*(renderThread->advancePathsKernel), cl::NullRange,
cl::NDRange(taskCount), cl::NDRange(renderThread->advancePathsWorkGroupSize));
}
oclQueue.flush();
event.wait();
const double elapsedTime = WallClockTime() - startTime;
/*if(renderThread->threadIndex == 0)
cerr<< "[DEBUG] Elapsed time: " << elapsedTime * 1000.0 <<
"ms (screenRefreshInterval: " << renderThread->renderEngine->screenRefreshInterval << ")");*/
if ((elapsedTime * 1000.0 > (double)screenRefreshInterval) ||
boost::this_thread::interruption_requested())
break;
}
startTime = WallClockTime();
}
//LM_LOG_ENGINE("[PathOCLRenderThread::" << renderThread->threadIndex << "] Rendering thread halted");
} catch (boost::thread_interrupted) {
LM_LOG_ENGINE("[PathOCLRenderThread::" << renderThread->threadIndex << "] Rendering thread halted");
} catch (cl::Error err) {
LM_LOG_ENGINE("[PathOCLRenderThread::" << renderThread->threadIndex << "] Rendering thread ERROR: " << err.what() <<
"(" << luxrays::utils::oclErrorString(err.err()) << ")");
}
oclQueue.enqueueReadBuffer(
*(renderThread->frameBufferBuff),
CL_TRUE,
0,
renderThread->frameBufferBuff->getInfo<CL_MEM_SIZE>(),
renderThread->frameBuffer);
}
u_int64_t CPU_Worker::AdvancePhotonPath(u_int64_t photonTarget) {
uint todoPhotonCount = 0;
PhotonPath* livePhotonPaths = new PhotonPath[rayBuffer->GetSize()];
rayBuffer->Reset();
size_t initc = min((int) rayBuffer->GetSize(), (int) photonTarget);
double start = WallClockTime();
for (size_t i = 0; i < initc; ++i) {
int p = rayBuffer->ReserveRay();
Ray * b = &(rayBuffer->GetRayBuffer())[p];
engine->InitPhotonPath(engine->ss, &livePhotonPaths[i], b, seedBuffer[i]);
}
while (todoPhotonCount < photonTarget) {
Intersect(rayBuffer);
#ifndef __DEBUG
omp_set_num_threads(config->max_threads);
#pragma omp parallel for schedule(guided)
#endif
for (unsigned int i = 0; i < rayBuffer->GetRayCount(); ++i) {
PhotonPath *photonPath = &livePhotonPaths[i];
Ray *ray = &rayBuffer->GetRayBuffer()[i];
RayHit *rayHit = &rayBuffer->GetHitBuffer()[i];
if (photonPath->done == true) {
continue;
}
if (rayHit->Miss()) {
photonPath->done = true;
} else { // Something was hit
Point hitPoint;
Spectrum surfaceColor;
Normal N, shadeN;
if (engine->GetHitPointInformation(engine->ss, ray, rayHit, hitPoint, surfaceColor,
N, shadeN))
continue;
const unsigned int currentTriangleIndex = rayHit->index;
const unsigned int currentMeshIndex = engine->ss->meshIDs[currentTriangleIndex];
POINTERFREESCENE::Material *hitPointMat =
&engine->ss->materials[engine->ss->meshMats[currentMeshIndex]];
uint matType = hitPointMat->type;
if (matType == MAT_AREALIGHT) {
photonPath->done = true;
} else {
float fPdf;
Vector wi;
Vector wo = -ray->d;
bool specularBounce = true;
float u0 = getFloatRNG(seedBuffer[i]);
float u1 = getFloatRNG(seedBuffer[i]);
float u2 = getFloatRNG(seedBuffer[i]);
Spectrum f;
switch (matType) {
case MAT_MATTE:
engine->ss->Matte_Sample_f(&hitPointMat->param.matte, &wo, &wi, &fPdf, &f,
&shadeN, u0, u1, &specularBounce);
f *= surfaceColor;
break;
case MAT_MIRROR:
engine->ss->Mirror_Sample_f(&hitPointMat->param.mirror, &wo, &wi, &fPdf,
&f, &shadeN, &specularBounce);
f *= surfaceColor;
break;
case MAT_GLASS:
engine->ss->Glass_Sample_f(&hitPointMat->param.glass, &wo, &wi, &fPdf, &f,
&N, &shadeN, u0, &specularBounce);
f *= surfaceColor;
break;
case MAT_MATTEMIRROR:
engine->ss->MatteMirror_Sample_f(&hitPointMat->param.matteMirror, &wo, &wi,
&fPdf, &f, &shadeN, u0, u1, u2, &specularBounce);
f *= surfaceColor;
break;
case MAT_METAL:
engine->ss->Metal_Sample_f(&hitPointMat->param.metal, &wo, &wi, &fPdf, &f,
&shadeN, u0, u1, &specularBounce);
f *= surfaceColor;
break;
case MAT_MATTEMETAL:
engine->ss->MatteMetal_Sample_f(&hitPointMat->param.matteMetal, &wo, &wi,
&fPdf, &f, &shadeN, u0, u1, u2, &specularBounce);
f *= surfaceColor;
break;
case MAT_ALLOY:
engine->ss->Alloy_Sample_f(&hitPointMat->param.alloy, &wo, &wi, &fPdf, &f,
&shadeN, u0, u1, u2, &specularBounce);
f *= surfaceColor;
break;
case MAT_ARCHGLASS:
engine->ss->ArchGlass_Sample_f(&hitPointMat->param.archGlass, &wo, &wi,
&fPdf, &f, &N, &shadeN, u0, &specularBounce);
f *= surfaceColor;
break;
case MAT_NULL:
wi = ray->d;
specularBounce = 1;
fPdf = 1.f;
break;
default:
// Huston, we have a problem...
specularBounce = 1;
fPdf = 0.f;
break;
}
if (!specularBounce) // if difuse
lookupA->AddFlux(engine->ss, engine->alpha, hitPoint, shadeN, -ray->d,
photonPath->flux, currentPhotonRadius2);
if (photonPath->depth < MAX_PHOTON_PATH_DEPTH) {
// Build the next vertex path ray
if ((fPdf <= 0.f) || f.Black()) {
photonPath->done = true;
} else {
photonPath->depth++;
photonPath->flux *= f / fPdf;
// Russian Roulette
const float p = 0.75f;
if (photonPath->depth < 3) {
*ray = Ray(hitPoint, wi);
} else if (getFloatRNG(seedBuffer[i]) < p) {
photonPath->flux /= p;
*ray = Ray(hitPoint, wi);
} else {
photonPath->done = true;
}
}
} else {
photonPath->done = true;
}
}
}
}
uint oldc = rayBuffer->GetRayCount();
rayBuffer->Reset();
for (unsigned int i = 0; i < oldc; ++i) {
PhotonPath *photonPath = &livePhotonPaths[i];
Ray *ray = &rayBuffer->GetRayBuffer()[i];
if (photonPath->done && todoPhotonCount < photonTarget) {
todoPhotonCount++;
Ray n;
engine->InitPhotonPath(engine->ss, photonPath, &n, seedBuffer[i]);
livePhotonPaths[i].done = false;
size_t p = rayBuffer->AddRay(n);
livePhotonPaths[p] = *photonPath;
} else if (!photonPath->done) {
rayBuffer->AddRay(*ray);
}
}
}
// float MPhotonsSec = todoPhotonCount / ((WallClockTime()-start) * 1000000.f);
//printf("\nRate: %.3f MPhotons/sec\n",MPhotonsSec);
profiler->addPhotonTracingTime(WallClockTime() - start);
profiler->addPhotonsTraced(todoPhotonCount);
rayBuffer->Reset();
return todoPhotonCount;
}
void Worker::ProcessIterations(PPM* engine) {
u_int64_t photonPerIteration = engine->photonsFirstIteration;
uint iterationCount;
resetRayBuffer();
UpdateBBox();
LookupSetHitPoints(hitPointsStaticInfo_iterationCopy, hitPoints_iterationCopy);
uint iter = 0;
double previousIterTime = WallClockTime();
fprintf(stdout, "iteration, photons_iter, photons_total, photons_sec, total_time, radius, device\n");
while (iter < config->max_iters) {
++iter;
double start = WallClockTime();
iterationCount = engine->IncIteration();
if (engine->GetIterationNumber() > MAX_ITERATIONS) {
break;
}
photonPerIteration = engine->photonsFirstIteration;
#if defined USE_SPPMPA || defined USE_SPPM
BuildHitPoints(iterationCount);
UpdateBBox();
#endif
#if defined USE_SPPM || defined USE_PPM
if (iterationCount == 1)
InitRadius(iterationCount);
#else
InitRadius(iterationCount);
#endif
updateDeviceHitPoints();
ReHash(currentPhotonRadius2);//argument ignored in non-PA
updateDeviceLookupAcc();
photonPerIteration = AdvancePhotonPath(photonPerIteration);
getDeviceHitpoints();
#if defined USE_PPM
AccumulateFluxPPM(iterationCount, photonPerIteration);
#endif
#if defined USE_SPPM
AccumulateFluxSPPM(iterationCount, photonPerIteration);
#endif
#if defined USE_SPPMPA
AccumulateFluxSPPMPA(iterationCount, photonPerIteration);
#endif
#if defined USE_PPMPA
AccumulateFluxPPMPA(iterationCount, photonPerIteration);
#endif
UpdateSampleFrameBuffer(photonPerIteration);
/**
* iteration lock required in PhotonTracedTotal
*/
engine->incPhotonTracedTotal(photonPerIteration);
//PushHitPoints();
profiler->additeratingTime(WallClockTime() - start);
profiler->addIteration(1);
if (profiler->iterationCount % 100 == 0)
profiler->printStats(deviceID);
// if (iterationCount % 50 == 0)
// engine->SaveImpl(to_string<uint> (iterationCount, std::dec) + engine->fileName);
#if defined USE_SPPM || defined USE_PPM
const float radius = hitPoints_iterationCopy[0].accumPhotonRadius2;
#else
const float radius = currentPhotonRadius2;
#endif
const double time = WallClockTime();
const double totalTime = time - engine->startTime;
const double iterTime = time - previousIterTime;
// const float itsec = engine->GetIterationNumber() / totalTime;
const uint photonTotal = engine->getPhotonTracedTotal();
const float photonSec = photonTotal / (totalTime * 1000.f);
fprintf(stdout, "%d, %lu, %u, %f, %f, %f, %f, %d\n", iterationCount, photonPerIteration, photonTotal, photonSec, iterTime, totalTime, radius, getDeviceID());
previousIterTime = time;
}
}
static int BatchMode(double stopTime, unsigned int stopSPP) {
const double startTime = WallClockTime();
#if !defined(LUXRAYS_DISABLE_OPENCL)
double lastFilmUpdate = WallClockTime();
#endif
double sampleSec = 0.0;
char buf[512];
const vector<IntersectionDevice *> interscetionDevices = config->GetIntersectionDevices();
for (;;) {
boost::this_thread::sleep(boost::posix_time::millisec(1000));
const double now = WallClockTime();
const double elapsedTime =now - startTime;
const unsigned int pass = config->GetRenderEngine()->GetPass();
if ((stopTime > 0) && (elapsedTime >= stopTime))
break;
if ((stopSPP > 0) && (pass >= stopSPP))
break;
// Check if periodic save is enabled
if (config->NeedPeriodicSave()) {
#if !defined(LUXRAYS_DISABLE_OPENCL)
if (config->GetRenderEngine()->GetEngineType() == PATHGPU) {
// I need to update the Film
PathGPURenderEngine *pre = (PathGPURenderEngine *)config->GetRenderEngine();
pre->UpdateFilm();
}
#endif
// Time to save the image and film
config->SaveFilmImage();
}
// Print some information about the rendering progress
double raysSec = 0.0;
for (size_t i = 0; i < interscetionDevices.size(); ++i)
raysSec += interscetionDevices[i]->GetPerformance();
switch (config->GetRenderEngine()->GetEngineType()) {
case DIRECTLIGHT:
case PATH: {
sampleSec = config->film->GetAvgSampleSec();
sprintf(buf, "[Elapsed time: %3d/%dsec][Samples %4d/%d][Avg. samples/sec % 4dK][Avg. rays/sec % 4dK on %.1fK tris]",
int(elapsedTime), int(stopTime), pass, stopSPP, int(sampleSec/ 1000.0),
int(raysSec / 1000.0), config->scene->dataSet->GetTotalTriangleCount() / 1000.0);
break;
}
case SPPM: {
SPPMRenderEngine *sre = (SPPMRenderEngine *)config->GetRenderEngine();
sprintf(buf, "[Elapsed time: %3d/%dsec][Pass %3d][Photon %.1fM][Avg. photon/sec % 4dK][Avg. rays/sec % 4dK on %.1fK tris]",
int(elapsedTime), int(stopTime), pass, sre->GetTotalPhotonCount() / 1000000.0, int(sre->GetTotalPhotonSec() / 1000.0),
int(raysSec / 1000.0), config->scene->dataSet->GetTotalTriangleCount() / 1000.0);
break;
}
#if !defined(LUXRAYS_DISABLE_OPENCL)
case PATHGPU: {
PathGPURenderEngine *pre = (PathGPURenderEngine *)config->GetRenderEngine();
sampleSec = pre->GetTotalSamplesSec();
sprintf(buf, "[Elapsed time: %3d/%dsec][Samples %4d/%d][Avg. samples/sec % 3.2fM][Avg. rays/sec % 4dK on %.1fK tris]",
int(elapsedTime), int(stopTime), pass, stopSPP, sampleSec / 1000000.0,
int(raysSec / 1000.0), config->scene->dataSet->GetTotalTriangleCount() / 1000.0);
if (WallClockTime() - lastFilmUpdate > 5.0) {
pre->UpdateFilm();
lastFilmUpdate = WallClockTime();
}
break;
}
#endif
default:
assert (false);
}
cerr << buf << endl;
}
// Stop the rendering
config->StopAllRenderThreads();
// Save the rendered image
config->SaveFilmImage();
sprintf(buf, "LuxMark index: %.3f", sampleSec / 1000000.0);
cerr << buf << endl;
delete config;
cerr << "Done." << endl;
return EXIT_SUCCESS;
}
void Worker::ProcessIterations(PPM* engine) {
u_int64_t photonPerIteration = engine->photonsFirstIteration;
uint iterationCount;
resetRayBuffer();
UpdateBBox();
while (!boost::this_thread::interruption_requested()) {
double start = WallClockTime();
if (engine->GetIterationNumber() > config->max_iters) {
break;
}
iterationCount = engine->IncIteration();
photonPerIteration = engine->photonsFirstIteration;
// #if defined USE_SPPMPA || defined USE_SPPM
BuildHitPoints(iterationCount);
UpdateBBox();
// #endif
// #if defined USE_SPPM || defined USE_PPM
// if (iterationCount == 1)
// InitRadius(iterationCount);
// #else
InitRadius(iterationCount);
// #endif
updateDeviceHitPoints();
#ifndef REBUILD_HASH
if (iterationCount == 1)
#endif
ReHash(currentPhotonRadius2);//argument ignored in non-PA
#ifndef REBUILD_HASH
if (iterationCount == 1)
#endif
updateDeviceLookupAcc();
photonPerIteration = AdvancePhotonPath(photonPerIteration);
// #if defined USE_PPM
// getDeviceHitpoints();
// AccumulateFluxPPM(iterationCount, photonPerIteration);
// #endif
// #if defined USE_SPPM
// AccumulateFluxSPPM(iterationCount, photonPerIteration);
// #endif
// #if defined USE_SPPMPA
AccumulateFluxSPPMPA(iterationCount, photonPerIteration);
// #endif
// #if defined USE_PPMPA
// AccumulateFluxPPMPA(iterationCount, photonPerIteration);
// getDeviceHitpoints();
// #endif
UpdateSampleFrameBuffer(photonPerIteration);
/**
* iteration lock required in PhotonTracedTotal
*/
engine->incPhotonTracedTotal(photonPerIteration);
profiler->additeratingTime(WallClockTime() - start);
profiler->addIteration(1);
//if (profiler->iterationCount % 20 == 0)
// profiler->printStats(deviceID);
//if (iterationCount % 100 == 0)
// engine->SaveImpl(to_string<uint> (iterationCount, std::dec) + engine->fileName);
printf("iteration %d finished\n", iterationCount);
}
//profiler->printStats(deviceID);
}
Scene::Scene(const std::string &fileName,uint width, uint height ,const int aType ) {
accelType = aType;
extMeshCache = new ExtMeshCache();
texMapCache = new TextureMapCache();
SDL_LOG("Reading scene: " << fileName);
scnProp = new Properties(fileName);
//--------------------------------------------------------------------------
// Read camera position and target
//--------------------------------------------------------------------------
std::vector<float> vf = GetParameters(*scnProp, "scene.camera.lookat", 6, "10.0 0.0 0.0 0.0 0.0 0.0");
Point o(vf.at(0), vf.at(1), vf.at(2));
Point t(vf.at(3), vf.at(4), vf.at(5));
SDL_LOG("Camera postion: " << o);
SDL_LOG("Camera target: " << t);
vf = GetParameters(*scnProp, "scene.camera.up", 3, "0.0 0.0 0.1");
const Vector up(vf.at(0), vf.at(1), vf.at(2));
camera = new PerspectiveCamera(o, t, up);
camera->lensRadius = scnProp->GetFloat("scene.camera.lensradius", 0.f);
camera->focalDistance = scnProp->GetFloat("scene.camera.focaldistance", 10.f);
camera->fieldOfView = scnProp->GetFloat("scene.camera.fieldofview", 45.f);
// Check if camera motion blur is enabled
if (scnProp->GetInt("scene.camera.motionblur.enable", 0)) {
camera->motionBlur = true;
vf = GetParameters(*scnProp, "scene.camera.motionblur.lookat", 6, "10.0 1.0 0.0 0.0 1.0 0.0");
camera->mbOrig = Point(vf.at(0), vf.at(1), vf.at(2));
camera->mbTarget = Point(vf.at(3), vf.at(4), vf.at(5));
vf = GetParameters(*scnProp, "scene.camera.motionblur.up", 3, "0.0 0.0 0.1");
camera->mbUp = Vector(vf.at(0), vf.at(1), vf.at(2));
}
//--------------------------------------------------------------------------
// Read all materials
//--------------------------------------------------------------------------
std::vector<std::string> matKeys = scnProp->GetAllKeys("scene.materials.");
if (matKeys.size() == 0)
throw std::runtime_error("No material definition found");
for (std::vector<std::string>::const_iterator matKey = matKeys.begin(); matKey != matKeys.end(); ++matKey) {
const std::string &key = *matKey;
const std::string matType = Properties::ExtractField(key, 2);
if (matType == "")
throw std::runtime_error("Syntax error in " + key);
const std::string matName = Properties::ExtractField(key, 3);
if (matName == "")
throw std::runtime_error("Syntax error in " + key);
SDL_LOG("Material definition: " << matName << " [" << matType << "]");
Material *mat = CreateMaterial(key, *scnProp);
materialIndices[matName] = materials.size();
materials.push_back(mat);
}
//--------------------------------------------------------------------------
// Read all objects .ply file
//--------------------------------------------------------------------------
std::vector<std::string> objKeys = scnProp->GetAllKeys("scene.objects.");
if (objKeys.size() == 0)
throw std::runtime_error("Unable to find object definitions");
double lastPrint = WallClockTime();
unsigned int objCount = 0;
for (std::vector<std::string>::const_iterator objKey = objKeys.begin(); objKey != objKeys.end(); ++objKey) {
const std::string &key = *objKey;
// Check if it is the root of the definition of an object otherwise skip
const size_t dot1 = key.find(".", std::string("scene.objects.").length());
if (dot1 == std::string::npos)
continue;
const size_t dot2 = key.find(".", dot1 + 1);
if (dot2 != std::string::npos)
continue;
const std::string objName = Properties::ExtractField(key, 3);
if (objName == "")
throw std::runtime_error("Syntax error in " + key);
// Build the object
const std::vector<std::string> args = scnProp->GetStringVector(key, "");
const std::string plyFileName = args.at(0);
const double now = WallClockTime();
if (now - lastPrint > 2.0) {
SDL_LOG("PLY object count: " << objCount);
lastPrint = now;
}
++objCount;
//SDL_LOG("PLY object [" << objName << "] file name: " << plyFileName);
// Check if I have to calculate normal or not
const bool usePlyNormals = (scnProp->GetInt(key + ".useplynormals", 0) != 0);
// Check if I have to use an instance mesh or not
ExtMesh *meshObject;
if (scnProp->IsDefined(key + ".transformation")) {
const std::vector<float> vf = GetParameters(*scnProp, key + ".transformation", 16, "1.0 0.0 0.0 0.0 0.0 1.0 0.0 0.0 0.0 0.0 1.0 0.0 0.0 0.0 0.0 1.0");
const Matrix4x4 mat(
vf.at(0), vf.at(4), vf.at(8), vf.at(12),
vf.at(1), vf.at(5), vf.at(9), vf.at(13),
vf.at(2), vf.at(6), vf.at(10), vf.at(14),
vf.at(3), vf.at(7), vf.at(11), vf.at(15));
const Transform trans(mat);
meshObject = extMeshCache->GetExtMesh(plyFileName, usePlyNormals, trans);
} else
meshObject = extMeshCache->GetExtMesh(plyFileName, usePlyNormals);
objectIndices[objName] = objects.size();
objects.push_back(meshObject);
// Get the material
const std::string matName = Properties::ExtractField(key, 2);
if (matName == "")
throw std::runtime_error("Syntax error in material name: " + matName);
if (materialIndices.count(matName) < 1)
throw std::runtime_error("Unknown material: " + matName);
Material *mat = materials[materialIndices[matName]];
// Check if it is a light sources
if (mat->IsLightSource()) {
SDL_LOG("The " << objName << " object is a light sources with " << meshObject->GetTotalTriangleCount() << " triangles");
AreaLightMaterial *light = (AreaLightMaterial *)mat;
objectMaterials.push_back(mat);
for (unsigned int i = 0; i < meshObject->GetTotalTriangleCount(); ++i) {
TriangleLight *tl = new TriangleLight(light, static_cast<unsigned int>(objects.size()) - 1, i, objects);
lights.push_back(tl);
}
} else {
SurfaceMaterial *surfMat = (SurfaceMaterial *)mat;
objectMaterials.push_back(surfMat);
}
// [old deprecated syntax] Check if there is a texture map associated to the object
if (args.size() > 1) {
// Check if the object has UV coords
if (!meshObject->HasUVs())
throw std::runtime_error("PLY object " + plyFileName + " is missing UV coordinates for texture mapping");
TexMapInstance *tm = texMapCache->GetTexMapInstance(args.at(1), 2.2f);
objectTexMaps.push_back(tm);
objectBumpMaps.push_back(NULL);
objectNormalMaps.push_back(NULL);
} else {
// Check for if there is a texture map associated to the object with the new syntax
const std::string texMap = scnProp->GetString(key + ".texmap", "");
if (texMap != "") {
// Check if the object has UV coords
if (!meshObject->HasUVs())
throw std::runtime_error("PLY object " + plyFileName + " is missing UV coordinates for texture mapping");
const float gamma = scnProp->GetFloat(key + ".texmap.gamma", 2.2f);
TexMapInstance *tm = texMapCache->GetTexMapInstance(texMap, gamma);
objectTexMaps.push_back(tm);
} else
objectTexMaps.push_back(NULL);
/**
* Check if there is an alpha map associated to the object
* If there is, the map is added to a previously added texturemap.
* If no texture map (diffuse map) is detected, a black texture
* is created and the alpha map is added to it. --PC
*/
const std::string alphaMap = scnProp->GetString(key + ".alphamap", "");
if (alphaMap != "") {
// Got an alpha map, retrieve the textureMap and add the alpha channel to it.
const std::string texMap = scnProp->GetString(key + ".texmap", "");
const float gamma = scnProp->GetFloat(key + ".texmap.gamma", 2.2f);
TextureMap *tm;
if (!(tm = texMapCache->FindTextureMap(texMap, gamma))) {
SDL_LOG("Alpha map " << alphaMap << " is for a materials without texture. A black texture has been created for support!");
// We have an alpha map without a diffuse texture. In this case we need to create
// a texture map filled with black
tm = new TextureMap(alphaMap, gamma, 1.0, 1.0, 1.0);
tm->AddAlpha(alphaMap);
TexMapInstance *tmi = texMapCache->AddTextureMap(alphaMap, tm);
// Remove the NULL inserted above, when no texmap was found. Without doing this the whole thing will not work
objectTexMaps.pop_back();
// Add the new texture to the chain
objectTexMaps.push_back(tmi);
} else {
// Add an alpha map to the pre-existing diffuse texture
tm->AddAlpha(alphaMap);
}
}
// Check for if there is a bump map associated to the object
const std::string bumpMap = scnProp->GetString(key + ".bumpmap", "");
if (bumpMap != "") {
// Check if the object has UV coords
if (!meshObject->HasUVs())
throw std::runtime_error("PLY object " + plyFileName + " is missing UV coordinates for bump mapping");
const float scale = scnProp->GetFloat(key + ".bumpmap.scale", 1.f);
BumpMapInstance *bm = texMapCache->GetBumpMapInstance(bumpMap, scale);
objectBumpMaps.push_back(bm);
} else
objectBumpMaps.push_back(NULL);
// Check for if there is a normal map associated to the object
const std::string normalMap = scnProp->GetString(key + ".normalmap", "");
if (normalMap != "") {
// Check if the object has UV coords
if (!meshObject->HasUVs())
throw std::runtime_error("PLY object " + plyFileName + " is missing UV coordinates for normal mapping");
NormalMapInstance *nm = texMapCache->GetNormalMapInstance(normalMap);
objectNormalMaps.push_back(nm);
} else
objectNormalMaps.push_back(NULL);
}
}
SDL_LOG("PLY object count: " << objCount);
//--------------------------------------------------------------------------
// Check if there is an infinitelight source defined
//--------------------------------------------------------------------------
const std::vector<std::string> ilParams = scnProp->GetStringVector("scene.infinitelight.file", "");
if (ilParams.size() > 0) {
const float gamma = scnProp->GetFloat("scene.infinitelight.gamma", 2.2f);
TexMapInstance *tex = texMapCache->GetTexMapInstance(ilParams.at(0), gamma);
// Check if I have to use InfiniteLightBF method
if (scnProp->GetInt("scene.infinitelight.usebruteforce", 0)) {
SDL_LOG("Using brute force infinite light sampling");
infiniteLight = new InfiniteLightBF(tex);
useInfiniteLightBruteForce = true;
} else {
if (ilParams.size() == 2)
infiniteLight = new InfiniteLightPortal(tex, ilParams.at(1));
else
infiniteLight = new InfiniteLightIS(tex);
// Add the infinite light to the list of light sources
lights.push_back(infiniteLight);
useInfiniteLightBruteForce = false;
}
std::vector<float> vf = GetParameters(*scnProp, "scene.infinitelight.gain", 3, "1.0 1.0 1.0");
infiniteLight->SetGain(Spectrum(vf.at(0), vf.at(1), vf.at(2)));
vf = GetParameters(*scnProp, "scene.infinitelight.shift", 2, "0.0 0.0");
infiniteLight->SetShift(vf.at(0), vf.at(1));
infiniteLight->Preprocess();
} else {
infiniteLight = NULL;
useInfiniteLightBruteForce = false;
}
//--------------------------------------------------------------------------
// Check if there is a SkyLight defined
//--------------------------------------------------------------------------
const std::vector<std::string> silParams = scnProp->GetStringVector("scene.skylight.dir", "");
if (silParams.size() > 0) {
if (infiniteLight)
throw std::runtime_error("Can not define a skylight when there is already an infinitelight defined");
std::vector<float> sdir = GetParameters(*scnProp, "scene.skylight.dir", 3, "0.0 0.0 1.0");
const float turb = scnProp->GetFloat("scene.skylight.turbidity", 2.2f);
std::vector<float> gain = GetParameters(*scnProp, "scene.skylight.gain", 3, "1.0 1.0 1.0");
SkyLight *sl = new SkyLight(turb, Vector(sdir.at(0), sdir.at(1), sdir.at(2)));
infiniteLight = sl;
sl->SetGain(Spectrum(gain.at(0), gain.at(1), gain.at(2)));
sl->Init();
useInfiniteLightBruteForce = true;
}
//--------------------------------------------------------------------------
// Check if there is a SunLight defined
//--------------------------------------------------------------------------
const std::vector<std::string> sulParams = scnProp->GetStringVector("scene.sunlight.dir", "");
if (sulParams.size() > 0) {
std::vector<float> sdir = GetParameters(*scnProp, "scene.sunlight.dir", 3, "0.0 0.0 1.0");
const float turb = scnProp->GetFloat("scene.sunlight.turbidity", 2.2f);
const float relSize = scnProp->GetFloat("scene.sunlight.relsize", 1.0f);
std::vector<float> gain = GetParameters(*scnProp, "scene.sunlight.gain", 3, "1.0 1.0 1.0");
SunLight *sunLight = new SunLight(turb, relSize, Vector(sdir.at(0), sdir.at(1), sdir.at(2)));
sunLight->SetGain(Spectrum(gain.at(0), gain.at(1), gain.at(2)));
sunLight->Init();
lights.push_back(sunLight);
}
//--------------------------------------------------------------------------
camera->Update(width, height);
}
void PathOCLRenderThread::InitRender() {
Scene *scene = renderEngine->renderConfig->scene;
cl::Context &oclContext = intersectionDevice->GetOpenCLContext();
cl::Device &oclDevice = intersectionDevice->GetOpenCLDevice();
const OpenCLDeviceDescription *deviceDesc = intersectionDevice->GetDeviceDesc();
double tStart, tEnd;
//--------------------------------------------------------------------------
// FrameBuffer definition
//--------------------------------------------------------------------------
InitFrameBuffer();
//--------------------------------------------------------------------------
// Camera definition
//--------------------------------------------------------------------------
InitCamera();
//--------------------------------------------------------------------------
// Scene geometry
//--------------------------------------------------------------------------
InitGeometry();
//--------------------------------------------------------------------------
// Translate material definitions
//--------------------------------------------------------------------------
InitMaterials();
//--------------------------------------------------------------------------
// Translate area lights
//--------------------------------------------------------------------------
InitAreaLights();
//--------------------------------------------------------------------------
// Check if there is an infinite light source
//--------------------------------------------------------------------------
InitInfiniteLight();
//--------------------------------------------------------------------------
// Check if there is an sun light source
//--------------------------------------------------------------------------
InitSunLight();
//--------------------------------------------------------------------------
// Check if there is an sky light source
//--------------------------------------------------------------------------
InitSkyLight();
const unsigned int areaLightCount = renderEngine->compiledScene->areaLights.size();
if (!skyLightBuff && !sunLightBuff && !infiniteLightBuff && (areaLightCount == 0))
throw runtime_error("There are no light sources supported by PathOCL in the scene");
//--------------------------------------------------------------------------
// Translate mesh texture maps
//--------------------------------------------------------------------------
InitTextureMaps();
//--------------------------------------------------------------------------
// Allocate Ray/RayHit buffers
//--------------------------------------------------------------------------
const unsigned int taskCount = renderEngine->taskCount;
tStart = WallClockTime();
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Ray buffer size: " << (sizeof(Ray) * taskCount / 1024) << "Kbytes");
raysBuff = new cl::Buffer(oclContext,
CL_MEM_READ_WRITE,
sizeof(Ray) * taskCount);
deviceDesc->AllocMemory(raysBuff->getInfo<CL_MEM_SIZE>());
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] RayHit buffer size: " << (sizeof(RayHit) * taskCount / 1024) << "Kbytes");
hitsBuff = new cl::Buffer(oclContext,
CL_MEM_READ_WRITE,
sizeof(RayHit) * taskCount);
deviceDesc->AllocMemory(hitsBuff->getInfo<CL_MEM_SIZE>());
tEnd = WallClockTime();
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] OpenCL buffer creation time: " << int((tEnd - tStart) * 1000.0) << "ms");
//--------------------------------------------------------------------------
// Allocate GPU task buffers
//--------------------------------------------------------------------------
// TODO: clenup all this mess
const size_t gpuTaksSizePart1 =
// Seed size
sizeof(PathOCL::Seed);
const size_t uDataEyePathVertexSize =
// IDX_SCREEN_X, IDX_SCREEN_Y
sizeof(float) * 2 +
// IDX_DOF_X, IDX_DOF_Y
((scene->camera->lensRadius > 0.f) ? (sizeof(float) * 2) : 0);
const size_t uDataPerPathVertexSize =
// IDX_TEX_ALPHA,
((texMapAlphaBuff) ? sizeof(float) : 0) +
// IDX_BSDF_X, IDX_BSDF_Y, IDX_BSDF_Z
sizeof(float) * 3 +
// IDX_DIRECTLIGHT_X, IDX_DIRECTLIGHT_Y, IDX_DIRECTLIGHT_Z
(((areaLightCount > 0) || sunLightBuff) ? (sizeof(float) * 3) : 0) +
// IDX_RR
sizeof(float);
const size_t uDataSize = (renderEngine->sampler->type == PathOCL::INLINED_RANDOM) ?
// Only IDX_SCREEN_X, IDX_SCREEN_Y
(sizeof(float) * 2) :
((renderEngine->sampler->type == PathOCL::METROPOLIS) ?
(sizeof(float) * 2 + sizeof(unsigned int) * 5 + sizeof(Spectrum) + 2 * (uDataEyePathVertexSize + uDataPerPathVertexSize * renderEngine->maxPathDepth)) :
(uDataEyePathVertexSize + uDataPerPathVertexSize * renderEngine->maxPathDepth));
size_t sampleSize =
// uint pixelIndex;
((renderEngine->sampler->type == PathOCL::METROPOLIS) ? 0 : sizeof(unsigned int)) +
uDataSize +
// Spectrum radiance;
sizeof(Spectrum);
stratifiedDataSize = 0;
if (renderEngine->sampler->type == PathOCL::STRATIFIED) {
PathOCL::StratifiedSampler *s = (PathOCL::StratifiedSampler *)renderEngine->sampler;
stratifiedDataSize =
// stratifiedScreen2D
sizeof(float) * s->xSamples * s->ySamples * 2 +
// stratifiedDof2D
((scene->camera->lensRadius > 0.f) ? (sizeof(float) * s->xSamples * s->ySamples * 2) : 0) +
// stratifiedAlpha1D
((texMapAlphaBuff) ? (sizeof(float) * s->xSamples) : 0) +
// stratifiedBSDF2D
sizeof(float) * s->xSamples * s->ySamples * 2 +
// stratifiedBSDF1D
sizeof(float) * s->xSamples +
// stratifiedLight2D
// stratifiedLight1D
(((areaLightCount > 0) || sunLightBuff) ? (sizeof(float) * s->xSamples * s->ySamples * 2 + sizeof(float) * s->xSamples) : 0);
sampleSize += stratifiedDataSize;
}
const size_t gpuTaksSizePart2 = sampleSize;
const size_t gpuTaksSizePart3 =
// PathState size
((((areaLightCount > 0) || sunLightBuff) ? sizeof(PathOCL::PathStateDL) : sizeof(PathOCL::PathState)) +
//unsigned int diffuseVertexCount;
((renderEngine->maxDiffusePathVertexCount < renderEngine->maxPathDepth) ? sizeof(unsigned int) : 0));
const size_t gpuTaksSize = gpuTaksSizePart1 + gpuTaksSizePart2 + gpuTaksSizePart3;
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Size of a GPUTask: " << gpuTaksSize <<
"bytes (" << gpuTaksSizePart1 << " + " << gpuTaksSizePart2 << " + " << gpuTaksSizePart3 << ")");
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Tasks buffer size: " << (gpuTaksSize * taskCount / 1024) << "Kbytes");
// Check if the task buffer is too big
if (oclDevice.getInfo<CL_DEVICE_MAX_MEM_ALLOC_SIZE>() < gpuTaksSize * taskCount) {
stringstream ss;
ss << "The GPUTask buffer is too big for this device (i.e. CL_DEVICE_MAX_MEM_ALLOC_SIZE=" <<
oclDevice.getInfo<CL_DEVICE_MAX_MEM_ALLOC_SIZE>() <<
"): try to reduce opencl.task.count and/or path.maxdepth and/or to change Sampler";
throw std::runtime_error(ss.str());
}
tasksBuff = new cl::Buffer(oclContext,
CL_MEM_READ_WRITE,
gpuTaksSize * taskCount);
deviceDesc->AllocMemory(tasksBuff->getInfo<CL_MEM_SIZE>());
//--------------------------------------------------------------------------
// Allocate GPU task statistic buffers
//--------------------------------------------------------------------------
LM_LOG_ENGINE("[PathOCLRenderThread::" << threadIndex << "] Task Stats buffer size: " << (sizeof(PathOCL::GPUTaskStats) * taskCount / 1024) << "Kbytes");
taskStatsBuff = new cl::Buffer(oclContext,
CL_MEM_READ_WRITE,
sizeof(PathOCL::GPUTaskStats) * taskCount);
deviceDesc->AllocMemory(taskStatsBuff->getInfo<CL_MEM_SIZE>());
//--------------------------------------------------------------------------
// Compile kernels
//--------------------------------------------------------------------------
InitKernels();
//--------------------------------------------------------------------------
// Initialize
//--------------------------------------------------------------------------
// Set kernel arguments
SetKernelArgs();
cl::CommandQueue &oclQueue = intersectionDevice->GetOpenCLQueue();
// Clear the frame buffer
oclQueue.enqueueNDRangeKernel(*initFBKernel, cl::NullRange,
cl::NDRange(RoundUp<unsigned int>(frameBufferPixelCount, initFBWorkGroupSize)),
cl::NDRange(initFBWorkGroupSize));
// Initialize the tasks buffer
oclQueue.enqueueNDRangeKernel(*initKernel, cl::NullRange,
cl::NDRange(taskCount), cl::NDRange(initWorkGroupSize));
oclQueue.finish();
// Reset statistics in order to be more accurate
intersectionDevice->ResetPerformaceStats();
}
static int BatchSimpleMode(const double haltTime, const unsigned int haltSpp, const float haltThreshold) {
RenderConfig *config = session->renderConfig;
RenderEngine *engine = session->renderEngine;
// Force the film update at 2.5secs (mostly used by PathOCL)
config->SetScreenRefreshInterval(2500);
// Start the rendering
session->Start();
// Nothing to do if it is the FileSaverRenderEngine
if (!dynamic_cast<FileSaverRenderEngine *>(engine)) {
const double startTime = WallClockTime();
double lastFilmUpdate = WallClockTime();
char buf[512];
for (;;) {
boost::this_thread::sleep(boost::posix_time::millisec(1000));
// Check if periodic save is enabled
if (session->NeedPeriodicSave()) {
// Time to save the image and film
session->SaveFilmImage();
lastFilmUpdate = WallClockTime();
} else {
// Film update may be required by some render engine to
// update statistics, convergence test and more
if (WallClockTime() - lastFilmUpdate > 5.0) {
session->renderEngine->UpdateFilm();
lastFilmUpdate = WallClockTime();
}
}
const double now = WallClockTime();
const double elapsedTime = now - startTime;
if ((haltTime > 0) && (elapsedTime >= haltTime))
break;
const unsigned int pass = engine->GetPass();
if ((haltSpp > 0) && (pass >= haltSpp))
break;
// Convergence test is update inside UpdateFilm()
const float convergence = engine->GetConvergence();
if ((haltThreshold >= 0.f) && (1.f - convergence <= haltThreshold))
break;
// Print some information about the rendering progress
sprintf(buf, "[Elapsed time: %3d/%dsec][Samples %4d/%d][Convergence %f%%][Avg. samples/sec % 3.2fM on %.1fK tris]",
int(elapsedTime), int(haltTime), pass, haltSpp, 100.f * convergence, engine->GetTotalSamplesSec() / 1000000.0,
config->scene->dataSet->GetTotalTriangleCount() / 1000.0);
SLG_LOG(buf);
}
}
// Stop the rendering
session->Stop();
// Save the rendered image
session->SaveFilmImage();
delete session;
SLG_LOG("Done.");
return EXIT_SUCCESS;
}
static void motionFunc(int x, int y) {
const double minInterval = 0.2;
if (mouseButton0) {
// Check elapsed time since last update
if (WallClockTime() - lastMouseUpdate > minInterval) {
const int distX = x - mouseGrabLastX;
const int distY = y - mouseGrabLastY;
session->BeginEdit();
if (mouseGrabMode) {
session->renderConfig->scene->camera->RotateUp(0.04f * distY * ROTATE_STEP);
session->renderConfig->scene->camera->RotateLeft(0.04f * distX * ROTATE_STEP);
}
else {
session->renderConfig->scene->camera->RotateDown(0.04f * distY * ROTATE_STEP);
session->renderConfig->scene->camera->RotateRight(0.04f * distX * ROTATE_STEP);
};
session->renderConfig->scene->camera->Update(
session->film->GetWidth(), session->film->GetHeight());
session->editActions.AddAction(CAMERA_EDIT);
session->EndEdit();
mouseGrabLastX = x;
mouseGrabLastY = y;
displayFunc();
lastMouseUpdate = WallClockTime();
}
} else if (mouseButton2) {
// Check elapsed time since last update
if (WallClockTime() - lastMouseUpdate > minInterval) {
const int distX = x - mouseGrabLastX;
const int distY = y - mouseGrabLastY;
session->BeginEdit();
if (mouseGrabMode) {
session->renderConfig->scene->camera->TranslateLeft(0.04f * distX * MOVE_STEP);
session->renderConfig->scene->camera->TranslateForward(0.04f * distY * MOVE_STEP);
}
else {
session->renderConfig->scene->camera->TranslateRight(0.04f * distX * MOVE_STEP);
session->renderConfig->scene->camera->TranslateBackward(0.04f * distY * MOVE_STEP);
}
session->renderConfig->scene->camera->Update(
session->film->GetWidth(), session->film->GetHeight());
session->editActions.AddAction(CAMERA_EDIT);
session->EndEdit();
mouseGrabLastX = x;
mouseGrabLastY = y;
displayFunc();
lastMouseUpdate = WallClockTime();
}
}
}
void HashGrid::RefreshMutex() {
const unsigned int hitPointsCount = hitPoints->GetSize();
const BBox &hpBBox = hitPoints->GetBBox();
// Calculate the size of the grid cell
const float maxPhotonRadius2 = hitPoints->GetMaxPhotonRaidus2();
const float cellSize = sqrtf(maxPhotonRadius2) * 2.f;
std::cerr << "Hash grid cell size: " << cellSize << std::endl;
invCellSize = 1.f / cellSize;
// TODO: add a tunable parameter for hashgrid size
gridSize = hitPointsCount;
if (!grid) {
grid = new std::list<HitPoint *>*[gridSize];
for (unsigned int i = 0; i < gridSize; ++i)
grid[i] = NULL;
} else {
for (unsigned int i = 0; i < gridSize; ++i) {
delete grid[i];
grid[i] = NULL;
}
}
std::cerr << "Building hit points hash grid:" << std::endl;
std::cerr << " 0k/" << hitPointsCount / 1000 << "k" <<std::endl;
//unsigned int maxPathCount = 0;
double lastPrintTime = WallClockTime();
unsigned long long entryCount = 0;
for (unsigned int i = 0; i < hitPointsCount; ++i) {
if (WallClockTime() - lastPrintTime > 2.0) {
std::cerr << " " << i / 1000 << "k/" << hitPointsCount / 1000 << "k" <<std::endl;
lastPrintTime = WallClockTime();
}
HitPoint *hp = hitPoints->GetHitPoint(i);
if (hp->type == SURFACE) {
const float photonRadius = sqrtf(hp->accumPhotonRadius2);
const Vector rad(photonRadius, photonRadius, photonRadius);
const Vector bMin = ((hp->position - rad) - hpBBox.pMin) * invCellSize;
const Vector bMax = ((hp->position + rad) - hpBBox.pMin) * invCellSize;
for (int iz = abs(int(bMin.z)); iz <= abs(int(bMax.z)); iz++) {
for (int iy = abs(int(bMin.y)); iy <= abs(int(bMax.y)); iy++) {
for (int ix = abs(int(bMin.x)); ix <= abs(int(bMax.x)); ix++) {
int hv = Hash(ix, iy, iz);
if (grid[hv] == NULL)
grid[hv] = new std::list<HitPoint *>();
grid[hv]->push_front(hp);
++entryCount;
/*// grid[hv]->size() is very slow to execute
if (grid[hv]->size() > maxPathCount)
maxPathCount = grid[hv]->size();*/
}
}
}
}
}
//std::cerr << "Max. hit points in a single hash grid entry: " << maxPathCount << std::endl;
std::cerr << "Total hash grid entry: " << entryCount << std::endl;
std::cerr << "Avg. hit points in a single hash grid entry: " << entryCount / gridSize << std::endl;
// HashGrid debug code
/*for (unsigned int i = 0; i < hashGridSize; ++i) {
if (grid[i]) {
if (grid[i]->size() > 10) {
std::cerr << "HashGrid[" << i << "].size() = " <<grid[i]->size() << std::endl;
}
}
}*/
}
void PointerFreeHashGrid::ReHash(float currentPhotonRadius2) {
#else
void PointerFreeHashGrid::ReHash(float /*currentPhotonRadius2*/) {
#endif
const unsigned int hitPointsCount = engine->hitPointTotal;
const BBox &hpBBox = hitPointsbbox;
// Calculate the size of the grid cell
#if defined USE_SPPMPA || defined USE_PPMPA
float maxPhotonRadius2 = currentPhotonRadius2;
#else
float maxPhotonRadius2 = 0.f;
for (unsigned int i = 0; i < hitPointsCount; ++i) {
HitPointStaticInfo *ihp = &workerHitPointsInfo[i];
HitPoint *hp = &workerHitPoints[i];
if (ihp->type == SURFACE)
maxPhotonRadius2 = Max(maxPhotonRadius2, hp->accumPhotonRadius2);
}
#endif
const float cellSize = sqrtf(maxPhotonRadius2) * 2.f;
//std::cerr << "Hash grid cell size: " << cellSize << std::endl;
invCellSize = 1.f / cellSize;
// TODO: add a tunable parameter for hashgrid size
//hashGridSize = hitPointsCount;
if (!hashGrid) {
hashGrid = new std::list<uint>*[hashGridSize];
for (unsigned int i = 0; i < hashGridSize; ++i)
hashGrid[i] = NULL;
} else {
for (unsigned int i = 0; i < hashGridSize; ++i) {
delete hashGrid[i];
hashGrid[i] = NULL;
}
}
//std::cerr << "Building hit points hash grid:" << std::endl;
//std::cerr << " 0k/" << hitPointsCount / 1000 << "k" << std::endl;
//unsigned int maxPathCount = 0;
double lastPrintTime = WallClockTime();
unsigned long long entryCount = 0;
for (unsigned int i = 0; i < hitPointsCount; ++i) {
if (WallClockTime() - lastPrintTime > 2.0) {
std::cerr << " " << i / 1000 << "k/" << hitPointsCount / 1000 << "k" << std::endl;
lastPrintTime = WallClockTime();
}
//HitPointInfo *hp = engine->GetHitPointInfo(i);
HitPointStaticInfo *hp = &workerHitPointsInfo[i];
if (hp->type == SURFACE) {
#if defined USE_SPPMPA || defined USE_PPMPA
const float photonRadius = sqrtf(currentPhotonRadius2);
#else
HitPoint *hpp = &workerHitPoints[i];
const float photonRadius = sqrtf(hpp->accumPhotonRadius2);
#endif
const Vector rad(photonRadius, photonRadius, photonRadius);
const Vector bMin = ((hp->position - rad) - hpBBox.pMin) * invCellSize;
const Vector bMax = ((hp->position + rad) - hpBBox.pMin) * invCellSize;
for (int iz = abs(int(bMin.z)); iz <= abs(int(bMax.z)); iz++) {
for (int iy = abs(int(bMin.y)); iy <= abs(int(bMax.y)); iy++) {
for (int ix = abs(int(bMin.x)); ix <= abs(int(bMax.x)); ix++) {
int hv = Hash(ix, iy, iz);
//if (hv == engine->hitPointTotal - 1)
if (hashGrid[hv] == NULL)
hashGrid[hv] = new std::list<uint>();
hashGrid[hv]->push_front(i);
++entryCount;
/*// hashGrid[hv]->size() is very slow to execute
if (hashGrid[hv]->size() > maxPathCount)
maxPathCount = hashGrid[hv]->size();*/
}
}
}
}
}
hashGridEntryCount = entryCount;
//std::cerr << "Max. hit points in a single hash grid entry: " << maxPathCount << std::endl;
std::cerr << "Total hash grid entry: " << entryCount << std::endl;
std::cerr << "Avg. hit points in a single hash grid entry: " << entryCount
/ hashGridSize << std::endl;
//printf("Sizeof %d\n", sizeof(HitPoint*));
// HashGrid debug code
/*for (unsigned int i = 0; i < hashGridSize; ++i) {
if (hashGrid[i]) {
if (hashGrid[i]->size() > 10) {
std::cerr << "HashGrid[" << i << "].size() = " <<hashGrid[i]->size() << std::endl;
}
}
}*/
}
void PointerFreeHashGrid::updateLookupTable() {
if (hashGridLists)
delete[] hashGridLists;
hashGridLists = new uint[hashGridEntryCount];
if (hashGridLenghts)
delete[] hashGridLenghts;
hashGridLenghts = new uint[engine->hitPointTotal];
if (hashGridListsIndex)
delete[] hashGridListsIndex;
hashGridListsIndex = new uint[engine->hitPointTotal];
uint listIndex = 0;
for (unsigned int i = 0; i < engine->hitPointTotal; ++i) {
std::list<uint> *hps = hashGrid[i];
hashGridListsIndex[i] = listIndex;
if (hps) {
hashGridLenghts[i] = hps->size();
std::list<uint>::iterator iter = hps->begin();
while (iter != hps->end()) {
hashGridLists[listIndex++] = *iter++;
}
} else {
hashGridLenghts[i] = 0;
}
}
}
int main(int argc, char *argv[]) {
// load configurations
config = new Config("Options", argc, argv);
srand(1000);
float alpha = 0.7f;
uint width;
uint height;
uint superSampling;
unsigned long long photonsFirstIteration;
//alpha = alpha;
width = config->width;//640;
height = config->height;//480;
superSampling = config->spp;//1;
photonsFirstIteration = 1 << config->photons_first_iter_exp;//0.5M;
#if defined USE_SPPM || defined USE_SPPMPA
superSampling=1;
#endif
size_t hitPointTotal = width * height * superSampling * superSampling;
uint ndvices = 0;
#ifdef GPU0
ndvices++;
#endif
#ifdef GPU2
ndvices++;
#endif
#ifdef CPU
ndvices++;
#endif
engine = new PPM(alpha, width, height, superSampling, photonsFirstIteration, ndvices);
std::string sceneFileName = config->scene_file.c_str(); //"scenes/kitchen/kitchen.scn";
engine->fileName = config->output_file;//"kitchen.png";
// std::string sceneFileName = "scenes/alloy/alloy.scn";
// std::string sceneFileName = "scenes/bigmonkey/bigmonkey.scn";
// std::string sceneFileName = "scenes/psor-cube/psor-cube.scn";
// std::string sceneFileName = "scenes/classroom/classroom.scn";
// std::string sceneFileName = "scenes/luxball/luxball.scn";
// std::string sceneFileName = "scenes/cornell/cornell.scn";
//std::string sceneFileName = "scenes/simple/simple.scn";
// std::string sceneFileName = "scenes/simple-mat/simple-mat.scn";
// std::string sceneFileName = "scenes/sky/sky.scn";
// std::string sceneFileName = "scenes/studiotest/studiotest.scn";
engine->ss = new PointerFreeScene(width, height, sceneFileName);
engine->startTime = WallClockTime();
Seed* seedBuffer;
uint devID;
uint c;
bool build_hit = false;
#ifdef GPU0
devID = 0;
size_t WORK_BUCKET_SIZE_GPU0 = SM * FACTOR * BLOCKSIZE; // SMs*FACTOR*THEADSBLOCK
c = max(hitPointTotal, WORK_BUCKET_SIZE_GPU0);
seedBuffer = new Seed[c];
for (uint i = 0; i < c; i++)
seedBuffer[i] = mwc(i+devID);
build_hit = true;
CUDA_Worker* gpuWorker0 = new CUDA_Worker(0, engine->ss, WORK_BUCKET_SIZE_GPU0, seedBuffer,
build_hit);
#endif
#ifdef GPU2
devID = 2;
size_t WORK_BUCKET_SIZE_GPU2 = SM * FACTOR * BLOCKSIZE; // SMs*FACTOR*THEADSBLOCK
c = max(hitPointTotal, WORK_BUCKET_SIZE_GPU2);
seedBuffer = new Seed[c];
for (uint i = 0; i < c; i++)
seedBuffer[i] = mwc(i+devID);
if (build_hit)
build_hit = false;
else
build_hit = true;
CUDA_Worker* gpuWorker2 = new CUDA_Worker(2, engine->ss, WORK_BUCKET_SIZE_GPU2, seedBuffer,
build_hit);
build_hit = true;
#endif
#ifdef CPU
devID = 100;
size_t WORK_BUCKET_SIZE_CPU = 1024 * 256;
c = max((int)hitPointTotal, (int)WORK_BUCKET_SIZE_CPU);
seedBuffer = new Seed[c];
for (uint i = 0; i < c; i++) {
seedBuffer[i] = mwc(i+devID);
}
if (build_hit)
build_hit = false;
else
build_hit = true;
CPU_Worker* cpuWorker = new CPU_Worker(devID, engine->ss, WORK_BUCKET_SIZE_CPU, seedBuffer,
build_hit);
#endif
if (config->use_display)
engine->draw_thread = new boost::thread(boost::bind(Draw, argc, argv));
#ifdef GPU0
gpuWorker0->thread->join();
#endif
#ifdef GPU2
gpuWorker2->thread->join();
#endif
#ifdef CPU
cpuWorker->thread->join();
#endif
const double elapsedTime = WallClockTime() - engine->startTime;
// float MPhotonsSec = engine->getPhotonTracedTotal() / (elapsedTime * 1000000.f);
const float itsec = engine->GetIterationNumber() / elapsedTime;
if (config->use_display) {
cout << "Done. waiting for display" << endl;
engine->draw_thread->join();
}
engine->SaveImpl(config->output_file.c_str());
fprintf(stderr, "Avg. %.3f iteration/sec\n", itsec);
fprintf(stderr, "Total photons: %.2fM\n", engine->getPhotonTracedTotal() / 1000000.f);
fprintf(stderr, "Total time:\n%f\n", elapsedTime);
fflush(stdout);
fflush(stderr);
return EXIT_SUCCESS;
}
static int BatchTileMode(const unsigned int haltSpp, const float haltThreshold) {
RenderConfig *config = session->renderConfig;
RenderEngine *engine = session->renderEngine;
// batch.halttime condition doesn't make sense in the context
// of tile rendering
if (config->cfg.IsDefined("batch.halttime") && (config->cfg.GetInt("batch.halttime", 0) > 0))
throw runtime_error("batch.halttime parameter can not be used with batch.tile");
// image.subregion condition doesn't make sense in the context
// of tile rendering
if (config->cfg.IsDefined("image.subregion"))
throw runtime_error("image.subregion parameter can not be used with batch.tile");
// Force the film update at 2.5secs (mostly used by PathOCL)
config->SetScreenRefreshInterval(2500);
const u_int filterBorder = 2;
// Original film size
const u_int originalFilmWidth = session->film->GetWidth();
const u_int originalFilmHeight = session->film->GetHeight();
// Allocate a film where to merge all tiles
Film mergeTileFilm(originalFilmWidth, originalFilmWidth,
session->film->HasPerPixelNormalizedBuffer(), session->film->HasPerScreenNormalizedBuffer(),
true);
mergeTileFilm.CopyDynamicSettings(*(session->film));
mergeTileFilm.EnableOverlappedScreenBufferUpdate(false);
mergeTileFilm.Init(originalFilmWidth, originalFilmWidth);
// Get the tile size
vector<int> tileSize = config->cfg.GetIntVector("batch.tile", "256 256");
if (tileSize.size() != 2)
throw runtime_error("Syntax error in batch.tile (required 2 parameters)");
tileSize[0] = Max<int>(64, Min<int>(tileSize[0], originalFilmWidth));
tileSize[1] = Max<int>(64, Min<int>(tileSize[1], originalFilmHeight));
SLG_LOG("Tile size: " << tileSize[0] << " x " << tileSize[1]);
// Get the loop count
int loopCount = config->cfg.GetInt("batch.tile.loops", 1);
bool sessionStarted = false;
for (int loopIndex = 0; loopIndex < loopCount; ++loopIndex) {
u_int tileX = 0;
u_int tileY = 0;
const double startTime = WallClockTime();
// To setup new rendering parameters
if (loopIndex > 0) {
// I can begin an edit only after the start
session->BeginEdit();
// I have to use a new seed or I will render exactly the same images
session->renderEngine->GenerateNewSeed();
}
for (;;) {
SLG_LOG("Rendering tile offset: (" << tileX << "/" << originalFilmWidth + 2 * filterBorder << ", " <<
tileY << "/" << originalFilmHeight + 2 * filterBorder << ")");
// Set the film subregion to render
u_int filmSubRegion[4];
filmSubRegion[0] = tileX;
filmSubRegion[1] = Min(tileX + tileSize[0] - 1, originalFilmWidth - 1) + 2 * filterBorder;
filmSubRegion[2] = tileY;
filmSubRegion[3] = Min(tileY + tileSize[1] - 1, originalFilmHeight - 1) + 2 * filterBorder;
SLG_LOG("Tile subregion: " <<
boost::lexical_cast<string>(filmSubRegion[0]) << " " <<
boost::lexical_cast<string>(filmSubRegion[1]) << " " <<
boost::lexical_cast<string>(filmSubRegion[2]) << " " <<
boost::lexical_cast<string>(filmSubRegion[3]));
// Update the camera and resize the film
session->renderConfig->scene->camera->Update(
originalFilmWidth + 2 * filterBorder,
originalFilmHeight + 2 * filterBorder,
filmSubRegion);
session->film->Init(session->renderConfig->scene->camera->GetFilmWeight(),
session->renderConfig->scene->camera->GetFilmHeight());
if (sessionStarted) {
session->editActions.AddAction(CAMERA_EDIT);
session->editActions.AddAction(FILM_EDIT);
session->EndEdit();
} else {
session->Start();
sessionStarted = true;
}
double lastFilmUpdate = WallClockTime();
char buf[512];
for (;;) {
boost::this_thread::sleep(boost::posix_time::millisec(1000));
// Film update may be required by some render engine to
// update statistics, convergence test and more
if (WallClockTime() - lastFilmUpdate > 5.0) {
session->renderEngine->UpdateFilm();
lastFilmUpdate = WallClockTime();
}
const unsigned int pass = engine->GetPass();
if ((haltSpp > 0) && (pass >= haltSpp))
break;
// Convergence test is update inside UpdateFilm()
const float convergence = engine->GetConvergence();
if ((haltThreshold >= 0.f) && (1.f - convergence <= haltThreshold))
break;
// Print some information about the rendering progress
const double now = WallClockTime();
const double elapsedTime = now - startTime;
sprintf(buf, "[Loop step: %d/%d][Elapsed time: %3d][Samples %4d/%d][Convergence %f%%][Avg. samples/sec % 3.2fM on %.1fK tris]",
loopIndex + 1, loopCount,
int(elapsedTime), pass, haltSpp, 100.f * convergence, engine->GetTotalSamplesSec() / 1000000.0,
config->scene->dataSet->GetTotalTriangleCount() / 1000.0);
SLG_LOG(buf);
}
// Splat the current tile on the merge film
session->renderEngine->UpdateFilm();
{
boost::unique_lock<boost::mutex> lock(session->filmMutex);
mergeTileFilm.AddFilm(*(session->film), filterBorder, filterBorder,
filmSubRegion[1] - filmSubRegion[0] - 2 * filterBorder + 1,
filmSubRegion[3] - filmSubRegion[2] - 2 * filterBorder + 1,
tileX, tileY);
}
// Save the merge film
const string fileName = config->cfg.GetString("image.filename", "image.png");
SLG_LOG("Saving merged tiles to: " << fileName);
mergeTileFilm.UpdateScreenBuffer();
mergeTileFilm.SaveScreenBuffer(fileName);
// Advance to the next tile
tileX += tileSize[0];
if (tileX >= originalFilmWidth) {
tileX = 0;
tileY += tileSize[1];
if (tileY >= originalFilmHeight) {
// Rendering done
break;
}
}
// To setup new rendering parameters
session->BeginEdit();
}
}
// Stop the rendering
session->Stop();
delete session;
SLG_LOG("Done.");
return EXIT_SUCCESS;
}
PathIntegrator::PathIntegrator(PathRenderEngine *re, Sampler *samp) :
renderEngine(re), sampler(samp) {
sampleBuffer = renderEngine->film->GetFreeSampleBuffer();
statsRenderingStart = WallClockTime();
statsTotalSampleCount = 0;
}
