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));
}
