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