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