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