void VolumeMaxCLProcessor::executeVolumeOperation(const Volume* volume,
const VolumeCLBase* volumeCL,
VolumeCLBase* volumeOutCL, const size3_t& outDim,
const size3_t& globalWorkGroupSize,
const size3_t& localWorkgroupSize) {
cl::Event events[2];
try {
BufferCL* tmpVolumeCL;
int argIndex = 0;
kernel_->setArg(argIndex++, *volumeCL);
kernel_->setArg(argIndex++,
*(volumeCL->getVolumeStruct(volume)
.getRepresentation<BufferCL>())); // Scaling for 12-bit data
if (supportsVolumeWrite_) {
kernel_->setArg(argIndex++, *volumeOutCL);
} else {
size_t outDimFlattened = outDim.x * outDim.y * outDim.z;
if (tmpVolume_ == nullptr || tmpVolume_->getSize() != outDimFlattened) {
delete tmpVolume_;
tmpVolume_ = new Buffer<unsigned char>(outDimFlattened);
}
tmpVolumeCL = tmpVolume_->getEditableRepresentation<BufferCL>();
kernel_->setArg(argIndex++, *tmpVolumeCL);
}
kernel_->setArg(argIndex++, ivec4(outDim, 0));
kernel_->setArg(argIndex++, ivec4(volumeRegionSize_.get()));
OpenCL::getPtr()->getQueue().enqueueNDRangeKernel(
*kernel_, cl::NullRange, globalWorkGroupSize, localWorkgroupSize, nullptr, &events[0]);
if (!supportsVolumeWrite_) {
std::vector<cl::Event> waitFor(1, events[0]);
OpenCL::getPtr()->getQueue().enqueueCopyBufferToImage(
tmpVolumeCL->get(), volumeOutCL->getEditable(), 0, size3_t(0), size3_t(outDim),
&waitFor, &events[1]);
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
#if IVW_PROFILING
try {
if (supportsVolumeWrite_) {
events[0].wait();
LogInfo("Exec time: " << events[0].getElapsedTime() << " ms");
} else {
// Measure both computation and copy (only need to wait for copy)
events[1].wait();
LogInfo("Exec time (computation, copy): "
<< events[0].getElapsedTime() << " + " << events[1].getElapsedTime() << " = "
<< events[0].getElapsedTime() + events[1].getElapsedTime() << " ms");
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
#endif
}
void OpenCLKernel::run(int numDimensions, size_t *globalSize, size_t *localSize, cl_uint eventsInWaitList_, const cl_event* eventWaitList_, cl_event* runEvent_) {
if (clKernel== NULL) return;
cl_int err=CL_SUCCESS;
bindOpenGLInterOp();
err = clEnqueueNDRangeKernel(pOpenCL->getQueue(), clKernel, numDimensions, NULL, globalSize, localSize, eventsInWaitList_, eventWaitList_, runEvent_);
if (err != CL_SUCCESS) {
ofLogNotice() << getCLErrorString(err);
}
unbindOpenGLInterOp();
}
void VolumeRaycasterCL::samplingRate(float samplingRate) {
samplingRate_ = samplingRate;
if (kernel_) {
try {
kernel_->setArg(7, samplingRate);
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
}
bool OpenCLKernel::setArg(int argNumber, void* argp_, size_t size_){
if ( !clKernel ) return false;
// ----------| invariant: we have a valid kernel.
cl_int err = clSetKernelArg(clKernel, argNumber, size_, argp_);
if (err != CL_SUCCESS) {
ofLogNotice() << getCLErrorString(err);
}
return (err == CL_SUCCESS);
}
void VolumeRaycasterCL::outputSize(ivec2 val) {
if (kernel_) {
try {
kernel_->setArg(11, val);
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
outputSize_ = val;
}
void DirectionalLightSamplerCL::sampleLightSource(const Mesh* mesh, LightSamples& lightSamplesOut, const VECTOR_CLASS<cl::Event>* waitForEvents /*= nullptr*/, cl::Event* event /*= nullptr*/) {
const LightSource* light = lightSource_.get();
const BufferRAMPrecision<vec3>* vertices = dynamic_cast<const BufferRAMPrecision<vec3>*>(mesh->getBuffer(0)->getRepresentation<BufferRAM>());
if (vertices == nullptr || sampleGenerator_ == nullptr) {
return ;
}
if (samples_.getSize() != lightSamplesOut.getSize()) {
samples_.setSize(lightSamplesOut.getSize());
}
std::vector<cl::Event> sampleGenEvents(1);
sampleGenerator_->setUseGLSharing(false);
sampleGenerator_->generateNextSamples(samples_, waitForEvents, &sampleGenEvents[0]);
//const DirectionalLight* light = lights_.getData().get();
PackedLightSource lightBase = baseLightToPackedLight(light, 1.f, mesh->getCoordinateTransformer().getWorldToDataMatrix());
vec3 lightDirection = glm::normalize((lightBase.tm * vec4(0.f, 0.f, 1.f, 0.f)).xyz());
vec3 u, v;
vec3 lightOrigin{ (lightBase.tm*vec4(0.f, 0.f, 0.f, 1.f)).xyz() };
std::tie(lightOrigin, u, v) = geometry::fitPlaneAlignedOrientedBoundingBox2D(*vertices->getDataContainer(), Plane(lightOrigin.xyz(), lightDirection.xyz()));
float area = glm::length(u) * glm::length(v);
//LogInfo("Bounding box center: " << lightOrigin + 0.5f*(u + v));
//LogInfo("direction, o, lightU, lightV:" << lightDirection << o << u << v);
bool useGLSharing = true;
IVW_OPENCL_PROFILING(profilingEvent, "Light sampling")
//IVW_OPENCL_PROFILING(intersectionEvent, "Intersection computation")
try {
auto samplesCL = samples_.getRepresentation<BufferCL>();
if (useGLSharing) {
SyncCLGL glSync;
BufferCLGL* lightSamplesCL = lightSamplesOut.getLightSamples()->getEditableRepresentation<BufferCLGL>();
// Acquire shared representations before using them in OpenGL
// The SyncCLGL object will take care of synchronization between OpenGL and OpenCL
glSync.addToAquireGLObjectList(lightSamplesCL);
glSync.aquireAllObjects();
sampleLightSource(samplesCL, lightBase.radiance.xyz(), lightDirection, lightOrigin, u, v, area, samples_.getSize(), lightSamplesCL, &sampleGenEvents, profilingEvent);
} else {
BufferCL* lightSamplesCL = lightSamplesOut.getLightSamples()->getEditableRepresentation<BufferCL>();
sampleLightSource(samplesCL, lightBase.radiance.xyz(), lightDirection, lightOrigin, u, v, area, samples_.getSize(), lightSamplesCL, &sampleGenEvents, profilingEvent);
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
};
lightSamplesOut.advanceIteration();
}
bool RunningImageMeanAndStandardDeviationCL::computeMeanAndStandardDeviation(const Layer* newSamples, int iteration, Layer*& outMean, Layer*& outStandardDeviation, bool useGLSharing, const VECTOR_CLASS<cl::Event> *waitForEvents, cl::Event *event) {
if (kernel_ == nullptr) {
return false;
}
if (glm::any(glm::notEqual(newSamples->getDimensions(), standardDeviation_[0].getDimensions()))) {
standardDeviation_[0].setDimensions(newSamples->getDimensions());
standardDeviation_[1].setDimensions(newSamples->getDimensions());
mean_[0].setDimensions(newSamples->getDimensions());
mean_[1].setDimensions(newSamples->getDimensions());
}
//IVW_OPENCL_PROFILING(profilingEvent, "")
int prevStdId = pingPongIndex_;
int nextStdId = (pingPongIndex_ + 1) % 2;
try {
if (useGLSharing) {
SyncCLGL glSync;
const LayerCLGL* samples = newSamples->getRepresentation<LayerCLGL>();
LayerCLGL* prevMeanCL = mean_[prevStdId].getEditableRepresentation<LayerCLGL>();
LayerCLGL* nextMeanCL = mean_[nextStdId].getEditableRepresentation<LayerCLGL>();
LayerCLGL* prevStandardDeviation = standardDeviation_[prevStdId].getEditableRepresentation<LayerCLGL>();
LayerCLGL* nextStandardDeviation = standardDeviation_[nextStdId].getEditableRepresentation<LayerCLGL>();
// Acquire shared representations before using them in OpenGL
// The SyncCLGL object will take care of synchronization between OpenGL and OpenCL
glSync.addToAquireGLObjectList(samples);
glSync.addToAquireGLObjectList(prevMeanCL);
glSync.addToAquireGLObjectList(nextMeanCL);
glSync.addToAquireGLObjectList(prevStandardDeviation);
glSync.addToAquireGLObjectList(nextStandardDeviation);
glSync.aquireAllObjects();
computeMeanAndStandardDeviation(newSamples->getDimensions(), samples, iteration, prevMeanCL, nextMeanCL, prevStandardDeviation, nextStandardDeviation, workGroupSize_, waitForEvents, event);
} else {
LayerCL* prevMeanCL = mean_[prevStdId].getEditableRepresentation<LayerCL>();
LayerCL* nextMeanCL = mean_[nextStdId].getEditableRepresentation<LayerCL>();
const LayerCL* samples = newSamples->getRepresentation<LayerCL>();
LayerCL* prevStandardDeviation = standardDeviation_[prevStdId].getEditableRepresentation<LayerCL>();
LayerCL* nextStandardDeviation = standardDeviation_[nextStdId].getEditableRepresentation<LayerCL>();
computeMeanAndStandardDeviation(newSamples->getDimensions(), samples, iteration, prevMeanCL, nextMeanCL, prevStandardDeviation, nextStandardDeviation, workGroupSize_, waitForEvents, event);
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
return false;
}
pingPongIndex_ = nextStdId;
outMean = &mean_[nextStdId];
outStandardDeviation = &standardDeviation_[nextStdId];
return true;
}
void MWC64XSeedGenerator::generateSeeds(BufferCLBase* randomSeedBufferCL, int nRandomSeeds, size_t localWorkGroupSize) {
try
{
kernel_->setArg(0, *randomSeedBufferCL);
kernel_->setArg(1, nRandomSeeds);
size_t globalWorkSizeX = getGlobalWorkGroupSize(nRandomSeeds, localWorkGroupSize);
OpenCL::getPtr()->getQueue().enqueueNDRangeKernel(*kernel_, 0, globalWorkSizeX, localWorkGroupSize);
}
catch (cl::Error& err)
{
LogError(getCLErrorString(err));
}
}
ScopedClockCL::~ScopedClockCL() {
try {
profilingEvent_->wait();
std::stringstream message;
message << logMessage_ << ": " << profilingEvent_->getElapsedTime() << " ms";
LogCentral::getPtr()->log(logSource_, LogLevel::Info, LogAudience::Developer, __FILE__,
__FUNCTION__, __LINE__, message.str());
// LogInfo("Exec time: " << profilingEvent->getElapsedTime() << " ms");
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
delete profilingEvent_;
}
bool LayerCLGL::copyRepresentationsTo(DataRepresentation* targetRep) const {
// ivwAssert(false, "Not implemented");
// Make sure that the OpenCL layer is deleted before resizing the texture
// TODO: Implement copying in addition to the resizing
LayerCLGL* target = dynamic_cast<LayerCLGL*>(targetRep);
const LayerCLGL* source = this;
try {
SyncCLGL glSync;
glSync.addToAquireGLObjectList(target);
glSync.addToAquireGLObjectList(source);
glSync.aquireAllObjects();
LayerCLResizer::resize(source->get(), target->get(), target->getDimensions());
} catch (cl::Error err) {
LogError(getCLErrorString(err));
return false;
}
return true;
}
void VolumeRaycasterCLProcessor::process() {
try {
// This macro will create an event called profilingEvent if IVW_PROFILING is enabled,
// otherwise the profilingEvent will be declared as a null pointer
IVW_OPENCL_PROFILING(profilingEvent, "")
if (backgroundPort_.isReady()) {
volumeRaycaster_.setBackground(backgroundPort_.getData()->getColorLayer());
} else {
// Use default background
volumeRaycaster_.setBackground(nullptr);
}
volumeRaycaster_.outputSize(outport_.getDimensions());
volumeRaycaster_.volumeRaycast(
volumePort_.getData().get(), entryPort_.getData()->getColorLayer(),
exitPort_.getData()->getColorLayer(), transferFunction_.get().getData(),
outport_.getEditableData()->getColorLayer(), nullptr, profilingEvent);
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
void VolumeRaycasterCL::setLightingProperties(ShadingMode::Modes mode, const vec3& lightPosition,
const vec3& ambientColor, const vec3& diffuseColor,
const vec3& specularColor, float specularExponent) {
light_.position = vec4(lightPosition, 1.f);
light_.ambientColor = vec4(ambientColor, 1.f);
light_.diffuseColor = vec4(diffuseColor, 1.f);
light_.specularColor = vec4(specularColor, 1.f);
light_.specularExponent = specularExponent;
if (mode != light_.shadingMode) {
light_.shadingMode = mode;
compileKernel();
}
if (kernel_) {
try {
// Update data before returning it
lightStruct_.upload(&light_, sizeof(utilcl::LightParameters));
kernel_->setArg(8, lightStruct_);
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
}
void GrayscaleCLProcessor::process() {
if (kernel_ == nullptr) {
return;
}
auto outImage = outport_.getEditableData();
// outImage->resize(inImage->getDimensions());
uvec2 outportDim = outImage->getDimensions();
auto inImage = input_.getData();
try {
if (useGLSharing_.get()) {
SyncCLGL glSync;
const ImageCLGL* colorImageCL = inImage->getRepresentation<ImageCLGL>();
ImageCLGL* outImageCL = outImage->getEditableRepresentation<ImageCLGL>();
glSync.addToAquireGLObjectList(colorImageCL);
glSync.addToAquireGLObjectList(outImageCL);
glSync.aquireAllObjects();
cl_uint arg = 0;
kernel_->setArg(arg++, *colorImageCL);
kernel_->setArg(arg++, *outImageCL);
OpenCL::getPtr()->getQueue().enqueueNDRangeKernel(
*kernel_, cl::NullRange, static_cast<glm::size2_t>(outportDim));
} else {
const ImageCL* colorImageCL = inImage->getRepresentation<ImageCL>();
ImageCL* outImageCL = outImage->getEditableRepresentation<ImageCL>();
cl_uint arg = 0;
kernel_->setArg(arg++, *colorImageCL);
kernel_->setArg(arg++, *outImageCL);
OpenCL::getPtr()->getQueue().enqueueNDRangeKernel(
*kernel_, cl::NullRange, static_cast<glm::size2_t>(outportDim));
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
void MeshEntryExitPointsCL::computeEntryExitPoints(
const mat4& NDCToTextureMat, const mat4& worldToTextureMat, const BufferCLBase* vertices,
const BufferCLBase* indices, int nIndices, const LayerCLBase* entryPointsCL,
const LayerCLBase* exitPointsCL, const uvec2& outportDim,
const VECTOR_CLASS<cl::Event>* waitForEvents /*= nullptr*/, cl::Event* event /*= nullptr*/) {
size2_t localWorkGroupSize(workGroupSize_);
size2_t globalWorkGroupSize(getGlobalWorkGroupSize(outportDim.x, localWorkGroupSize.x),
getGlobalWorkGroupSize(outportDim.y, localWorkGroupSize.y));
try {
cl_uint arg = 0;
kernel_->setArg(arg++, NDCToTextureMat);
kernel_->setArg(arg++, worldToTextureMat);
kernel_->setArg(arg++, *vertices);
kernel_->setArg(arg++, *indices);
kernel_->setArg(arg++, nIndices);
kernel_->setArg(arg++, *entryPointsCL);
kernel_->setArg(arg++, *exitPointsCL);
OpenCL::getPtr()->getQueue().enqueueNDRangeKernel(
*kernel_, cl::NullRange, globalWorkGroupSize, localWorkGroupSize, waitForEvents, event);
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
// reimplemented from QApplication so we can throw exceptions in slots
virtual bool notify(QObject *receiver, QEvent *event) {
try {
return QApplication::notify(receiver, event);
} catch(Exception &e) {
Reporter::error() << "FAST exception caught in Qt event handler " << e.what() << Reporter::end();
throw e;
} catch(cl::Error &e) {
Reporter::error() << "OpenCL exception caught in Qt event handler " << e.what() << "(" << getCLErrorString(e.err()) << ")" << Reporter::end();
throw e;
} catch(std::exception &e) {
Reporter::error() << "Std exception caught in Qt event handler " << e.what() << Reporter::end();
throw e;
}
return false;
}
void PhotonTracerCL::tracePhotons(const Volume* volume, const TransferFunction& transferFunction, const BufferCL* axisAlignedBoundingBoxCL, const AdvancedMaterialProperty& material, const Camera* camera, float stepSize, const LightSamples* lightSamples, const Buffer<unsigned int>* photonsToRecomputeIndices, int nInvalidPhotons, int photonOffset, int batch, int maxInteractions, PhotonData* photonOutData, const VECTOR_CLASS<cl::Event> *waitForEvents, cl::Event *event /*= nullptr*/) {
if (!photonTracerKernel_) {
return;
}
if (randomState_.getSize() != photonOutData->getNumberOfPhotons()) {
setRandomSeedSize(photonOutData->getNumberOfPhotons());
}
auto volumeDim = volume->getDimensions();
// Texture space spacing
const mat4 volumeTextureToWorld = volume->getCoordinateTransformer().getTextureToWorldMatrix();
const mat4 textureToIndexMatrix = volume->getCoordinateTransformer().getTextureToIndexMatrix();
vec3 voxelSpacing(1.f / glm::length(textureToIndexMatrix[0]), 1.f / glm::length(textureToIndexMatrix[1]), 1.f / glm::length(textureToIndexMatrix[2]));
try {
if (useGLSharing_) {
SyncCLGL glSync;
auto volumeCL = volume->getRepresentation<VolumeCLGL>();
const BufferCLGL* lightSamplesCL = lightSamples->getLightSamples()->getRepresentation<BufferCLGL>();
const BufferCLGL* intersectionPointsCL = lightSamples->getIntersectionPoints()->getRepresentation<BufferCLGL>();
BufferCLGL* photonCL = photonOutData->photons_.getEditableRepresentation<BufferCLGL>();
const LayerCLGL* transferFunctionCL = transferFunction.getData()->getRepresentation<LayerCLGL>();
const ElementBufferCLGL* photonsToRecomputeIndicesCL = nullptr;
// Acquire shared representations before using them in OpenGL
// The SyncCLGL object will take care of synchronization between OpenGL and OpenCL
glSync.addToAquireGLObjectList(volumeCL);
glSync.addToAquireGLObjectList(lightSamplesCL);
glSync.addToAquireGLObjectList(intersectionPointsCL);
glSync.addToAquireGLObjectList(photonCL);
glSync.addToAquireGLObjectList(transferFunctionCL);
//{IVW_CPU_PROFILING("aquireAllObjects")
if (photonsToRecomputeIndices) {
photonsToRecomputeIndicesCL = photonsToRecomputeIndices->getRepresentation<ElementBufferCLGL>();
glSync.addToAquireGLObjectList(photonsToRecomputeIndicesCL);
}
glSync.aquireAllObjects();
//}
//{IVW_CPU_PROFILING("tracePhotons")
tracePhotons(photonOutData, volumeCL, volumeCL->getVolumeStruct(volume), axisAlignedBoundingBoxCL
, transferFunctionCL, material, stepSize, lightSamplesCL, intersectionPointsCL, lightSamples->getSize(), photonsToRecomputeIndicesCL, nInvalidPhotons
, photonCL, photonOffset, batch, maxInteractions
, waitForEvents, event);
//}
} else {
const VolumeCL* volumeCL = volume->getRepresentation<VolumeCL>();
const BufferCL* lightSamplesCL = lightSamples->getLightSamples()->getRepresentation<BufferCL>();
const BufferCL* intersectionPointsCL = lightSamples->getIntersectionPoints()->getRepresentation<BufferCL>();
BufferCL* photonCL = photonOutData->photons_.getEditableRepresentation<BufferCL>();
const LayerCL* transferFunctionCL = transferFunction.getData()->getRepresentation<LayerCL>();
const BufferCL* photonsToRecomputeIndicesCL = nullptr;
if (photonsToRecomputeIndices) {
photonsToRecomputeIndicesCL = photonsToRecomputeIndices->getRepresentation<BufferCL>();
}
tracePhotons(photonOutData, volumeCL, volumeCL->getVolumeStruct(volume), axisAlignedBoundingBoxCL
, transferFunctionCL, material, stepSize, lightSamplesCL, intersectionPointsCL, lightSamples->getSize(), photonsToRecomputeIndicesCL, nInvalidPhotons
, photonCL, photonOffset, batch, maxInteractions
, waitForEvents, event);
}
} catch (cl::Error& err) {
LogError(getCLErrorString(err));
}
}
