void BloomFilterPlugin::InitFilterTable(const Film &film) { const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); // Compute image-space extent of bloom effect const u_int bloomSupport = Float2UInt(radius * Max(width, height)); bloomWidth = bloomSupport / 2; // Initialize bloom filter table delete[] bloomFilter; bloomFilterSize = 2 * bloomWidth * bloomWidth + 1; bloomFilter = new float[bloomFilterSize]; for (u_int i = 0; i < bloomFilterSize; ++i) bloomFilter[i] = 0.f; for (u_int i = 0; i < bloomWidth * bloomWidth; ++i) { const float z0 = 3.8317f; const float dist = z0 * sqrtf(i) / bloomWidth; if (dist == 0.f) bloomFilter[i] = 1.f; else if (dist >= z0) bloomFilter[i] = 0.f; else { // Airy function //const float b = boost::math::cyl_bessel_j(1, dist); //bloomFilter[i] = powf(2*b/dist, 2.f); // Gaussian approximation // best-fit sigma^2 for above airy function, based on RMSE // depends on choice of zero const float sigma2 = 1.698022698724f; bloomFilter[i] = expf(-dist * dist / sigma2); } } }
void BloomFilterPlugin::Apply(Film &film, const u_int index) { //const double t1 = WallClockTime(); Spectrum *pixels = (Spectrum *)film.channel_IMAGEPIPELINEs[index]->GetPixels(); const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); // Allocate the temporary buffer if required if ((!bloomBuffer) || (width * height != bloomBufferSize)) { delete[] bloomBuffer; delete[] bloomBufferTmp; bloomBufferSize = width * height; bloomBuffer = new Spectrum[bloomBufferSize]; bloomBufferTmp = new Spectrum[bloomBufferSize]; InitFilterTable(film); } // Apply separable filter BloomFilter(film, pixels); for (u_int i = 0; i < bloomBufferSize; ++i) { if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(i))) pixels[i] = Lerp(weight, pixels[i], bloomBuffer[i]); } //const double t2 = WallClockTime(); //SLG_LOG("Bloom time: " << int((t2 - t1) * 1000.0) << "ms"); }
void PathHybridState::Init(const PathHybridRenderThread *thread) { PathHybridRenderEngine *renderEngine = (PathHybridRenderEngine *)thread->renderEngine; Scene *scene = renderEngine->renderConfig->scene; depth = 1; lastPdfW = 1.f; throuput = Spectrum(1.f); directLightRadiance = Spectrum(); // Initialize eye ray PerspectiveCamera *camera = scene->camera; Film *film = thread->threadFilm; const u_int filmWidth = film->GetWidth(); const u_int filmHeight = film->GetHeight(); sampleResults[0].screenX = std::min(sampler->GetSample(0) * filmWidth, (float)(filmWidth - 1)); sampleResults[0].screenY = std::min(sampler->GetSample(1) * filmHeight, (float)(filmHeight - 1)); camera->GenerateRay(sampleResults[0].screenX, sampleResults[0].screenY, &nextPathVertexRay, sampler->GetSample(2), sampler->GetSample(3)); sampleResults[0].alpha = 1.f; sampleResults[0].radiance = Spectrum(0.f); lastSpecular = true; }
void TileRepository::Tile::AddPass(const Film &tileFilm) { // Increase the pass count ++pass; // Update the done flag if (tileRepository->enableMultipassRendering) { // Check if convergence test is enable if (tileRepository->convergenceTestThreshold > 0.f) { // Add the tile to the all pass film allPassFilm->AddFilm(tileFilm, 0, 0, tileFilm.GetWidth(), tileFilm.GetHeight(), 0, 0); if (pass % 2 == 1) { // If it is an odd pass, add also to the even pass film evenPassFilm->AddFilm(tileFilm, 0, 0, tileFilm.GetWidth(), tileFilm.GetHeight(), 0, 0); } else { // Update tileRepository->tileMaxPixelValue before to check the // convergence UpdateMaxPixelValue(); // Update linear tone mapping plugin LinearToneMap *allLT = (LinearToneMap *)allPassFilm->GetImagePipeline()->GetPlugin(typeid(LinearToneMap)); allLT->scale = 1.f / tileRepository->tileMaxPixelValue; LinearToneMap *evenLT = (LinearToneMap *)evenPassFilm->GetImagePipeline()->GetPlugin(typeid(LinearToneMap)); evenLT->scale = allLT->scale; // If it is an even pass, check convergence status CheckConvergence(); } } if ((tileRepository->maxPassCount > 0) && (pass >= tileRepository->maxPassCount)) done = true; } else done = true; }
void BackgroundImgPlugin::UpdateFilmImageMap(const Film &film) { const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); // Check if I have to resample the image map if ((!filmImageMap) || (filmImageMap->GetWidth() != width) || (filmImageMap->GetHeight() != height)) { delete filmImageMap; filmImageMap = NULL; filmImageMap = imgMap->Copy(); filmImageMap->Resize(width, height); } }
void BloomFilterPlugin::BloomFilterY(const Film &film) { const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); // Apply bloom filter to image pixels #pragma omp parallel for for ( // Visual C++ 2013 supports only OpenMP 2.5 #if _OPENMP >= 200805 unsigned #endif int x = 0; x < width; ++x) { for (u_int y = 0; y < height; ++y) { if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(x, y))) { // Compute bloom for pixel (x, y) // Compute extent of pixels contributing bloom const u_int y0 = Max<u_int>(y, bloomWidth) - bloomWidth; const u_int y1 = Min<u_int>(y + bloomWidth, height - 1); float sumWt = 0.f; const u_int bx = x; Spectrum &pixel(bloomBuffer[x + y * width]); pixel = Spectrum(); for (u_int by = y0; by <= y1; ++by) { if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(bx, by))) { // Accumulate bloom from pixel (bx, by) const u_int dist2 = (x - bx) * (x - bx) + (y - by) * (y - by); const float wt = bloomFilter[dist2]; if (wt == 0.f) continue; const u_int bloomOffset = bx + by * width; sumWt += wt; pixel += wt * bloomBufferTmp[bloomOffset]; } } if (sumWt > 0.f) pixel /= sumWt; } } } }
void Reinhard02ToneMap::Apply(Film &film, const u_int index) { Spectrum *pixels = (Spectrum *)film.channel_IMAGEPIPELINEs[index]->GetPixels(); RGBColor *rgbPixels = (RGBColor *)pixels; const float alpha = .1f; const u_int pixelCount = film.GetWidth() * film.GetHeight(); float Ywa = 0.f; for (u_int i = 0; i < pixelCount; ++i) { if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(i)) && !rgbPixels[i].IsInf()) Ywa += logf(Max(rgbPixels[i].Y(), 1e-6f)); } if (pixelCount > 0) Ywa = expf(Ywa / pixelCount); // Avoid division by zero if (Ywa == 0.f) Ywa = 1.f; const float invB2 = (burn > 0.f) ? 1.f / (burn * burn) : 1e5f; const float scale = alpha / Ywa; const float preS = scale / preScale; const float postS = scale * postScale; #pragma omp parallel for for ( // Visual C++ 2013 supports only OpenMP 2.5 #if _OPENMP >= 200805 unsigned #endif int i = 0; i < pixelCount; ++i) { if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(i))) { const float ys = rgbPixels[i].Y() * preS; // Note: I don't need to convert to XYZ and back because I'm only // scaling the value. rgbPixels[i] *= postS * (1.f + ys * invB2) / (1.f + ys); } } }
void BackgroundImgPlugin::Apply(Film &film, const u_int index) { if (!film.HasChannel(Film::ALPHA)) { // I can not work without alpha channel return; } // Check if I have to resample the image map UpdateFilmImageMap(film); Spectrum *pixels = (Spectrum *)film.channel_IMAGEPIPELINEs[index]->GetPixels(); const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); #pragma omp parallel for for ( // Visual C++ 2013 supports only OpenMP 2.5 #if _OPENMP >= 200805 unsigned #endif int y = 0; y < height; ++y) { for (u_int x = 0; x < width; ++x) { const u_int filmPixelIndex = x + y * width; if (*(film.channel_FRAMEBUFFER_MASK->GetPixel(filmPixelIndex))) { float alpha; film.channel_ALPHA->GetWeightedPixel(x, y, &alpha); // Need to flip the along the Y axis for the image const u_int imgPixelIndex = x + (height - y - 1) * width; pixels[filmPixelIndex] = Lerp(alpha, filmImageMap->GetStorage()->GetSpectrum(imgPixelIndex), pixels[filmPixelIndex]); } } } }
void BiDirVMCPURenderThread::RenderFuncVM() { //SLG_LOG("[BiDirVMCPURenderThread::" << threadIndex << "] Rendering thread started"); //-------------------------------------------------------------------------- // Initialization //-------------------------------------------------------------------------- BiDirVMCPURenderEngine *engine = (BiDirVMCPURenderEngine *)renderEngine; RandomGenerator *rndGen = new RandomGenerator(engine->seedBase + threadIndex); Scene *scene = engine->renderConfig->scene; Camera *camera = scene->camera; Film *film = threadFilm; const u_int filmWidth = film->GetWidth(); const u_int filmHeight = film->GetHeight(); pixelCount = filmWidth * filmHeight; // Setup the samplers vector<Sampler *> samplers(engine->lightPathsCount, NULL); const u_int sampleSize = sampleBootSizeVM + // To generate the initial light vertex and trace eye ray engine->maxLightPathDepth * sampleLightStepSize + // For each light vertex engine->maxEyePathDepth * sampleEyeStepSize; // For each eye vertex // metropolisSharedTotalLuminance and metropolisSharedSampleCount are // initialized inside MetropolisSampler::RequestSamples() double metropolisSharedTotalLuminance, metropolisSharedSampleCount; for (u_int i = 0; i < samplers.size(); ++i) { Sampler *sampler = engine->renderConfig->AllocSampler(rndGen, film, &metropolisSharedTotalLuminance, &metropolisSharedSampleCount); sampler->RequestSamples(sampleSize); samplers[i] = sampler; } u_int iteration = 0; vector<vector<SampleResult> > samplesResults(samplers.size()); vector<vector<PathVertexVM> > lightPathsVertices(samplers.size()); vector<Point> lensPoints(samplers.size()); HashGrid hashGrid; const u_int haltDebug = engine->renderConfig->GetProperty("batch.haltdebug").Get<u_int>(); for(u_int steps = 0; !boost::this_thread::interruption_requested(); ++steps) { // Clear the arrays for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { samplesResults[samplerIndex].clear(); lightPathsVertices[samplerIndex].clear(); } // Setup vertex merging float radius = engine->baseRadius; radius /= powf(float(iteration + 1), .5f * (1.f - engine->radiusAlpha)); radius = Max(radius, DEFAULT_EPSILON_STATIC); const float radius2 = radius * radius; const float vmFactor = M_PI * radius2 * engine->lightPathsCount; vmNormalization = 1.f / vmFactor; const float etaVCM = vmFactor; misVmWeightFactor = MIS(etaVCM); misVcWeightFactor = MIS(1.f / etaVCM); // Using the same time for all rays in the same pass is required by the // current implementation (i.e. I can not mix paths with different // times). However this is detrimental for the Metropolis sampler. const float time = rndGen->floatValue(); //---------------------------------------------------------------------- // Trace all light paths //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { Sampler *sampler = samplers[samplerIndex]; // Sample a point on the camera lens if (!camera->SampleLens(time, sampler->GetSample(3), sampler->GetSample(4), &lensPoints[samplerIndex])) continue; TraceLightPath(time, sampler, lensPoints[samplerIndex], lightPathsVertices[samplerIndex], samplesResults[samplerIndex]); } //---------------------------------------------------------------------- // Store all light path vertices in the k-NN accelerator //---------------------------------------------------------------------- hashGrid.Build(lightPathsVertices, radius); //cout << "==========================================\n"; //cout << "Iteration: " << iteration << " Paths: " << engine->lightPathsCount << " Light path vertices: "<< hashGrid.GetVertexCount() <<"\n"; //---------------------------------------------------------------------- // Trace all eye paths //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { Sampler *sampler = samplers[samplerIndex]; PathVertexVM eyeVertex; SampleResult eyeSampleResult(Film::RADIANCE_PER_PIXEL_NORMALIZED | Film::ALPHA, 1); eyeSampleResult.alpha = 1.f; Ray eyeRay; eyeSampleResult.filmX = min(sampler->GetSample(0) * filmWidth, (float)(filmWidth - 1)); eyeSampleResult.filmY = min(sampler->GetSample(1) * filmHeight, (float)(filmHeight - 1)); camera->GenerateRay(eyeSampleResult.filmX, eyeSampleResult.filmY, &eyeRay, sampler->GetSample(9), sampler->GetSample(10), time); eyeVertex.bsdf.hitPoint.fixedDir = -eyeRay.d; eyeVertex.throughput = Spectrum(1.f); const float cosAtCamera = Dot(scene->camera->GetDir(), eyeRay.d); const float cameraPdfW = 1.f / (cosAtCamera * cosAtCamera * cosAtCamera * scene->camera->GetPixelArea()); eyeVertex.dVCM = MIS(1.f / cameraPdfW); eyeVertex.dVC = 1.f; eyeVertex.dVM = 1.f; eyeVertex.depth = 1; while (eyeVertex.depth <= engine->maxEyePathDepth) { const u_int sampleOffset = sampleBootSizeVM + engine->maxLightPathDepth * sampleLightStepSize + (eyeVertex.depth - 1) * sampleEyeStepSize; RayHit eyeRayHit; Spectrum connectionThroughput, connectEmission; const bool hit = scene->Intersect(device, false, &eyeVertex.volInfo, sampler->GetSample(sampleOffset), &eyeRay, &eyeRayHit, &eyeVertex.bsdf, &connectionThroughput, NULL, NULL, &connectEmission); // I account for volume emission only with path tracing (i.e. here and // not in any other place) eyeSampleResult.radiancePerPixelNormalized[0] += connectEmission; if (!hit) { // Nothing was hit, look for infinitelight // This is a trick, you can not have a BSDF of something that has // not been hit. DirectHitInfiniteLight must be aware of this. eyeVertex.bsdf.hitPoint.fixedDir = -eyeRay.d; eyeVertex.throughput *= connectionThroughput; DirectHitLight(false, eyeVertex, &eyeSampleResult.radiancePerPixelNormalized[0]); if (eyeVertex.depth == 1) eyeSampleResult.alpha = 0.f; break; } eyeVertex.throughput *= connectionThroughput; // Something was hit // Update MIS constants const float factor = 1.f / MIS(AbsDot(eyeVertex.bsdf.hitPoint.shadeN, eyeVertex.bsdf.hitPoint.fixedDir)); eyeVertex.dVCM *= MIS(eyeRayHit.t * eyeRayHit.t) * factor; eyeVertex.dVC *= factor; eyeVertex.dVM *= factor; // Check if it is a light source if (eyeVertex.bsdf.IsLightSource()) DirectHitLight(true, eyeVertex, &eyeSampleResult.radiancePerPixelNormalized[0]); // Note: pass-through check is done inside Scene::Intersect() //-------------------------------------------------------------- // Direct light sampling //-------------------------------------------------------------- DirectLightSampling(time, sampler->GetSample(sampleOffset + 1), sampler->GetSample(sampleOffset + 2), sampler->GetSample(sampleOffset + 3), sampler->GetSample(sampleOffset + 4), sampler->GetSample(sampleOffset + 5), eyeVertex, &eyeSampleResult.radiancePerPixelNormalized[0]); if (!eyeVertex.bsdf.IsDelta()) { //---------------------------------------------------------- // Connect vertex path ray with all light path vertices //---------------------------------------------------------- const vector<PathVertexVM> &lightPathVertices = lightPathsVertices[samplerIndex]; for (vector<PathVertexVM>::const_iterator lightPathVertex = lightPathVertices.begin(); lightPathVertex < lightPathVertices.end(); ++lightPathVertex) ConnectVertices(time, eyeVertex, *lightPathVertex, &eyeSampleResult, sampler->GetSample(sampleOffset + 6)); //---------------------------------------------------------- // Vertex Merging step //---------------------------------------------------------- hashGrid.Process(this, eyeVertex, &eyeSampleResult.radiancePerPixelNormalized[0]); } //-------------------------------------------------------------- // Build the next vertex path ray //-------------------------------------------------------------- if (!Bounce(time, sampler, sampleOffset + 7, &eyeVertex, &eyeRay)) break; } samplesResults[samplerIndex].push_back(eyeSampleResult); } //---------------------------------------------------------------------- // Splat all samples //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) samplers[samplerIndex]->NextSample(samplesResults[samplerIndex]); ++iteration; #ifdef WIN32 // Work around Windows bad scheduling renderThread->yield(); #endif //hashGrid.PrintStatistics(); if ((haltDebug > 0u) && (steps >= haltDebug)) break; } for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) delete samplers[samplerIndex]; delete rndGen; //SLG_LOG("[BiDirVMCPURenderThread::" << renderThread->threadIndex << "] Rendering thread halted"); }
void BloomFilterPlugin::ApplyOCL(Film &film, const u_int index) { const u_int width = film.GetWidth(); const u_int height = film.GetHeight(); if ((!bloomFilter) || (width * height != bloomBufferSize)) { bloomBufferSize = width * height; InitFilterTable(film); } if (!bloomFilterXKernel) { oclIntersectionDevice = film.oclIntersectionDevice; // Allocate OpenCL buffers film.ctx->SetVerbose(true); oclIntersectionDevice->AllocBufferRW(&oclBloomBuffer, bloomBufferSize * sizeof(Spectrum), "Bloom buffer"); oclIntersectionDevice->AllocBufferRW(&oclBloomBufferTmp, bloomBufferSize * sizeof(Spectrum), "Bloom temporary buffer"); oclIntersectionDevice->AllocBufferRO(&oclBloomFilter, bloomFilter, bloomFilterSize * sizeof(float), "Bloom filter table"); film.ctx->SetVerbose(false); // Compile sources const double tStart = WallClockTime(); cl::Program *program = ImagePipelinePlugin::CompileProgram( film, "-D LUXRAYS_OPENCL_KERNEL -D SLG_OPENCL_KERNEL", slg::ocl::KernelSource_plugin_bloom_funcs, "BloomFilterPlugin"); //---------------------------------------------------------------------- // BloomFilterPlugin_FilterX kernel //---------------------------------------------------------------------- SLG_LOG("[BloomFilterPlugin] Compiling BloomFilterPlugin_FilterX Kernel"); bloomFilterXKernel = new cl::Kernel(*program, "BloomFilterPlugin_FilterX"); // Set kernel arguments u_int argIndex = 0; bloomFilterXKernel->setArg(argIndex++, film.GetWidth()); bloomFilterXKernel->setArg(argIndex++, film.GetHeight()); bloomFilterXKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); bloomFilterXKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); bloomFilterXKernel->setArg(argIndex++, *oclBloomBuffer); bloomFilterXKernel->setArg(argIndex++, *oclBloomBufferTmp); bloomFilterXKernel->setArg(argIndex++, *oclBloomFilter); bloomFilterXKernel->setArg(argIndex++, bloomWidth); //---------------------------------------------------------------------- // BloomFilterPlugin_FilterY kernel //---------------------------------------------------------------------- SLG_LOG("[BloomFilterPlugin] Compiling BloomFilterPlugin_FilterY Kernel"); bloomFilterYKernel = new cl::Kernel(*program, "BloomFilterPlugin_FilterY"); // Set kernel arguments argIndex = 0; bloomFilterYKernel->setArg(argIndex++, film.GetWidth()); bloomFilterYKernel->setArg(argIndex++, film.GetHeight()); bloomFilterYKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); bloomFilterYKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); bloomFilterYKernel->setArg(argIndex++, *oclBloomBuffer); bloomFilterYKernel->setArg(argIndex++, *oclBloomBufferTmp); bloomFilterYKernel->setArg(argIndex++, *oclBloomFilter); bloomFilterYKernel->setArg(argIndex++, bloomWidth); //---------------------------------------------------------------------- // BloomFilterPlugin_Merge kernel //---------------------------------------------------------------------- SLG_LOG("[BloomFilterPlugin] Compiling BloomFilterPlugin_Merge Kernel"); bloomFilterMergeKernel = new cl::Kernel(*program, "BloomFilterPlugin_Merge"); // Set kernel arguments argIndex = 0; bloomFilterMergeKernel->setArg(argIndex++, film.GetWidth()); bloomFilterMergeKernel->setArg(argIndex++, film.GetHeight()); bloomFilterMergeKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); bloomFilterMergeKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); bloomFilterMergeKernel->setArg(argIndex++, *oclBloomBuffer); bloomFilterMergeKernel->setArg(argIndex++, weight); //---------------------------------------------------------------------- delete program; const double tEnd = WallClockTime(); SLG_LOG("[BloomFilterPlugin] Kernels compilation time: " << int((tEnd - tStart) * 1000.0) << "ms"); } oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*bloomFilterXKernel, cl::NullRange, cl::NDRange(RoundUp(film.GetWidth() * film.GetHeight(), 256u)), cl::NDRange(256)); oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*bloomFilterYKernel, cl::NullRange, cl::NDRange(RoundUp(film.GetWidth() * film.GetHeight(), 256u)), cl::NDRange(256)); oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*bloomFilterMergeKernel, cl::NullRange, cl::NDRange(RoundUp(film.GetWidth() * film.GetHeight(), 256u)), cl::NDRange(256)); }
void HybridRenderThread::RenderFunc() { //SLG_LOG("[HybridRenderThread::" << threadIndex << "] Rendering thread started"); boost::this_thread::disable_interruption di; Film *film = threadFilm; const u_int filmWidth = film->GetWidth(); const u_int filmHeight = film->GetHeight(); pixelCount = filmWidth * filmHeight; RandomGenerator *rndGen = new RandomGenerator(threadIndex + renderEngine->seedBase); const u_int incrementStep = 4096; vector<HybridRenderState *> states(incrementStep); try { // Initialize the first states for (u_int i = 0; i < states.size(); ++i) states[i] = AllocRenderState(rndGen); u_int generateIndex = 0; u_int collectIndex = 0; while (!boost::this_thread::interruption_requested()) { // Generate new rays up to the point to have 3 pending buffers while (pendingRayBuffers < 3) { states[generateIndex]->GenerateRays(this); generateIndex = (generateIndex + 1) % states.size(); if (generateIndex == collectIndex) { //SLG_LOG("[HybridRenderThread::" << threadIndex << "] Increasing states size by " << incrementStep); //SLG_LOG("[HybridRenderThread::" << threadIndex << "] State size: " << states.size()); // Insert a set of new states and continue states.insert(states.begin() + generateIndex, incrementStep, NULL); for (u_int i = generateIndex; i < generateIndex + incrementStep; ++i) states[i] = AllocRenderState(rndGen); collectIndex += incrementStep; } } //SLG_LOG("[HybridRenderThread::" << threadIndex << "] State size: " << states.size()); //SLG_LOG("[HybridRenderThread::" << threadIndex << "] generateIndex: " << generateIndex); //SLG_LOG("[HybridRenderThread::" << threadIndex << "] collectIndex: " << collectIndex); //SLG_LOG("[HybridRenderThread::" << threadIndex << "] pendingRayBuffers: " << pendingRayBuffers); // Collect rays up to the point to have only 1 pending buffer while (pendingRayBuffers > 1) { samplesCount += states[collectIndex]->CollectResults(this); const u_int newCollectIndex = (collectIndex + 1) % states.size(); // A safety-check, it should never happen if (newCollectIndex == generateIndex) break; collectIndex = newCollectIndex; } } //SLG_LOG("[HybridRenderThread::" << threadIndex << "] Rendering thread halted"); } catch (boost::thread_interrupted) { SLG_LOG("[HybridRenderThread::" << threadIndex << "] Rendering thread halted"); } #ifndef LUXRAYS_DISABLE_OPENCL catch (cl::Error err) { SLG_LOG("[HybridRenderThread::" << threadIndex << "] Rendering thread ERROR: " << err.what() << "(" << luxrays::utils::oclErrorString(err.err()) << ")"); } #endif // Clean current ray buffers if (currentRayBufferToSend) { currentRayBufferToSend->Reset(); freeRayBuffers.push_back(currentRayBufferToSend); currentRayBufferToSend = NULL; } if (currentReiceivedRayBuffer) { currentReiceivedRayBuffer->Reset(); freeRayBuffers.push_back(currentReiceivedRayBuffer); currentReiceivedRayBuffer = NULL; } // Free all states for (u_int i = 0; i < states.size(); ++i) delete states[i]; delete rndGen; // Remove all pending ray buffers while (pendingRayBuffers > 0) { RayBuffer *rayBuffer = device->PopRayBuffer(); --(pendingRayBuffers); rayBuffer->Reset(); freeRayBuffers.push_back(rayBuffer); } }
void LightCPURenderThread::TraceEyePath(Sampler *sampler, vector<SampleResult> *sampleResults) { LightCPURenderEngine *engine = (LightCPURenderEngine *)renderEngine; Scene *scene = engine->renderConfig->scene; PerspectiveCamera *camera = scene->camera; Film *film = threadFilm; const u_int filmWidth = film->GetWidth(); const u_int filmHeight = film->GetHeight(); // Sample offsets const u_int sampleBootSize = 11; const u_int sampleEyeStepSize = 3; Ray eyeRay; const float screenX = min(sampler->GetSample(0) * filmWidth, (float)(filmWidth - 1)); const float screenY = min(sampler->GetSample(1) * filmHeight, (float)(filmHeight - 1)); camera->GenerateRay(screenX, screenY, &eyeRay, sampler->GetSample(9), sampler->GetSample(10)); Spectrum radiance, eyePathThroughput(1.f, 1.f, 1.f); int depth = 1; while (depth <= engine->maxPathDepth) { const u_int sampleOffset = sampleBootSize + (depth - 1) * sampleEyeStepSize; RayHit eyeRayHit; BSDF bsdf; Spectrum connectionThroughput; const bool somethingWasHit = scene->Intersect(device, false, sampler->GetSample(sampleOffset), &eyeRay, &eyeRayHit, &bsdf, &connectionThroughput); if (!somethingWasHit) { // Nothing was hit, check infinite lights (including sun) const Spectrum throughput = eyePathThroughput * connectionThroughput; if (scene->envLight) radiance += throughput * scene->envLight->GetRadiance(scene, -eyeRay.d); if (scene->sunLight) radiance += throughput * scene->sunLight->GetRadiance(scene, -eyeRay.d); break; } else { // Something was hit, check if it is a light source if (bsdf.IsLightSource()) radiance = eyePathThroughput * connectionThroughput * bsdf.GetEmittedRadiance(scene); else { // Check if it is a specular bounce float bsdfPdf; Vector sampledDir; BSDFEvent event; float cosSampleDir; const Spectrum bsdfSample = bsdf.Sample(&sampledDir, sampler->GetSample(sampleOffset + 1), sampler->GetSample(sampleOffset + 2), &bsdfPdf, &cosSampleDir, &event); if (bsdfSample.Black() || ((depth == 1) && !(event & SPECULAR))) break; // If depth = 1 and it is a specular bounce, I continue to trace the // eye path looking for a light source eyePathThroughput *= connectionThroughput * bsdfSample * (cosSampleDir / bsdfPdf); assert (!eyePathThroughput.IsNaN() && !eyePathThroughput.IsInf()); eyeRay = Ray(bsdf.hitPoint, sampledDir); } ++depth; } } // Add a sample even if it is black in order to avoid aliasing problems // between sampled pixel and not sampled one (in PER_PIXEL_NORMALIZED buffer) AddSampleResult(*sampleResults, PER_PIXEL_NORMALIZED, screenX, screenY, radiance, (depth == 1) ? 1.f : 0.f); }
void BiDirVMCPURenderThread::RenderFuncVM() { //SLG_LOG("[BiDirVMCPURenderThread::" << threadIndex << "] Rendering thread started"); //-------------------------------------------------------------------------- // Initialization //-------------------------------------------------------------------------- BiDirVMCPURenderEngine *engine = (BiDirVMCPURenderEngine *)renderEngine; RandomGenerator *rndGen = new RandomGenerator(engine->seedBase + threadIndex); Scene *scene = engine->renderConfig->scene; PerspectiveCamera *camera = scene->camera; Film *film = threadFilm; const unsigned int filmWidth = film->GetWidth(); const unsigned int filmHeight = film->GetHeight(); pixelCount = filmWidth * filmHeight; // Setup the samplers vector<Sampler *> samplers(engine->lightPathsCount, NULL); const unsigned int sampleSize = sampleBootSize + // To generate the initial light vertex and trace eye ray engine->maxLightPathDepth * sampleLightStepSize + // For each light vertex engine->maxEyePathDepth * sampleEyeStepSize; // For each eye vertex double metropolisSharedTotalLuminance, metropolisSharedSampleCount; for (u_int i = 0; i < samplers.size(); ++i) { Sampler *sampler = engine->renderConfig->AllocSampler(rndGen, film, &metropolisSharedTotalLuminance, &metropolisSharedSampleCount); sampler->RequestSamples(sampleSize); samplers[i] = sampler; } u_int iteration = 0; vector<vector<SampleResult> > samplesResults(samplers.size()); vector<vector<PathVertexVM> > lightPathsVertices(samplers.size()); vector<Point> lensPoints(samplers.size()); HashGrid hashGrid; while (!boost::this_thread::interruption_requested()) { // Clear the arrays for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { samplesResults[samplerIndex].clear(); lightPathsVertices[samplerIndex].clear(); } // Setup vertex merging float radius = engine->baseRadius; radius /= powf(float(iteration + 1), .5f * (1.f - engine->radiusAlpha)); radius = Max(radius, DEFAULT_EPSILON_STATIC); const float radius2 = radius * radius; const float vmFactor = M_PI * radius2 * engine->lightPathsCount; vmNormalization = 1.f / vmFactor; const float etaVCM = vmFactor; misVmWeightFactor = MIS(etaVCM); misVcWeightFactor = MIS(1.f / etaVCM); //---------------------------------------------------------------------- // Trace all light paths //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { Sampler *sampler = samplers[samplerIndex]; // Sample a point on the camera lens if (!camera->SampleLens(sampler->GetSample(3), sampler->GetSample(4), &lensPoints[samplerIndex])) continue; TraceLightPath(sampler, lensPoints[samplerIndex], lightPathsVertices[samplerIndex], samplesResults[samplerIndex]); } //---------------------------------------------------------------------- // Store all light path vertices in the k-NN accelerator //---------------------------------------------------------------------- hashGrid.Build(lightPathsVertices, radius); //---------------------------------------------------------------------- // Trace all eye paths //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) { Sampler *sampler = samplers[samplerIndex]; const vector<PathVertexVM> &lightPathVertices = lightPathsVertices[samplerIndex]; PathVertexVM eyeVertex; SampleResult eyeSampleResult; eyeSampleResult.type = PER_PIXEL_NORMALIZED; eyeSampleResult.alpha = 1.f; Ray eyeRay; eyeSampleResult.screenX = min(sampler->GetSample(0) * filmWidth, (float)(filmWidth - 1)); eyeSampleResult.screenY = min(sampler->GetSample(1) * filmHeight, (float)(filmHeight - 1)); camera->GenerateRay(eyeSampleResult.screenX, eyeSampleResult.screenY, &eyeRay, sampler->GetSample(9), sampler->GetSample(10)); eyeVertex.bsdf.hitPoint.fixedDir = -eyeRay.d; eyeVertex.throughput = Spectrum(1.f, 1.f, 1.f); const float cosAtCamera = Dot(scene->camera->GetDir(), eyeRay.d); const float cameraPdfW = 1.f / (cosAtCamera * cosAtCamera * cosAtCamera * scene->camera->GetPixelArea()); eyeVertex.dVCM = MIS(1.f / cameraPdfW); eyeVertex.dVC = 1.f; eyeVertex.dVM = 1.f; eyeVertex.depth = 1; while (eyeVertex.depth <= engine->maxEyePathDepth) { const unsigned int sampleOffset = sampleBootSize + engine->maxLightPathDepth * sampleLightStepSize + (eyeVertex.depth - 1) * sampleEyeStepSize; RayHit eyeRayHit; Spectrum connectionThroughput; if (!scene->Intersect(device, false, sampler->GetSample(sampleOffset), &eyeRay, &eyeRayHit, &eyeVertex.bsdf, &connectionThroughput)) { // Nothing was hit, look for infinitelight // This is a trick, you can not have a BSDF of something that has // not been hit. DirectHitInfiniteLight must be aware of this. eyeVertex.bsdf.hitPoint.fixedDir = -eyeRay.d; eyeVertex.throughput *= connectionThroughput; DirectHitLight(false, eyeVertex, &eyeSampleResult.radiance); if (eyeVertex.depth == 1) eyeSampleResult.alpha = 0.f; break; } eyeVertex.throughput *= connectionThroughput; // Something was hit // Update MIS constants const float factor = 1.f / MIS(AbsDot(eyeVertex.bsdf.hitPoint.shadeN, eyeVertex.bsdf.hitPoint.fixedDir)); eyeVertex.dVCM *= MIS(eyeRayHit.t * eyeRayHit.t) * factor; eyeVertex.dVC *= factor; eyeVertex.dVM *= factor; // Check if it is a light source if (eyeVertex.bsdf.IsLightSource()) DirectHitLight(true, eyeVertex, &eyeSampleResult.radiance); // Note: pass-through check is done inside SceneIntersect() //-------------------------------------------------------------- // Direct light sampling //-------------------------------------------------------------- DirectLightSampling(sampler->GetSample(sampleOffset + 1), sampler->GetSample(sampleOffset + 2), sampler->GetSample(sampleOffset + 3), sampler->GetSample(sampleOffset + 4), sampler->GetSample(sampleOffset + 5), eyeVertex, &eyeSampleResult.radiance); if (!eyeVertex.bsdf.IsDelta()) { //---------------------------------------------------------- // Connect vertex path ray with all light path vertices //---------------------------------------------------------- for (vector<PathVertexVM>::const_iterator lightPathVertex = lightPathVertices.begin(); lightPathVertex < lightPathVertices.end(); ++lightPathVertex) ConnectVertices(eyeVertex, *lightPathVertex, &eyeSampleResult, sampler->GetSample(sampleOffset + 6)); //---------------------------------------------------------- // Vertex Merging step //---------------------------------------------------------- hashGrid.Process(this, eyeVertex, &eyeSampleResult.radiance); } //-------------------------------------------------------------- // Build the next vertex path ray //-------------------------------------------------------------- if (!Bounce(sampler, sampleOffset + 7, &eyeVertex, &eyeRay)) break; ++(eyeVertex.depth); } samplesResults[samplerIndex].push_back(eyeSampleResult); } //---------------------------------------------------------------------- // Splat all samples //---------------------------------------------------------------------- for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) samplers[samplerIndex]->NextSample(samplesResults[samplerIndex]); ++iteration; #ifdef WIN32 // Work around Windows bad scheduling renderThread->yield(); #endif } for (u_int samplerIndex = 0; samplerIndex < samplers.size(); ++samplerIndex) delete samplers[samplerIndex]; delete rndGen; //SLG_LOG("[BiDirVMCPURenderThread::" << renderThread->threadIndex << "] Rendering thread halted"); }
void PathCPURenderThread::RenderFunc() { //SLG_LOG("[PathCPURenderEngine::" << threadIndex << "] Rendering thread started"); //-------------------------------------------------------------------------- // Initialization //-------------------------------------------------------------------------- PathCPURenderEngine *engine = (PathCPURenderEngine *)renderEngine; RandomGenerator *rndGen = new RandomGenerator(engine->seedBase + threadIndex); Scene *scene = engine->renderConfig->scene; PerspectiveCamera *camera = scene->camera; Film * film = threadFilm; const unsigned int filmWidth = film->GetWidth(); const unsigned int filmHeight = film->GetHeight(); // Setup the sampler double metropolisSharedTotalLuminance, metropolisSharedSampleCount; Sampler *sampler = engine->renderConfig->AllocSampler(rndGen, film, &metropolisSharedTotalLuminance, &metropolisSharedSampleCount); const unsigned int sampleBootSize = 4; const unsigned int sampleStepSize = 9; const unsigned int sampleSize = sampleBootSize + // To generate eye ray engine->maxPathDepth * sampleStepSize; // For each path vertex sampler->RequestSamples(sampleSize); //-------------------------------------------------------------------------- // Trace paths //-------------------------------------------------------------------------- vector<SampleResult> sampleResults(1); sampleResults[0].type = PER_PIXEL_NORMALIZED; while (!boost::this_thread::interruption_requested()) { float alpha = 1.f; Ray eyeRay; const float screenX = min(sampler->GetSample(0) * filmWidth, (float)(filmWidth - 1)); const float screenY = min(sampler->GetSample(1) * filmHeight, (float)(filmHeight - 1)); camera->GenerateRay(screenX, screenY, &eyeRay, sampler->GetSample(2), sampler->GetSample(3)); int depth = 1; bool lastSpecular = true; float lastPdfW = 1.f; Spectrum radiance; Spectrum pathThrouput(1.f, 1.f, 1.f); BSDF bsdf; while (depth <= engine->maxPathDepth) { const unsigned int sampleOffset = sampleBootSize + (depth - 1) * sampleStepSize; RayHit eyeRayHit; Spectrum connectionThroughput; if (!scene->Intersect(device, false, sampler->GetSample(sampleOffset), &eyeRay, &eyeRayHit, &bsdf, &connectionThroughput)) { // Nothing was hit, look for infinitelight DirectHitInfiniteLight(lastSpecular, pathThrouput * connectionThroughput, eyeRay.d, lastPdfW, &radiance); if (depth == 1) alpha = 0.f; break; } pathThrouput *= connectionThroughput; // Something was hit // Check if it is a light source if (bsdf.IsLightSource()) { DirectHitFiniteLight(lastSpecular, pathThrouput, eyeRayHit.t, bsdf, lastPdfW, &radiance); } // Note: pass-through check is done inside SceneIntersect() //------------------------------------------------------------------ // Direct light sampling //------------------------------------------------------------------ DirectLightSampling(sampler->GetSample(sampleOffset + 1), sampler->GetSample(sampleOffset + 2), sampler->GetSample(sampleOffset + 3), sampler->GetSample(sampleOffset + 4), sampler->GetSample(sampleOffset + 5), pathThrouput, bsdf, depth, &radiance); //------------------------------------------------------------------ // Build the next vertex path ray //------------------------------------------------------------------ Vector sampledDir; BSDFEvent event; float cosSampledDir; const Spectrum bsdfSample = bsdf.Sample(&sampledDir, sampler->GetSample(sampleOffset + 6), sampler->GetSample(sampleOffset + 7), &lastPdfW, &cosSampledDir, &event); if (bsdfSample.Black()) break; lastSpecular = ((event & SPECULAR) != 0); if ((depth >= engine->rrDepth) && !lastSpecular) { // Russian Roulette const float prob = Max(bsdfSample.Filter(), engine->rrImportanceCap); if (sampler->GetSample(sampleOffset + 8) < prob) lastPdfW *= prob; else break; } pathThrouput *= bsdfSample * (cosSampledDir / lastPdfW); assert (!pathThrouput.IsNaN() && !pathThrouput.IsInf()); eyeRay = Ray(bsdf.hitPoint, sampledDir); ++depth; } assert (!radiance.IsNaN() && !radiance.IsInf()); sampleResults[0].screenX = screenX; sampleResults[0].screenY = screenY; sampleResults[0].radiance = radiance; sampleResults[0].alpha = alpha; sampler->NextSample(sampleResults); #ifdef WIN32 // Work around Windows bad scheduling renderThread->yield(); #endif } delete sampler; delete rndGen; //SLG_LOG("[PathCPURenderEngine::" << threadIndex << "] Rendering thread halted"); }
void BackgroundImgPlugin::ApplyOCL(Film &film, const u_int index) { if (!film.HasChannel(Film::ALPHA)) { // I can not work without alpha channel return; } // Check if I have to resample the image map UpdateFilmImageMap(film); if (!applyKernel) { oclIntersectionDevice = film.oclIntersectionDevice; slg::ocl::ImageMap imgMapDesc; imgMapDesc.channelCount = filmImageMap->GetChannelCount(); imgMapDesc.width = filmImageMap->GetWidth(); imgMapDesc.height = filmImageMap->GetHeight(); imgMapDesc.pageIndex = 0; imgMapDesc.pixelsIndex = 0; imgMapDesc.storageType = (slg::ocl::ImageMapStorageType)filmImageMap->GetStorage()->GetStorageType(); // Allocate OpenCL buffers film.ctx->SetVerbose(true); oclIntersectionDevice->AllocBufferRO(&oclFilmImageMapDesc, &imgMapDesc, sizeof(slg::ocl::ImageMap), "BackgroundImg image map description"); oclIntersectionDevice->AllocBufferRO(&oclFilmImageMap, filmImageMap->GetStorage()->GetPixelsData(), filmImageMap->GetStorage()->GetMemorySize(), "BackgroundImg image map"); film.ctx->SetVerbose(false); // Compile sources const double tStart = WallClockTime(); // Set #define symbols stringstream ssParams; ssParams.precision(6); ssParams << scientific << " -D LUXRAYS_OPENCL_KERNEL" << " -D SLG_OPENCL_KERNEL"; ssParams << " -D PARAM_HAS_IMAGEMAPS"; ssParams << " -D PARAM_IMAGEMAPS_PAGE_0"; ssParams << " -D PARAM_IMAGEMAPS_COUNT=1"; switch (imgMapDesc.storageType) { case slg::ocl::BYTE: ssParams << " -D PARAM_HAS_IMAGEMAPS_BYTE_FORMAT"; break; case slg::ocl::HALF: ssParams << " -D PARAM_HAS_IMAGEMAPS_HALF_FORMAT"; break; case slg::ocl::FLOAT: ssParams << " -D PARAM_HAS_IMAGEMAPS_FLOAT_FORMAT"; break; default: throw runtime_error("Unknown storage type in BackgroundImgPlugin::ApplyOCL(): " + ToString(imgMapDesc.storageType)); } switch (imgMapDesc.channelCount) { case 1: ssParams << " -D PARAM_HAS_IMAGEMAPS_1xCHANNELS"; break; case 2: ssParams << " -D PARAM_HAS_IMAGEMAPS_2xCHANNELS"; break; case 3: ssParams << " -D PARAM_HAS_IMAGEMAPS_3xCHANNELS"; break; case 4: ssParams << " -D PARAM_HAS_IMAGEMAPS_4xCHANNELS"; break; default: throw runtime_error("Unknown channel count in BackgroundImgPlugin::ApplyOCL(): " + ToString(imgMapDesc.channelCount)); } cl::Program *program = ImagePipelinePlugin::CompileProgram( film, ssParams.str(), slg::ocl::KernelSource_utils_funcs + slg::ocl::KernelSource_color_types + slg::ocl::KernelSource_color_funcs + slg::ocl::KernelSource_imagemap_types + slg::ocl::KernelSource_imagemap_funcs + slg::ocl::KernelSource_plugin_backgroundimg_funcs, "BackgroundImgPlugin"); //---------------------------------------------------------------------- // BackgroundImgPlugin_Apply kernel //---------------------------------------------------------------------- SLG_LOG("[BackgroundImgPlugin] Compiling BackgroundImgPlugin_Apply Kernel"); applyKernel = new cl::Kernel(*program, "BackgroundImgPlugin_Apply"); // Set kernel arguments u_int argIndex = 0; applyKernel->setArg(argIndex++, film.GetWidth()); applyKernel->setArg(argIndex++, film.GetHeight()); applyKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); applyKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); applyKernel->setArg(argIndex++, *(film.ocl_ALPHA)); applyKernel->setArg(argIndex++, *oclFilmImageMapDesc); applyKernel->setArg(argIndex++, *oclFilmImageMap); //---------------------------------------------------------------------- delete program; // Because imgMapDesc is a local variable oclIntersectionDevice->GetOpenCLQueue().flush(); const double tEnd = WallClockTime(); SLG_LOG("[BackgroundImgPlugin] Kernels compilation time: " << int((tEnd - tStart) * 1000.0) << "ms"); } oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*applyKernel, cl::NullRange, cl::NDRange(RoundUp(film.GetWidth() * film.GetHeight(), 256u)), cl::NDRange(256)); }
void Reinhard02ToneMap::ApplyOCL(Film &film, const u_int index) { const u_int pixelCount = film.GetWidth() * film.GetHeight(); const u_int workSize = RoundUp((pixelCount + 1) / 2, 64u); if (!applyKernel) { // Allocate buffers oclIntersectionDevice = film.oclIntersectionDevice; film.ctx->SetVerbose(true); oclIntersectionDevice->AllocBufferRW(&oclAccumBuffer, (workSize / 64) * sizeof(float) * 3, "Accumulation buffer"); film.ctx->SetVerbose(false); // Compile sources const double tStart = WallClockTime(); cl::Program *program = ImagePipelinePlugin::CompileProgram( film, "-D LUXRAYS_OPENCL_KERNEL -D SLG_OPENCL_KERNEL", slg::ocl::KernelSource_luxrays_types + slg::ocl::KernelSource_color_types + slg::ocl::KernelSource_color_funcs + slg::ocl::KernelSource_tonemap_reinhard02_funcs + slg::ocl::KernelSource_tonemap_reduce_funcs, "Reinhard02ToneMap"); SLG_LOG("[Reinhard02ToneMap] Compiling OpRGBValuesReduce Kernel"); opRGBValuesReduceKernel = new cl::Kernel(*program, "OpRGBValuesReduce"); SLG_LOG("[Reinhard02ToneMap] Compiling OpRGBValueAccumulate Kernel"); opRGBValueAccumulateKernel = new cl::Kernel(*program, "OpRGBValueAccumulate"); SLG_LOG("[Reinhard02ToneMap] Compiling Reinhard02ToneMap_Apply Kernel"); applyKernel = new cl::Kernel(*program, "Reinhard02ToneMap_Apply"); delete program; // Set kernel arguments u_int argIndex = 0; opRGBValuesReduceKernel->setArg(argIndex++, film.GetWidth()); opRGBValuesReduceKernel->setArg(argIndex++, film.GetHeight()); opRGBValuesReduceKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); opRGBValuesReduceKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); opRGBValuesReduceKernel->setArg(argIndex++, *oclAccumBuffer); argIndex = 0; opRGBValueAccumulateKernel->setArg(argIndex++, workSize / 64); opRGBValueAccumulateKernel->setArg(argIndex++, *oclAccumBuffer); argIndex = 0; applyKernel->setArg(argIndex++, film.GetWidth()); applyKernel->setArg(argIndex++, film.GetHeight()); applyKernel->setArg(argIndex++, *(film.ocl_IMAGEPIPELINE)); applyKernel->setArg(argIndex++, *(film.ocl_FRAMEBUFFER_MASK)); const float gamma = GetGammaCorrectionValue(film, index); applyKernel->setArg(argIndex++, gamma); applyKernel->setArg(argIndex++, preScale); applyKernel->setArg(argIndex++, postScale); applyKernel->setArg(argIndex++, burn); applyKernel->setArg(argIndex++, *oclAccumBuffer); const double tEnd = WallClockTime(); SLG_LOG("[Reinhard02ToneMap] Kernels compilation time: " << int((tEnd - tStart) * 1000.0) << "ms"); } film.oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*opRGBValuesReduceKernel, cl::NullRange, cl::NDRange(workSize), cl::NDRange(64)); film.oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*opRGBValueAccumulateKernel, cl::NullRange, cl::NDRange(64), cl::NDRange(64)); film.oclIntersectionDevice->GetOpenCLQueue().enqueueNDRangeKernel(*applyKernel, cl::NullRange, cl::NDRange(RoundUp(pixelCount, 256u)), cl::NDRange(256)); }