# HG changeset patch # User Sebastien Jodogne # Date 1717097497 -7200 # Node ID 93258466311b6c03018a7a07be12fe456ea415b1 # Parent 2e050ec8537965d198e41cf0be3d9f9716067e0e# Parent 2965172977a45c5ca201c2a0f38057d65c46c6e9 integration mainline->deep-learning diff -r 2965172977a4 -r 93258466311b Applications/Samples/WebAssembly/CMakeLists.txt --- a/Applications/Samples/WebAssembly/CMakeLists.txt Thu May 30 21:31:07 2024 +0200 +++ b/Applications/Samples/WebAssembly/CMakeLists.txt Thu May 30 21:31:37 2024 +0200 @@ -45,12 +45,15 @@ set(WASM_FLAGS "${WASM_FLAGS} -s SAFE_HEAP=1") endif() +set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ENVIRONMENT=web") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s EXPORTED_RUNTIME_METHODS='[\"ccall\", \"cwrap\"]'") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ERROR_ON_UNDEFINED_SYMBOLS=1") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ALLOW_MEMORY_GROWTH=1 -s TOTAL_MEMORY=268435456") # 256MB + resize set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s DISABLE_DEPRECATED_FIND_EVENT_TARGET_BEHAVIOR=1") add_definitions( -DDISABLE_DEPRECATED_FIND_EVENT_TARGET_BEHAVIOR=1 + -DORTHANC_HAS_WASM_SIMD=0 + -DORTHANC_WEBGL2_HEAP_COMPAT=0 ) diff -r 2965172977a4 -r 93258466311b Applications/StoneWebViewer/WebApplication/app.js --- a/Applications/StoneWebViewer/WebApplication/app.js Thu May 30 21:31:07 2024 +0200 +++ b/Applications/StoneWebViewer/WebApplication/app.js Thu May 30 21:31:37 2024 +0200 @@ -603,7 +603,11 @@ series: [], studies: [], seriesIndex: {}, // Maps "SeriesInstanceUID" to "index in this.series" - virtualSeriesThumbnails: {} + virtualSeriesThumbnails: {}, + + deepLearningReady: false, + deepLearningProgress: 0, // Floating-point number in the range [0..1] + deepLearningStartTime: null } }, computed: { @@ -1301,6 +1305,11 @@ }); }, + ApplyDeepLearning: function() { + stone.ApplyDeepLearningModel(this.GetActiveCanvas()); + app.deepLearningStartTime = performance.now(); + }, + ChangeActiveSeries: function(offset) { var seriesTags = this.GetActiveViewportSeriesTags(); if (seriesTags !== null) { @@ -1726,3 +1735,18 @@ } } }); + + +window.addEventListener('DeepLearningInitialized', function() { + stone.LoadDeepLearningModel('../stone-deep-learning/model.message'); +}); + +window.addEventListener('DeepLearningModelReady', function() { + app.deepLearningReady = true; + app.deepLearningProgress = 0; +}); + +window.addEventListener('DeepLearningStep', function(args) { + app.deepLearningProgress = args.detail.progress; + console.log('Elapsed time: ' + Math.round(performance.now() - app.deepLearningStartTime) + 'ms'); +}); diff -r 2965172977a4 -r 93258466311b Applications/StoneWebViewer/WebApplication/index.html --- a/Applications/StoneWebViewer/WebApplication/index.html Thu May 30 21:31:07 2024 +0200 +++ b/Applications/StoneWebViewer/WebApplication/index.html Thu May 30 21:31:37 2024 +0200 @@ -316,11 +316,23 @@
+
+
+ +
+ +
+
+
+
+
+
+
-
+
stoneAnnotations_; - + bool linearInterpolation_; + boost::shared_ptr deepLearningMask_; + std::string deepLearningSopInstanceUid_; + unsigned int deepLearningFrameNumber_; void ScheduleNextPrefetch() { @@ -2330,6 +2334,26 @@ } } + std::unique_ptr deepLearningLayer; + + if (deepLearningMask_.get() != NULL && + deepLearningSopInstanceUid_ == instance.GetSopInstanceUid() && + deepLearningFrameNumber_ == frameIndex) + { + std::vector lut(4 * 256); + for (unsigned int v = 128; v < 256; v++) + { + lut[4 * v] = 196; + lut[4 * v + 1] = 0; + lut[4 * v + 2] = 0; + lut[4 * v + 3] = 196; + } + + deepLearningLayer.reset(new OrthancStone::LookupTableTextureSceneLayer(*deepLearningMask_)); + deepLearningLayer->SetLookupTable(lut); + deepLearningLayer->SetPixelSpacing(pixelSpacingX, pixelSpacingY); + } + StoneAnnotationsRegistry::GetInstance().Load(*stoneAnnotations_, instance.GetSopInstanceUid(), frameIndex); // Orientation markers, new in Stone Web viewer 2.4 @@ -2406,6 +2430,15 @@ scene.DeleteLayer(LAYER_ORIENTATION_MARKERS); } + if (deepLearningLayer.get() != NULL) + { + scene.SetLayer(LAYER_DEEP_LEARNING, deepLearningLayer.release()); + } + else + { + scene.DeleteLayer(LAYER_DEEP_LEARNING); + } + stoneAnnotations_->Render(scene); // Necessary for "FitContent()" to work if (fitNextContent_) @@ -2432,7 +2465,7 @@ { const size_t cursorIndex = cursor_->GetCurrentIndex(); const OrthancStone::DicomInstanceParameters& instance = frames_->GetInstanceOfFrame(cursorIndex); - const size_t frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); + const unsigned int frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); // Only change the scene if the loaded frame still corresponds to the current cursor if (instance.GetSopInstanceUid() == loadedSopInstanceUid && @@ -2739,7 +2772,7 @@ { const size_t cursorIndex = cursor_->GetCurrentIndex(); const OrthancStone::DicomInstanceParameters& instance = frames_->GetInstanceOfFrame(cursorIndex); - const size_t frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); + const unsigned int frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); StoneAnnotationsRegistry::GetInstance().Save(instance.GetSopInstanceUid(), frameNumber, *stoneAnnotations_); @@ -2965,7 +2998,7 @@ const size_t cursorIndex = cursor_->GetCurrentIndex(); const OrthancStone::DicomInstanceParameters& instance = frames_->GetInstanceOfFrame(cursorIndex); - const size_t frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); + const unsigned int frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); FramesCache::Accessor accessor(*framesCache_, instance.GetSopInstanceUid(), frameNumber); if (accessor.IsValid()) @@ -3539,7 +3572,7 @@ { const size_t cursorIndex = cursor_->GetCurrentIndex(); const OrthancStone::DicomInstanceParameters& instance = frames_->GetInstanceOfFrame(cursorIndex); - const size_t frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); + const unsigned int frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); if (instance.GetSopInstanceUid() == sopInstanceUid && frameNumber == frame) @@ -3555,6 +3588,7 @@ } } + void SetLinearInterpolation(bool linearInterpolation) { if (linearInterpolation_ != linearInterpolation) @@ -3564,6 +3598,7 @@ } } + void AddTextAnnotation(const std::string& label, const OrthancStone::ScenePoint2D& pointedPosition, const OrthancStone::ScenePoint2D& labelPosition) @@ -3573,6 +3608,43 @@ } + bool GetCurrentFrame(std::string& sopInstanceUid /* out */, + unsigned int& frameNumber /* out */) const + { + if (cursor_.get() != NULL && + frames_.get() != NULL) + { + const size_t cursorIndex = cursor_->GetCurrentIndex(); + const OrthancStone::DicomInstanceParameters& instance = frames_->GetInstanceOfFrame(cursorIndex); + sopInstanceUid = instance.GetSopInstanceUid(); + frameNumber = frames_->GetFrameNumberInInstance(cursorIndex); + return true; + } + else + { + return false; + } + } + + + void SetDeepLearningMask(const std::string& sopInstanceUid, + unsigned int frameNumber, + const Orthanc::ImageAccessor& mask) + { + std::string currentSopInstanceUid; + unsigned int currentFrameNumber; + if (GetCurrentFrame(currentSopInstanceUid, currentFrameNumber) && + sopInstanceUid == currentSopInstanceUid && + frameNumber == currentFrameNumber) + { + deepLearningSopInstanceUid_ = sopInstanceUid; + deepLearningFrameNumber_ = frameNumber; + deepLearningMask_.reset(Orthanc::Image::Clone(mask)); + Redraw(); + } + } + + void SignalSynchronizedBrowsing() { if (synchronizationEnabled_ && @@ -3938,6 +4010,218 @@ } +#include +#include + +enum DeepLearningState +{ + DeepLearningState_Waiting, + DeepLearningState_Pending, + DeepLearningState_Running +}; + +static DeepLearningState deepLearningState_ = DeepLearningState_Waiting; +static worker_handle deepLearningWorker_; +static std::string deepLearningPendingSopInstanceUid_; +static unsigned int deepLearningPendingFrameNumber_; + +// Forward declaration +static void DeepLearningCallback(char* data, + int size, + void* payload); + +static void SendRequestToWebWorker(const OrthancStone::Messages::Request& request) +{ + std::string s; + if (request.SerializeToString(&s) && + !s.empty()) + { + emscripten_call_worker(deepLearningWorker_, "Execute", &s[0], s.size(), DeepLearningCallback, NULL); + } + else + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot send command to the Web worker"); + } +} + +static void DeepLearningSchedule(const std::string& sopInstanceUid, + unsigned int frameNumber) +{ + if (deepLearningState_ == DeepLearningState_Waiting) + { + LOG(WARNING) << "Starting deep learning on: " << sopInstanceUid << " / " << frameNumber; + + FramesCache::Accessor accessor(*framesCache_, sopInstanceUid, frameNumber); + if (accessor.IsValid() && + accessor.GetImage().GetFormat() == Orthanc::PixelFormat_Float32) + { + const Orthanc::ImageAccessor& image = accessor.GetImage(); + + OrthancStone::Messages::Request request; + request.set_type(OrthancStone::Messages::RequestType::LOAD_IMAGE); + request.mutable_load_image()->set_sop_instance_uid(sopInstanceUid); + request.mutable_load_image()->set_frame_number(frameNumber); + request.mutable_load_image()->set_width(image.GetWidth()); + request.mutable_load_image()->set_height(image.GetHeight()); + + const unsigned int height = image.GetHeight(); + const unsigned int width = image.GetWidth(); + for (unsigned int y = 0; y < height; y++) + { + const float* p = reinterpret_cast(image.GetConstRow(y)); + for (unsigned int x = 0; x < width; x++, p++) + { + request.mutable_load_image()->mutable_values()->Add(*p); + } + } + + deepLearningState_ = DeepLearningState_Running; + SendRequestToWebWorker(request); + } + else + { + LOG(ERROR) << "Cannot access the frame content, maybe a color image?"; + + EM_ASM({ + const customEvent = document.createEvent("CustomEvent"); + customEvent.initCustomEvent("DeepLearningStep", false, false, + { "progress" : "0" }); + window.dispatchEvent(customEvent); + }); + } + } + else + { + deepLearningState_ = DeepLearningState_Pending; + deepLearningPendingSopInstanceUid_ = sopInstanceUid; + deepLearningPendingFrameNumber_ = frameNumber; + } +} + +static void DeepLearningNextStep() +{ + switch (deepLearningState_) + { + case DeepLearningState_Pending: + deepLearningState_ = DeepLearningState_Waiting; + DeepLearningSchedule(deepLearningPendingSopInstanceUid_, deepLearningPendingFrameNumber_); + break; + + case DeepLearningState_Running: + { + OrthancStone::Messages::Request request; + request.set_type(OrthancStone::Messages::RequestType::EXECUTE_STEP); + SendRequestToWebWorker(request); + break; + } + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, "Bad state for deep learning"); + } +} + +static void DeepLearningCallback(char* data, + int size, + void* payload) +{ + try + { + OrthancStone::Messages::Response response; + if (response.ParseFromArray(data, size)) + { + switch (response.type()) + { + case OrthancStone::Messages::ResponseType::INITIALIZED: + DISPATCH_JAVASCRIPT_EVENT("DeepLearningInitialized"); + break; + + case OrthancStone::Messages::ResponseType::PARSED_MODEL: + LOG(WARNING) << "Number of steps in the model: " << response.parse_model().number_of_steps(); + DISPATCH_JAVASCRIPT_EVENT("DeepLearningModelReady"); + break; + + case OrthancStone::Messages::ResponseType::LOADED_IMAGE: + DeepLearningNextStep(); + break; + + case OrthancStone::Messages::ResponseType::STEP_DONE: + { + EM_ASM({ + const customEvent = document.createEvent("CustomEvent"); + customEvent.initCustomEvent("DeepLearningStep", false, false, + { "progress" : $0 }); + window.dispatchEvent(customEvent); + }, + response.step().progress() + ); + + if (response.step().done()) + { + deepLearningState_ = DeepLearningState_Waiting; + + const unsigned int height = response.step().mask().height(); + const unsigned int width = response.step().mask().width(); + + LOG(WARNING) << "SUCCESS! Mask: " << width << "x" << height << " for frame " + << response.step().mask().sop_instance_uid() << " / " + << response.step().mask().frame_number(); + + Orthanc::Image mask(Orthanc::PixelFormat_Grayscale8, width, height, false); + + size_t pos = 0; + for (unsigned int y = 0; y < height; y++) + { + uint8_t* p = reinterpret_cast(mask.GetRow(y)); + for (unsigned int x = 0; x < width; x++, p++, pos++) + { + *p = response.step().mask().values(pos) ? 255 : 0; + } + } + + for (Viewports::iterator it = allViewports_.begin(); it != allViewports_.end(); ++it) + { + assert(it->second != NULL); + it->second->SetDeepLearningMask(response.step().mask().sop_instance_uid(), + response.step().mask().frame_number(), mask); + } + } + else + { + DeepLearningNextStep(); + } + + break; + } + + default: + LOG(ERROR) << "Unsupported response type from the deep learning worker"; + } + } + else + { + LOG(ERROR) << "Bad response received from the deep learning worker"; + } + } + EXTERN_CATCH_EXCEPTIONS; +} + +static void DeepLearningModelLoaded(emscripten_fetch_t *fetch) +{ + try + { + LOG(WARNING) << "Deep learning model loaded: " << fetch->numBytes; + + OrthancStone::Messages::Request request; + request.set_type(OrthancStone::Messages::RequestType::PARSE_MODEL); + request.mutable_parse_model()->mutable_content()->assign(fetch->data, fetch->numBytes); + + emscripten_fetch_close(fetch); // Don't use "fetch" below + SendRequestToWebWorker(request); + } + EXTERN_CATCH_EXCEPTIONS; +} + extern "C" { int main(int argc, char const *argv[]) @@ -3955,11 +4239,55 @@ framesCache_.reset(new FramesCache); osiriXAnnotations_.reset(new OrthancStone::OsiriX::CollectionOfAnnotations); + deepLearningWorker_ = emscripten_create_worker("../stone-deep-learning/DeepLearningWorker.js"); + emscripten_call_worker(deepLearningWorker_, "Initialize", NULL, 0, DeepLearningCallback, NULL); + DISPATCH_JAVASCRIPT_EVENT("StoneInitialized"); } EMSCRIPTEN_KEEPALIVE + void LoadDeepLearningModel(const char* uri) + { + try + { + LOG(WARNING) << "Loading deep learning model: " << uri; + + emscripten_fetch_attr_t attr; + emscripten_fetch_attr_init(&attr); + strcpy(attr.requestMethod, "GET"); + attr.attributes = EMSCRIPTEN_FETCH_LOAD_TO_MEMORY; + attr.onsuccess = DeepLearningModelLoaded; + attr.onerror = NULL; + emscripten_fetch(&attr, uri); + } + EXTERN_CATCH_EXCEPTIONS; + } + + + EMSCRIPTEN_KEEPALIVE + void ApplyDeepLearningModel(const char* canvas) + { + try + { + boost::shared_ptr viewport = GetViewport(canvas); + + std::string sopInstanceUid; + unsigned int frameNumber; + if (viewport->GetCurrentFrame(sopInstanceUid, frameNumber)) + { + DeepLearningSchedule(sopInstanceUid, frameNumber); + } + else + { + LOG(WARNING) << "No active frame"; + } + } + EXTERN_CATCH_EXCEPTIONS; + } + + + EMSCRIPTEN_KEEPALIVE void SetDicomWebRoot(const char* uri, int useRendered) { diff -r 2965172977a4 -r 93258466311b OrthancStone/Resources/CMake/OrthancStoneConfiguration.cmake --- a/OrthancStone/Resources/CMake/OrthancStoneConfiguration.cmake Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Resources/CMake/OrthancStoneConfiguration.cmake Thu May 30 21:31:37 2024 +0200 @@ -200,7 +200,6 @@ if (NOT ORTHANC_SANDBOXED AND ENABLE_THREADS AND ENABLE_WEB_CLIENT) list(APPEND ORTHANC_STONE_SOURCES ${ORTHANC_STONE_ROOT}/Loaders/GenericLoadersContext.cpp - ${ORTHANC_STONE_ROOT}/Loaders/GenericLoadersContext.h ${ORTHANC_STONE_ROOT}/Oracle/GenericOracleRunner.cpp ${ORTHANC_STONE_ROOT}/Oracle/ThreadedOracle.cpp ) @@ -210,7 +209,6 @@ if (ENABLE_PUGIXML) list(APPEND ORTHANC_STONE_SOURCES ${ORTHANC_STONE_ROOT}/Scene2D/OsiriXLayerFactory.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/OsiriXLayerFactory.h ${ORTHANC_STONE_ROOT}/Toolbox/OsiriX/AngleAnnotation.cpp ${ORTHANC_STONE_ROOT}/Toolbox/OsiriX/Annotation.cpp ${ORTHANC_STONE_ROOT}/Toolbox/OsiriX/ArrayValue.cpp @@ -237,36 +235,22 @@ ${ORTHANC_STONE_ROOT}/Fonts/TextBoundingBox.cpp ${ORTHANC_STONE_ROOT}/Loaders/BasicFetchingItemsSorter.cpp - ${ORTHANC_STONE_ROOT}/Loaders/BasicFetchingItemsSorter.h ${ORTHANC_STONE_ROOT}/Loaders/BasicFetchingStrategy.cpp - ${ORTHANC_STONE_ROOT}/Loaders/BasicFetchingStrategy.h ${ORTHANC_STONE_ROOT}/Loaders/DicomResourcesLoader.cpp ${ORTHANC_STONE_ROOT}/Loaders/DicomSource.cpp ${ORTHANC_STONE_ROOT}/Loaders/DicomStructureSetLoader.cpp - ${ORTHANC_STONE_ROOT}/Loaders/DicomStructureSetLoader.h ${ORTHANC_STONE_ROOT}/Loaders/DicomVolumeLoader.cpp - ${ORTHANC_STONE_ROOT}/Loaders/IFetchingItemsSorter.h - ${ORTHANC_STONE_ROOT}/Loaders/IFetchingStrategy.h ${ORTHANC_STONE_ROOT}/Loaders/LoadedDicomResources.cpp ${ORTHANC_STONE_ROOT}/Loaders/LoaderStateMachine.cpp - ${ORTHANC_STONE_ROOT}/Loaders/LoaderStateMachine.h ${ORTHANC_STONE_ROOT}/Loaders/OrthancMultiframeVolumeLoader.cpp - ${ORTHANC_STONE_ROOT}/Loaders/OrthancMultiframeVolumeLoader.h ${ORTHANC_STONE_ROOT}/Loaders/OracleScheduler.cpp ${ORTHANC_STONE_ROOT}/Loaders/OrthancSeriesVolumeProgressiveLoader.cpp - ${ORTHANC_STONE_ROOT}/Loaders/OrthancSeriesVolumeProgressiveLoader.h ${ORTHANC_STONE_ROOT}/Loaders/SeriesFramesLoader.cpp ${ORTHANC_STONE_ROOT}/Loaders/SeriesMetadataLoader.cpp ${ORTHANC_STONE_ROOT}/Loaders/SeriesOrderedFrames.cpp ${ORTHANC_STONE_ROOT}/Loaders/SeriesThumbnailsLoader.cpp - ${ORTHANC_STONE_ROOT}/Messages/ICallable.h - ${ORTHANC_STONE_ROOT}/Messages/IMessage.h - ${ORTHANC_STONE_ROOT}/Messages/IMessageEmitter.h ${ORTHANC_STONE_ROOT}/Messages/IObservable.cpp - ${ORTHANC_STONE_ROOT}/Messages/IObservable.h - ${ORTHANC_STONE_ROOT}/Messages/IObserver.h - ${ORTHANC_STONE_ROOT}/Messages/ObserverBase.h ${ORTHANC_STONE_ROOT}/Oracle/GetOrthancImageCommand.cpp ${ORTHANC_STONE_ROOT}/Oracle/GetOrthancWebViewerJpegCommand.cpp @@ -277,202 +261,102 @@ ${ORTHANC_STONE_ROOT}/Oracle/ParseDicomFromWadoCommand.cpp ${ORTHANC_STONE_ROOT}/Scene2D/AnnotationsSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/AnnotationsSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/ArrowSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/ArrowSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/CairoCompositor.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/CairoCompositor.h - ${ORTHANC_STONE_ROOT}/Scene2D/Color.h - ${ORTHANC_STONE_ROOT}/Scene2D/ColorSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/ColorTextureSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/ColorTextureSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/CopyStyleConfigurator.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/CopyStyleConfigurator.h ${ORTHANC_STONE_ROOT}/Scene2D/FloatTextureSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/FloatTextureSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/GrayscaleStyleConfigurator.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/GrayscaleStyleConfigurator.h ${ORTHANC_STONE_ROOT}/Scene2D/GrayscaleWindowingSceneTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/GrayscaleWindowingSceneTracker.h - ${ORTHANC_STONE_ROOT}/Scene2D/ICompositor.h - ${ORTHANC_STONE_ROOT}/Scene2D/ILayerStyleConfigurator.h - ${ORTHANC_STONE_ROOT}/Scene2D/ISceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/InfoPanelSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/InfoPanelSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/LookupTableStyleConfigurator.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/LookupTableStyleConfigurator.h ${ORTHANC_STONE_ROOT}/Scene2D/LookupTableTextureSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/LookupTableTextureSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/MacroSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/MacroSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/MagnifyingGlassTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/MagnifyingGlassTracker.h - ${ORTHANC_STONE_ROOT}/Scene2D/NullLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/PanSceneTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/PanSceneTracker.h ${ORTHANC_STONE_ROOT}/Scene2D/PointerEvent.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/PointerEvent.h ${ORTHANC_STONE_ROOT}/Scene2D/PolylineSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/PolylineSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/RotateSceneTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/RotateSceneTracker.h ${ORTHANC_STONE_ROOT}/Scene2D/Scene2D.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Scene2D.h ${ORTHANC_STONE_ROOT}/Scene2D/ScenePoint2D.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/ScenePoint2D.h ${ORTHANC_STONE_ROOT}/Scene2D/TextSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/TextSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/TextureBaseSceneLayer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/TextureBaseSceneLayer.h ${ORTHANC_STONE_ROOT}/Scene2D/ZoomSceneTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/ZoomSceneTracker.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoArrowRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoArrowRenderer.h - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoBaseRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoColorTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoColorTextureRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoFloatTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoFloatTextureRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoInfoPanelRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoInfoPanelRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoLookupTableTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoLookupTableTextureRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoPolylineRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoPolylineRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoTextRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CairoTextRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CompositorHelper.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/CompositorHelper.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/FixedPointAligner.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/FixedPointAligner.h - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/ICairoContextProvider.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/MacroLayerRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/MacroLayerRenderer.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/AngleMeasureTool.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/AngleMeasureTool.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateAngleMeasureCommand.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateAngleMeasureCommand.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateAngleMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateAngleMeasureTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateCircleMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateCircleMeasureTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateLineMeasureCommand.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateLineMeasureCommand.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateLineMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateLineMeasureTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/CreateMeasureTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditAngleMeasureCommand.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditAngleMeasureCommand.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditAngleMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditAngleMeasureTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditLineMeasureCommand.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditLineMeasureCommand.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditLineMeasureTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/EditLineMeasureTracker.h - ${ORTHANC_STONE_ROOT}/Scene2DViewport/IFlexiblePointerTracker.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/LayerHolder.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/LayerHolder.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/LineMeasureTool.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/LineMeasureTool.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureCommands.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureCommands.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureTool.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureTool.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureToolsToolbox.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureToolsToolbox.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureTrackers.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/MeasureTrackers.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/OneGesturePointerTracker.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/OneGesturePointerTracker.h - ${ORTHANC_STONE_ROOT}/Scene2DViewport/PredeclaredTypes.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/UndoStack.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/UndoStack.h ${ORTHANC_STONE_ROOT}/Scene2DViewport/ViewportController.cpp - ${ORTHANC_STONE_ROOT}/Scene2DViewport/ViewportController.h ${ORTHANC_STONE_ROOT}/StoneEnumerations.cpp - ${ORTHANC_STONE_ROOT}/StoneException.h ${ORTHANC_STONE_ROOT}/StoneInitialization.cpp ${ORTHANC_STONE_ROOT}/Toolbox/AffineTransform2D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/AffineTransform2D.h + ${ORTHANC_STONE_ROOT}/Toolbox/AlignedMatrix.cpp ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator1D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator1D.h ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator2D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator2D.h ${ORTHANC_STONE_ROOT}/Toolbox/CoordinateSystem3D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/CoordinateSystem3D.h ${ORTHANC_STONE_ROOT}/Toolbox/DicomInstanceParameters.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/DicomInstanceParameters.h ${ORTHANC_STONE_ROOT}/Toolbox/DicomStructureSet.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/DicomStructureSet.h ${ORTHANC_STONE_ROOT}/Toolbox/DynamicBitmap.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/DynamicBitmap.h ${ORTHANC_STONE_ROOT}/Toolbox/Extent2D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/Extent2D.h ${ORTHANC_STONE_ROOT}/Toolbox/FiniteProjectiveCamera.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/FiniteProjectiveCamera.h ${ORTHANC_STONE_ROOT}/Toolbox/GenericToolbox.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/GenericToolbox.h ${ORTHANC_STONE_ROOT}/Toolbox/GeometryToolbox.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/GeometryToolbox.h ${ORTHANC_STONE_ROOT}/Toolbox/ImageGeometry.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/ImageGeometry.h ${ORTHANC_STONE_ROOT}/Toolbox/ImageToolbox.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/ImageToolbox.h ${ORTHANC_STONE_ROOT}/Toolbox/Internals/BucketMapper.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/Internals/BucketMapper.h ${ORTHANC_STONE_ROOT}/Toolbox/Internals/OrientedIntegerLine2D.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/Internals/OrientedIntegerLine2D.h ${ORTHANC_STONE_ROOT}/Toolbox/Internals/RectanglesIntegerProjection.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/Internals/RectanglesIntegerProjection.h ${ORTHANC_STONE_ROOT}/Toolbox/LinearAlgebra.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/LinearAlgebra.h - ${ORTHANC_STONE_ROOT}/Toolbox/PixelTestPatterns.h ${ORTHANC_STONE_ROOT}/Toolbox/SegmentTree.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/SegmentTree.h ${ORTHANC_STONE_ROOT}/Toolbox/ShearWarpProjectiveTransform.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/ShearWarpProjectiveTransform.h ${ORTHANC_STONE_ROOT}/Toolbox/SlicesSorter.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/SlicesSorter.h ${ORTHANC_STONE_ROOT}/Toolbox/SortedFrames.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/SortedFrames.h - ${ORTHANC_STONE_ROOT}/Toolbox/SubpixelReader.h - ${ORTHANC_STONE_ROOT}/Toolbox/SubvoxelReader.h ${ORTHANC_STONE_ROOT}/Toolbox/TextRenderer.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/TextRenderer.h + ${ORTHANC_STONE_ROOT}/Toolbox/TimerLogger.cpp ${ORTHANC_STONE_ROOT}/Toolbox/UndoRedoStack.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/UndoRedoStack.h ${ORTHANC_STONE_ROOT}/Toolbox/UnionOfRectangles.cpp - ${ORTHANC_STONE_ROOT}/Toolbox/UnionOfRectangles.h ${ORTHANC_STONE_ROOT}/Viewport/DefaultViewportInteractor.cpp - ${ORTHANC_STONE_ROOT}/Viewport/IViewport.h ${ORTHANC_STONE_ROOT}/Viewport/ViewportLocker.cpp - ${ORTHANC_STONE_ROOT}/Volumes/IGeometryProvider.h ${ORTHANC_STONE_ROOT}/Volumes/IVolumeSlicer.cpp - ${ORTHANC_STONE_ROOT}/Volumes/IVolumeSlicer.h ${ORTHANC_STONE_ROOT}/Volumes/OrientedVolumeBoundingBox.cpp - ${ORTHANC_STONE_ROOT}/Volumes/OrientedVolumeBoundingBox.h ${ORTHANC_STONE_ROOT}/Volumes/VolumeImageGeometry.cpp - ${ORTHANC_STONE_ROOT}/Volumes/VolumeImageGeometry.h ${ORTHANC_STONE_ROOT}/Volumes/VolumeReslicer.cpp - ${ORTHANC_STONE_ROOT}/Volumes/VolumeReslicer.h ${ORTHANC_STONE_ROOT}/Volumes/VolumeSceneLayerSource.cpp - ${ORTHANC_STONE_ROOT}/Volumes/VolumeSceneLayerSource.h - ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImage.h ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImage.cpp - ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImage.h ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImageMPRSlicer.cpp - ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImageMPRSlicer.h ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImageReslicer.cpp - ${ORTHANC_STONE_ROOT}/Volumes/DicomVolumeImageReslicer.h ${ORTHANC_STONE_ROOT}/Volumes/ImageBuffer3D.cpp - ${ORTHANC_STONE_ROOT}/Volumes/ImageBuffer3D.h ${ORTHANC_STONE_ROOT}/Wrappers/CairoContext.cpp ${ORTHANC_STONE_ROOT}/Wrappers/CairoSurface.cpp @@ -494,44 +378,29 @@ if (ENABLE_OPENGL) list(APPEND ORTHANC_STONE_SOURCES - ${ORTHANC_STONE_ROOT}/Fonts/OpenGLTextCoordinates.h ${ORTHANC_STONE_ROOT}/Fonts/OpenGLTextCoordinates.cpp - ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLProgram.h ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLProgram.cpp - ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLShader.h ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLShader.cpp - ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLTexture.h ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLTexture.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/OpenGLCompositor.h + ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLTextureArray.cpp + ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLTextureVolume.cpp + ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLFramebuffer.cpp + ${ORTHANC_STONE_ROOT}/OpenGL/ImageProcessingProgram.cpp ${ORTHANC_STONE_ROOT}/Scene2D/OpenGLCompositor.cpp ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLAdvancedPolylineRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLAdvancedPolylineRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLArrowRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLArrowRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLBasicPolylineRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLBasicPolylineRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLColorTextureProgram.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLColorTextureProgram.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLColorTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLColorTextureRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLFloatTextureProgram.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLFloatTextureProgram.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLFloatTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLFloatTextureRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLInfoPanelRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLInfoPanelRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLLinesProgram.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLLinesProgram.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLLookupTableTextureRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLLookupTableTextureRenderer.h - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLShaderVersionDirective.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextProgram.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextProgram.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextRenderer.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextRenderer.h ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextureProgram.cpp - ${ORTHANC_STONE_ROOT}/Scene2D/Internals/OpenGLTextureProgram.h ) endif() diff -r 2965172977a4 -r 93258466311b OrthancStone/Resources/WebAssemblySharedLibrary/CMakeLists.txt --- a/OrthancStone/Resources/WebAssemblySharedLibrary/CMakeLists.txt Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Resources/WebAssemblySharedLibrary/CMakeLists.txt Thu May 30 21:31:37 2024 +0200 @@ -54,6 +54,8 @@ set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s DISABLE_DEPRECATED_FIND_EVENT_TARGET_BEHAVIOR=1") add_definitions( -DDISABLE_DEPRECATED_FIND_EVENT_TARGET_BEHAVIOR=1 + -DORTHANC_HAS_WASM_SIMD=0 + -DORTHANC_WEBGL2_HEAP_COMPAT=0 ) diff -r 2965172977a4 -r 93258466311b OrthancStone/Resources/WebAssemblyUnitTests/CMakeLists.txt --- a/OrthancStone/Resources/WebAssemblyUnitTests/CMakeLists.txt Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Resources/WebAssemblyUnitTests/CMakeLists.txt Thu May 30 21:31:37 2024 +0200 @@ -39,6 +39,7 @@ set(WASM_FLAGS "${WASM_FLAGS} -s SAFE_HEAP=1") endif() +set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ENVIRONMENT=web") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s EXPORTED_RUNTIME_METHODS='[\"ccall\", \"cwrap\"]'") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ERROR_ON_UNDEFINED_SYMBOLS=1") set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s ALLOW_MEMORY_GROWTH=1 -s TOTAL_MEMORY=268435456") # 256MB + resize @@ -46,6 +47,8 @@ add_definitions( -DDISABLE_DEPRECATED_FIND_EVENT_TARGET_BEHAVIOR=1 -DORTHANC_BUILD_UNIT_TESTS=1 + -DORTHANC_HAS_WASM_SIMD=0 + -DORTHANC_WEBGL2_HEAP_COMPAT=0 ) diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/ImageProcessingProgram.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/ImageProcessingProgram.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,213 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "ImageProcessingProgram.h" + +#include "OpenGLFramebuffer.h" + +#include + + +static const unsigned int DIMENSIONS = 2; // Number of dimensions (we draw in 2D) +static const unsigned int VERTICES = 6; // 2 triangles in 2D (each triangle has 3 vertices) + + +static const float TRIANGLES[DIMENSIONS * VERTICES] = { + // First triangle + -1.0f, -1.0f, + 1.0f, -1.0f, + -1.0f, 1.0f, + // Second triangle + -1.0f, 1.0f, + 1.0f, -1.0f, + 1.0f, 1.0f +}; + + +/** + * "varying" indicates variables that are shader by the vertex shader + * and the fragment shader. The reason for "v_position" is that + * "a_position" (position in the target frame buffer) ranges from -1 + * to 1, whereas texture samplers range from 0 to 1. + **/ +static const char* VERTEX_SHADER_2D = + "in vec2 a_position; \n" + "out vec2 v_position; \n" + "void main() { \n" + " v_position = (a_position + 1.0) / 2.0; \n" + " gl_Position = vec4(a_position, 0, 1.0); \n" + "} \n"; + + +/** + * VERTEX_SHADER_3D allows to sample a 3D texture by introducing the + * "u_z" uniform whose range is in [0,1] and that allows to scan a 3D + * texture along its Z axis. + **/ +static const char* VERTEX_SHADER_3D = + "in vec2 a_position; \n" + "out vec3 v_position; \n" + "uniform float u_z; \n" + "void main() { \n" + " v_position = vec3((a_position + 1.0) / 2.0, u_z); \n" + " gl_Position = vec4(a_position, u_z, 1.0); \n" + "} \n"; + + +namespace OrthancStone +{ + namespace OpenGL + { + void ImageProcessingProgram::SetupPosition() + { + glBindBuffer(GL_ARRAY_BUFFER, quad_vertexbuffer); + glBufferData(GL_ARRAY_BUFFER, sizeof(float) * DIMENSIONS * VERTICES, TRIANGLES, GL_STATIC_DRAW); + GLint positionLocation = program_.GetAttributeLocation("a_position"); + glVertexAttribPointer(positionLocation, 2, GL_FLOAT, GL_FALSE, 0, 0); + glEnableVertexAttribArray(positionLocation); + } + + + ImageProcessingProgram::ImageProcessingProgram(IOpenGLContext& context, + const std::string& fragmentShader, + bool addUniformZ) : + program_(context), + quad_vertexbuffer(0) + { + if (context.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + + context.MakeCurrent(); + + std::string version; + +#if ORTHANC_ENABLE_WASM == 1 + /** + * "#version 300 es" corresponds to: + * - OpenGL ES version 3.0: https://registry.khronos.org/OpenGL-Refpages/es3.0/ + * - WebGL version 2.0 + * - GLSL ES version 3.00.6 + * - Based on version GLSL version 3.0 + * + * Explanation for "highp": + * https://emscripetn.org/docs/optimizing/Optimizing-WebGL.html + * https://webglfundamentals.org/webgl/lessons/webgl-qna-when-to-choose-highp--mediump--lowp-in-shaders.html + **/ + version = ("#version 300 es\n" + "precision highp float;\n" + "precision highp sampler2D;\n" + "precision highp sampler2DArray;\n" + "precision highp sampler3D;\n"); +#else + /** + * "#version 130" corresponds to: + * - OpenGL version 3.0 + * - GLSL version 1.30.10 + **/ + version = "#version 130\n"; +#endif + + std::string vertexShader; + + if (addUniformZ) + { + vertexShader = version + VERTEX_SHADER_3D; + } + else + { + vertexShader = version + VERTEX_SHADER_2D; + } + + program_.CompileShaders(vertexShader, version + fragmentShader); + + glGenBuffers(1, &quad_vertexbuffer); + if (quad_vertexbuffer == 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot create OpenGL buffer"); + } + } + + + ImageProcessingProgram::~ImageProcessingProgram() + { + glDeleteBuffers(1, &quad_vertexbuffer); + } + + + void ImageProcessingProgram::Use(OpenGLTexture& target, + OpenGLFramebuffer& framebuffer, + bool checkStatus) + { + program_.Use(checkStatus); + framebuffer.SetTarget(target); + SetupPosition(); + } + + + void ImageProcessingProgram::Use(OpenGLTextureArray& target, + unsigned int targetLayer, + OpenGLFramebuffer& framebuffer, + bool checkStatus) + { + program_.Use(checkStatus); + framebuffer.SetTarget(target, targetLayer); + SetupPosition(); + } + + + void ImageProcessingProgram::Use(OpenGLTextureVolume& target, + unsigned int z, + OpenGLFramebuffer& framebuffer, + bool checkStatus) + { + program_.Use(checkStatus); + framebuffer.SetTarget(target, z); + SetupPosition(); + } + + + void ImageProcessingProgram::Render() + { + glClearColor(0.0, 0.0, 0.0, 1.0); + glClear(GL_COLOR_BUFFER_BIT); + +#if 1 + glDrawArrays(GL_TRIANGLES, 0, VERTICES); +#else + // Simpler, but not available in WebGL + glBegin(GL_QUADS); + glVertex2f(-1, 1); // vertex 1 + glVertex2f(-1, -1); // vertex 2 + glVertex2f( 1, -1); // vertex 3 + glVertex2f( 1, 1); // vertex 4 + glEnd(); +#endif + + ORTHANC_OPENGL_CHECK("glDrawArrays()"); + } + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/ImageProcessingProgram.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/ImageProcessingProgram.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,80 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include "OpenGLProgram.h" + + +namespace OrthancStone +{ + namespace OpenGL + { + class OpenGLFramebuffer; + class OpenGLTexture; + class OpenGLTextureArray; + class OpenGLTextureVolume; + + class ImageProcessingProgram : public boost::noncopyable + { + private: + OpenGLProgram program_; + GLuint quad_vertexbuffer; + + void SetupPosition(); + + public: + ImageProcessingProgram(IOpenGLContext& context, + const std::string& fragmentShader, + bool addUniformZ /* for 3D texture sampling */); + + ~ImageProcessingProgram(); + + void Use(OpenGLTexture& target, + OpenGLFramebuffer& framebuffer, + bool checkStatus); + + void Use(OpenGLTextureArray& target, + unsigned int targetLayer, + OpenGLFramebuffer& framebuffer, + bool checkStatus); + + void Use(OpenGLTextureVolume& volume, + unsigned int z, + OpenGLFramebuffer& framebuffer, + bool checkStatus); + + void Render(); + + GLint GetUniformLocation(const std::string& name) + { + return program_.GetUniformLocation(name); + } + + GLint GetAttributeLocation(const std::string& name) + { + return program_.GetAttributeLocation(name); + } + }; + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLFramebuffer.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLFramebuffer.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,313 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "OpenGLFramebuffer.h" + +#if defined(__EMSCRIPTEN__) +# if !defined(ORTHANC_WEBGL2_HEAP_COMPAT) +# error The macro ORTHANC_WEBGL2_HEAP_COMPAT must be defined +# endif +#endif + +#include "OpenGLTexture.h" +#include "OpenGLTextureArray.h" +#include "OpenGLTextureVolume.h" + +#include + + +namespace OrthancStone +{ + namespace OpenGL + { + void OpenGLFramebuffer::SetupTextureTarget() + { + GLenum drawBuffers[1] = { GL_COLOR_ATTACHMENT0 }; + glDrawBuffers(1, drawBuffers); + + if (glCheckFramebufferStatus(GL_FRAMEBUFFER) != GL_FRAMEBUFFER_COMPLETE) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Incomplete setup of an OpenGL framebuffer"); + } + } + + + void OpenGLFramebuffer::ReadContent(Orthanc::ImageAccessor& target) + { + if (target.GetPitch() != target.GetWidth() * Orthanc::GetBytesPerPixel(target.GetFormat()) || + target.GetBuffer() == NULL) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Image must have minimal pitch"); + } + + if (glCheckFramebufferStatus(GL_FRAMEBUFFER) == GL_FRAMEBUFFER_COMPLETE) + { + ORTHANC_OPENGL_CHECK("glCheckFramebufferStatus()"); + + glViewport(0, 0, target.GetWidth(), target.GetHeight()); + + GLenum sourceFormat, internalFormat, pixelType; + OpenGLTexture::ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, target.GetFormat()); + +#if defined(__EMSCRIPTEN__) && (ORTHANC_WEBGL2_HEAP_COMPAT == 1) + // Check out "OpenGLTexture.cpp" for an explanation + + int framebufferFormat, framebufferType; + glGetIntegerv(GL_IMPLEMENTATION_COLOR_READ_FORMAT, &framebufferFormat); + glGetIntegerv(GL_IMPLEMENTATION_COLOR_READ_TYPE, &framebufferType); + + switch (target.GetFormat()) + { + case Orthanc::PixelFormat_RGBA32: + if (sourceFormat != GL_RGBA || + internalFormat != GL_RGBA || + pixelType != GL_UNSIGNED_BYTE || + framebufferFormat != GL_RGBA || + framebufferType != GL_UNSIGNED_BYTE) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + + EM_ASM({ + var ptr = emscriptenWebGLGetTexPixelData(GLctx.UNSIGNED_BYTE, GLctx.RGBA, $1, $2, $0, GLctx.RGBA); + GLctx.readPixels(0, 0, $1, $2, GLctx.RGBA, GLctx.UNSIGNED_BYTE, ptr); + }, + target.GetBuffer(), // $0 + target.GetWidth(), // $1 + target.GetHeight()); // $2 + break; + + case Orthanc::PixelFormat_Float32: + // In Mozilla Firefox, "Float32" is not available as such. We + // have to download an RGBA image in Float32. + if (sourceFormat != GL_RED || + internalFormat != GL_R32F || + pixelType != GL_FLOAT || + framebufferType != GL_FLOAT) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + + switch (framebufferFormat) + { + case GL_RGBA: + // This is Mozilla Firefox + EM_ASM({ + var tmp = new Float32Array($1 * $2 * 4); + GLctx.readPixels(0, 0, $1, $2, GLctx.RGBA, GLctx.FLOAT, tmp); + + // From RGBA to RED + var ptr = emscriptenWebGLGetTexPixelData(GLctx.FLOAT, GLctx.RED, $1, $2, $0, GLctx.R32F); + for (var i = 0; i < $1 * $2; i++) { + ptr[i] = tmp[4 * i]; + } + }, + target.GetBuffer(), // $0 + target.GetWidth(), // $1 + target.GetHeight()); // $2 + break; + + case GL_RED: + // This is Chromium + EM_ASM({ + var ptr = emscriptenWebGLGetTexPixelData(GLctx.FLOAT, GLctx.RED, $1, $2, $0, GLctx.R32F); + GLctx.readPixels(0, 0, $1, $2, GLctx.RED, GLctx.FLOAT, ptr); + }, + target.GetBuffer(), // $0 + target.GetWidth(), // $1 + target.GetHeight()); // $2 + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + } + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + } +#else + glReadPixels(0, 0, target.GetWidth(), target.GetHeight(), sourceFormat, pixelType, target.GetBuffer()); +#endif + + ORTHANC_OPENGL_CHECK("glReadPixels()"); + } + else + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Incomplete setup of an OpenGL framebuffer"); + } + } + + + OpenGLFramebuffer::OpenGLFramebuffer(IOpenGLContext& context) : + context_(context), + framebuffer_(0) + { + if (context.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + + glGenFramebuffers(1, &framebuffer_); + if (framebuffer_ == 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot create an OpenGL framebuffer"); + } + + glBindFramebuffer(GL_FRAMEBUFFER, framebuffer_); + } + + + OpenGLFramebuffer::~OpenGLFramebuffer() + { + glDeleteFramebuffers(1, &framebuffer_); + } + + + void OpenGLFramebuffer::SetTarget(OpenGLTexture& target) + { + glFramebufferTexture2D(GL_DRAW_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, target.GetId(), 0); + ORTHANC_OPENGL_CHECK("glFramebufferTexture2D()"); + + SetupTextureTarget(); + glViewport(0, 0, target.GetWidth(), target.GetHeight()); + } + + + void OpenGLFramebuffer::SetTarget(OpenGLTextureArray& target, + unsigned int layer) + { + if (layer >= target.GetDepth()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else + { + glFramebufferTextureLayer(GL_DRAW_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, target.GetId(), 0, layer); + ORTHANC_OPENGL_CHECK("glFramebufferTextureLayer()"); + + SetupTextureTarget(); + glViewport(0, 0, target.GetWidth(), target.GetHeight()); + } + } + + + void OpenGLFramebuffer::SetTarget(OpenGLTextureVolume& target, + unsigned int z) + { + if (z >= target.GetDepth()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else + { + // Warning: "glFramebufferTexture3D()" is not available in WebGL 2 + glFramebufferTextureLayer(GL_DRAW_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, target.GetId(), 0, z); + ORTHANC_OPENGL_CHECK("glFramebufferTextureLayer()"); + + SetupTextureTarget(); + glViewport(0, 0, target.GetWidth(), target.GetHeight()); + } + } + + + void OpenGLFramebuffer::ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTexture& source) + { + if (target.GetWidth() != source.GetWidth() || + target.GetHeight() != source.GetHeight()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + else if (target.GetFormat() != source.GetFormat()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageFormat); + } + else + { + glFramebufferTexture2D(GL_READ_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, source.GetId(), 0); + ORTHANC_OPENGL_CHECK("glFramebufferTexture2D()"); + ReadContent(target); + } + } + + + void OpenGLFramebuffer::ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTextureArray& source, + unsigned int layer) + { + if (target.GetWidth() != source.GetWidth() || + target.GetHeight() != source.GetHeight()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + else if (target.GetFormat() != source.GetFormat()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageFormat); + } + else if (layer >= source.GetDepth()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else + { + glFramebufferTextureLayer(GL_READ_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, source.GetId(), 0, layer); + ORTHANC_OPENGL_CHECK("glFramebufferTextureLayer()"); + ReadContent(target); + } + } + + + void OpenGLFramebuffer::ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTextureVolume& source, + unsigned int z) + { + if (target.GetWidth() != source.GetWidth() || + target.GetHeight() != source.GetHeight()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + else if (target.GetFormat() != source.GetFormat()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageFormat); + } + else if (z >= source.GetDepth()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else + { + // Warning: "glFramebufferTexture3D()" is not available in WebGL 2 + glFramebufferTextureLayer(GL_READ_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, source.GetId(), 0, z); + ORTHANC_OPENGL_CHECK("glFramebufferTextureLayer()"); + + ReadContent(target); + } + } + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLFramebuffer.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLFramebuffer.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,74 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include "OpenGLIncludes.h" +#include "IOpenGLContext.h" + +#include + +namespace OrthancStone +{ + namespace OpenGL + { + class OpenGLTexture; + class OpenGLTextureArray; + class OpenGLTextureVolume; + + class OpenGLFramebuffer : public boost::noncopyable + { + private: + IOpenGLContext& context_; + GLuint framebuffer_; + + void SetupTextureTarget(); + + void ReadContent(Orthanc::ImageAccessor& target); + + public: + OpenGLFramebuffer(IOpenGLContext& context); + + ~OpenGLFramebuffer(); + + void SetTarget(OpenGLTexture& target); + + void SetTarget(OpenGLTextureArray& target, + unsigned int layer); + + void SetTarget(OpenGLTextureVolume& target, + unsigned int z); + + void ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTexture& source); + + void ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTextureArray& source, + unsigned int layer); + + void ReadTexture(Orthanc::ImageAccessor& target, + const OpenGLTextureVolume& source, + unsigned int z); + }; + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLProgram.cpp --- a/OrthancStone/Sources/OpenGL/OpenGLProgram.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLProgram.cpp Thu May 30 21:31:37 2024 +0200 @@ -82,10 +82,20 @@ } } - void OpenGLProgram::Use() + void OpenGLProgram::Use(bool checkStatus) { //ORTHANC_OPENGL_TRACE_CURRENT_CONTEXT("About to call glUseProgram"); + ORTHANC_OPENGL_CHECK("About to call glUseProgram()"); + glUseProgram(program_); + + if (checkStatus && + glGetError() != GL_NO_ERROR) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot use successfully compiled OpenGL shader"); + } + ORTHANC_OPENGL_CHECK("glUseProgram"); } diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLProgram.h --- a/OrthancStone/Sources/OpenGL/OpenGLProgram.h Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLProgram.h Thu May 30 21:31:37 2024 +0200 @@ -48,7 +48,9 @@ ~OpenGLProgram(); - void Use(); + // WARNING: Setting "checkStatus" to "true" impacts performance: + // Calling "glGetError()" seems like a costly operation in WebGL + void Use(bool checkStatus); // WARNING: A global OpenGL context must be active to run this method! void CompileShaders(const std::string& vertexCode, diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLShader.cpp --- a/OrthancStone/Sources/OpenGL/OpenGLShader.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLShader.cpp Thu May 30 21:31:37 2024 +0200 @@ -55,8 +55,10 @@ glCompileShader(shader); ORTHANC_OPENGL_CHECK("glCompileShader"); + GLenum error = glGetError(); + // Check if there were errors - int infoLen = 0; + GLint infoLen = 0; glGetShaderiv(shader, GL_INFO_LOG_LENGTH, &infoLen); ORTHANC_OPENGL_CHECK("glGetShaderiv"); @@ -65,18 +67,24 @@ std::string infoLog; infoLog.resize(infoLen + 1); glGetShaderInfoLog(shader, infoLen, NULL, &infoLog[0]); - ORTHANC_OPENGL_CHECK("glGetShaderInfoLog"); - ORTHANC_OPENGL_TRACE_CURRENT_CONTEXT("About to call glDeleteShader"); - glDeleteShader(shader); - ORTHANC_OPENGL_CHECK("glDeleteShader"); - throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, - "Error while creating an OpenGL shader: " + infoLog); + if (error) + { + ORTHANC_OPENGL_CHECK("glGetShaderInfoLog"); + ORTHANC_OPENGL_TRACE_CURRENT_CONTEXT("About to call glDeleteShader"); + glDeleteShader(shader); + ORTHANC_OPENGL_CHECK("glDeleteShader"); + + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Error while creating an OpenGL shader: " + infoLog); + } + else + { + LOG(WARNING) << "Warning while creating an OpenGL shader: " << infoLog; + } } - else - { - return shader; - } + + return shader; } } diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTexture.cpp --- a/OrthancStone/Sources/OpenGL/OpenGLTexture.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLTexture.cpp Thu May 30 21:31:37 2024 +0200 @@ -22,8 +22,14 @@ #include "OpenGLTexture.h" -#include "IOpenGLContext.h" +#if defined(__EMSCRIPTEN__) +# if !defined(ORTHANC_WEBGL2_HEAP_COMPAT) +# error The macro ORTHANC_WEBGL2_HEAP_COMPAT must be defined +# endif +#endif + +#include #include #include @@ -33,20 +39,28 @@ { namespace OpenGL { - OpenGLTexture::OpenGLTexture(OpenGL::IOpenGLContext& context) - : width_(0) - , height_(0) - , context_(context) + OpenGLTexture::OpenGLTexture(OpenGL::IOpenGLContext& context) : + context_(context), + texture_(0), + width_(0), + height_(0), + format_(Orthanc::PixelFormat_Grayscale8), + isLinearInterpolation_(false) { - if (!context_.IsContextLost()) + if (context.IsContextLost()) { - // Generate a texture object - glGenTextures(1, &texture_); - if (texture_ == 0) - { - throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, - "Cannot create an OpenGL program"); - } + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + + // Generate a texture object + glGenTextures(1, &texture_); + ORTHANC_OPENGL_CHECK("glGenTextures()"); + + if (texture_ == 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot create an OpenGL texture"); } } @@ -82,56 +96,90 @@ } } - void OpenGLTexture::Load(const Orthanc::ImageAccessor& image, - bool isLinearInterpolation) + void OpenGLTexture::Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + bool isLinearInterpolation, + const void* data) { - if (!context_.IsContextLost()) + if (context_.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + else { glPixelStorei(GL_UNPACK_ALIGNMENT, 1); // Disable byte-alignment restriction - if (image.GetPitch() != image.GetBytesPerPixel() * image.GetWidth()) - { - throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented, - "Unsupported non-zero padding"); - } - // Bind it glActiveTexture(GL_TEXTURE0); glBindTexture(GL_TEXTURE_2D, texture_); - GLenum sourceFormat, internalFormat; - - switch (image.GetFormat()) - { - case Orthanc::PixelFormat_Grayscale8: - sourceFormat = GL_RED; - internalFormat = GL_RED; - break; - - case Orthanc::PixelFormat_RGB24: - sourceFormat = GL_RGB; - internalFormat = GL_RGB; - break; + GLenum sourceFormat, internalFormat, pixelType; + ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, format); - case Orthanc::PixelFormat_RGBA32: - sourceFormat = GL_RGBA; - internalFormat = GL_RGBA; - break; - - default: - throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented, - "No support for this format in OpenGL textures: " + - std::string(EnumerationToString(image.GetFormat()))); - } - - width_ = image.GetWidth(); - height_ = image.GetHeight(); + format_ = format; + width_ = width; + height_ = height; + isLinearInterpolation_ = isLinearInterpolation; GLint interpolation = (isLinearInterpolation ? GL_LINEAR : GL_NEAREST); // Load the texture from the image buffer - glTexImage2D(GL_TEXTURE_2D, 0, internalFormat, image.GetWidth(), image.GetHeight(), - 0, sourceFormat, GL_UNSIGNED_BYTE, image.GetConstBuffer()); + +#if defined(__EMSCRIPTEN__) && (ORTHANC_WEBGL2_HEAP_COMPAT == 1) + /** + * This compatibility implementation seems to be necessary + * with WebGL2, at least in Web workers. In such a situation, + * the calls that are referred to as the "new garbage-free + * entry points" in the Emscripten source file + * "upstream/emscripten/src/library_webgl.js" seem to fail, + * because the "Uint8Array" and "Float32Array" seem to be + * incorrectly created. This compatibility reverts to the + * WebGL1 behavior of "library_webgl.js", which requires the + * function "emscriptenWebGLGetTexPixelData" that is defined + * in "upstream/emscripten/src/library_webgl.js" to be + * exported in the linker using option + * "EXTRA_EXPORTED_RUNTIME_METHODS" or + * "EXPORTED_RUNTIME_METHODS". + **/ + EM_ASM({ + var ptr = $0 ? emscriptenWebGLGetTexPixelData($5, $4, $2, $3, $0, $1) : null; + GLctx.texImage2D(GLctx.TEXTURE_2D, 0, $1, $2, $3, 0, $4, $5, ptr); + }, + data, // $0 + internalFormat, // $1 + width, // $2 + height, // $3 + sourceFormat, // $4 + pixelType); // $5 +#else + glTexImage2D(GL_TEXTURE_2D, 0, internalFormat, width, height, + 0, sourceFormat, pixelType, data); +#endif + + ORTHANC_OPENGL_CHECK("glTexImage2D()"); + +#if !defined(__EMSCRIPTEN__) + /** + * glGetTexLevelParameteriv() was introduced in OpenGL ES 3.1, + * but WebGL 2 only supports OpenGL ES 3.0, so it is not + * available in WebAssembly: + * https://registry.khronos.org/OpenGL-Refpages/es3.1/html/glGetTexLevelParameter.xhtml + **/ + GLint w, h; + glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_WIDTH, &w); + glGetTexLevelParameteriv(GL_TEXTURE_2D, 0, GL_TEXTURE_HEIGHT, &h); + if (width != static_cast(w) || + height != static_cast(h)) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Your GPU cannot create a texture of size " + + boost::lexical_cast(width) + " x " + + boost::lexical_cast(height)); + } +#endif + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, interpolation); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, interpolation); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); @@ -139,12 +187,157 @@ } } + void OpenGLTexture::Load(const Orthanc::ImageAccessor& image, + bool isLinearInterpolation) + { + if (image.GetPitch() != image.GetBytesPerPixel() * image.GetWidth()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented, + "Pitch is not the same as the row size"); + } + else + { + Setup(image.GetFormat(), image.GetWidth(), image.GetHeight(), + isLinearInterpolation, image.GetConstBuffer()); + } + } - void OpenGLTexture::Bind(GLint location) + + void OpenGLTexture::Bind(GLint location) const { glActiveTexture(GL_TEXTURE0); glBindTexture(GL_TEXTURE_2D, texture_); glUniform1i(location, 0 /* texture unit */); } + + + void OpenGLTexture::BindAsTextureUnit(GLint location, + unsigned int unit) const + { + if (unit >= 32) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + + assert(GL_TEXTURE0 + 1 == GL_TEXTURE1 && + GL_TEXTURE0 + 31 == GL_TEXTURE31); + + glActiveTexture(GL_TEXTURE0 + unit); + glBindTexture(GL_TEXTURE_2D, texture_); + glUniform1i(location, unit /* texture unit */); + } + + + Orthanc::ImageAccessor* OpenGLTexture::Download(Orthanc::PixelFormat format) const + { + if (context_.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context is lost"); + } + + std::unique_ptr target(new Orthanc::Image(format, width_, height_, true)); + assert(target->GetPitch() == width_ * Orthanc::GetBytesPerPixel(format)); + +#if defined(__EMSCRIPTEN__) + /** + * The "glGetTexImage()" is unavailable in WebGL, it is + * necessary to use a framebuffer: + * https://stackoverflow.com/a/15064957 + **/ + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + +#else + glBindTexture(GL_TEXTURE_2D, texture_); + + switch (format) + { + case Orthanc::PixelFormat_Grayscale8: + glGetTexImage(GL_TEXTURE_2D, 0 /* base level */, GL_RED, GL_UNSIGNED_BYTE, target->GetBuffer()); + break; + + case Orthanc::PixelFormat_RGB24: + glGetTexImage(GL_TEXTURE_2D, 0 /* base level */, GL_RGB, GL_UNSIGNED_BYTE, target->GetBuffer()); + break; + + case Orthanc::PixelFormat_RGBA32: + glGetTexImage(GL_TEXTURE_2D, 0 /* base level */, GL_RGBA, GL_UNSIGNED_BYTE, target->GetBuffer()); + break; + + case Orthanc::PixelFormat_Float32: + glGetTexImage(GL_TEXTURE_2D, 0 /* base level */, GL_RED, GL_FLOAT, target->GetBuffer()); + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + } +#endif + + return target.release(); + } + + + void OpenGLTexture::SetClampingToZero() + { +#if defined(__EMSCRIPTEN__) + /** + * This is because WebGL 2 derives from OpenGL ES 3.0, which + * doesn't support GL_CLAMP_TO_BORDER, as can be seen here: + * https://registry.khronos.org/OpenGL-Refpages/es3.0/html/glTexParameter.xhtml + **/ + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGLTextureArray::SetClampingToZero() is not available in WebGL 2"); +#else + ORTHANC_OPENGL_CHECK("Entering OpenGLTexture::SetClampingToZero()"); + + glBindTexture(GL_TEXTURE_2D, texture_); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); + + GLfloat colorfv[4] = { 0, 0, 0, 0 }; + glTextureParameterfv(texture_, GL_TEXTURE_BORDER_COLOR, colorfv); + + ORTHANC_OPENGL_CHECK("Exiting OpenGLTexture::SetClampingToZero()"); +#endif + } + + + void OpenGLTexture::ConvertToOpenGLFormats(GLenum& sourceFormat, + GLenum& internalFormat, + GLenum& pixelType, + Orthanc::PixelFormat format) + { + switch (format) + { + case Orthanc::PixelFormat_Grayscale8: + sourceFormat = GL_RED; + internalFormat = GL_RED; + pixelType = GL_UNSIGNED_BYTE; + break; + + case Orthanc::PixelFormat_RGB24: + sourceFormat = GL_RGB; + internalFormat = GL_RGB; + pixelType = GL_UNSIGNED_BYTE; + break; + + case Orthanc::PixelFormat_RGBA32: + sourceFormat = GL_RGBA; + internalFormat = GL_RGBA; + pixelType = GL_UNSIGNED_BYTE; + break; + + case Orthanc::PixelFormat_Float32: + sourceFormat = GL_RED; + internalFormat = GL_R32F; // Don't use "GL_RED" here, as it clamps to [0,1] + pixelType = GL_FLOAT; + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented, + "No support for this format in OpenGL textures: " + + std::string(EnumerationToString(format))); + } + } } } diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTexture.h --- a/OrthancStone/Sources/OpenGL/OpenGLTexture.h Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLTexture.h Thu May 30 21:31:37 2024 +0200 @@ -24,6 +24,7 @@ #pragma once #include "OpenGLIncludes.h" +#include "IOpenGLContext.h" #include @@ -34,21 +35,43 @@ { namespace OpenGL { - class IOpenGLContext; - class OpenGLTexture : public boost::noncopyable { + friend class OpenGLFramebuffer; + private: - GLuint texture_; - unsigned int width_; - unsigned int height_; OpenGL::IOpenGLContext& context_; + GLuint texture_; + unsigned int width_; + unsigned int height_; + Orthanc::PixelFormat format_; + bool isLinearInterpolation_; + + void Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + bool isLinearInterpolation, + const void* data); + + /** + * Returns the low-level OpenGL handle of the texture. Beware to + * never change the size of the texture using this handle! + **/ + GLuint GetId() const + { + return texture_; + } public: explicit OpenGLTexture(OpenGL::IOpenGLContext& context); ~OpenGLTexture(); + Orthanc::PixelFormat GetFormat() const + { + return format_; + } + unsigned int GetWidth() const { return width_; @@ -59,10 +82,39 @@ return height_; } + bool IsLinearInterpolation() const + { + return isLinearInterpolation_; + } + + void Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + bool isLinearInterpolation) + { + Setup(format, width, height, isLinearInterpolation, NULL); + } + void Load(const Orthanc::ImageAccessor& image, bool isLinearInterpolation); - void Bind(GLint location); + void Bind(GLint location) const; + + void BindAsTextureUnit(GLint location, + unsigned int unit) const; + + Orthanc::ImageAccessor* Download(Orthanc::PixelFormat format) const; + + /** + * By default, textures are mirrored at the borders. This + * function will set out-of-image access to zero. + **/ + void SetClampingToZero(); + + static void ConvertToOpenGLFormats(GLenum& sourceFormat, + GLenum& internalFormat, + GLenum& pixelType, + Orthanc::PixelFormat format); }; } } diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTextureArray.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureArray.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,308 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "OpenGLTextureArray.h" + +#if defined(__EMSCRIPTEN__) +# if !defined(ORTHANC_WEBGL2_HEAP_COMPAT) +# error The macro ORTHANC_WEBGL2_HEAP_COMPAT must be defined +# endif +#endif + +#include "OpenGLFramebuffer.h" +#include "OpenGLTexture.h" + +#include +#include + +#include +#include + +namespace OrthancStone +{ + namespace OpenGL + { + OpenGLTextureArray::OpenGLTextureArray(IOpenGLContext& context) : + context_(context), + texture_(0), + width_(0), + height_(0), + depth_(0), + format_(Orthanc::PixelFormat_Float32), + isLinearInterpolation_(false) + { + if (context.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + + glGenTextures(1, &texture_); + ORTHANC_OPENGL_CHECK("glGenTextures()"); + + if (texture_ == 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot create an OpenGL texture array"); + } + } + + + OpenGLTextureArray::~OpenGLTextureArray() + { + assert(texture_ != 0); + glDeleteTextures(1, &texture_); + } + + + void OpenGLTextureArray::Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + unsigned int depth, + bool isLinearInterpolation) + { + glActiveTexture(GL_TEXTURE0); + glBindTexture(GL_TEXTURE_2D_ARRAY, texture_); + ORTHANC_OPENGL_CHECK("glBindTexture(GL_TEXTURE_2D_ARRAY)"); + + GLenum sourceFormat, internalFormat, pixelType; + OpenGLTexture::ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, format); + + glTexImage3D(GL_TEXTURE_2D_ARRAY, 0, internalFormat, width, height, depth, + 0, sourceFormat, pixelType, NULL); + ORTHANC_OPENGL_CHECK("glTexImage3D()"); + +#if !defined(__EMSCRIPTEN__) + /** + * glGetTexLevelParameteriv() was introduced in OpenGL ES 3.1, + * but WebGL 2 only supports OpenGL ES 3.0, so it is not + * available in WebAssembly: + * https://registry.khronos.org/OpenGL-Refpages/es3.1/html/glGetTexLevelParameter.xhtml + **/ + GLint w, h, d; + glGetTexLevelParameteriv(GL_TEXTURE_2D_ARRAY, 0, GL_TEXTURE_WIDTH, &w); + glGetTexLevelParameteriv(GL_TEXTURE_2D_ARRAY, 0, GL_TEXTURE_HEIGHT, &h); + glGetTexLevelParameteriv(GL_TEXTURE_2D_ARRAY, 0, GL_TEXTURE_DEPTH, &d); + if (width != static_cast(w) || + height != static_cast(h) || + depth != static_cast(d)) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Your GPU cannot create an array of textures of size " + + boost::lexical_cast(width) + " x " + + boost::lexical_cast(height) + " x " + + boost::lexical_cast(depth)); + } +#endif + + format_ = format; + width_ = width; + height_ = height; + depth_ = depth; + isLinearInterpolation_ = isLinearInterpolation; + + GLint interpolation = (isLinearInterpolation ? GL_LINEAR : GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MIN_FILTER, interpolation); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MAG_FILTER, interpolation); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + } + + + void OpenGLTextureArray::SetClampingToZero() + { +#if defined(__EMSCRIPTEN__) + /** + * This is because WebGL 2 derives from OpenGL ES 3.0, which + * doesn't support GL_CLAMP_TO_BORDER, as can be seen here: + * https://registry.khronos.org/OpenGL-Refpages/es3.0/html/glTexParameter.xhtml + **/ + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGLTextureArray::SetClampingToZero() is not available in WebGL 2"); +#else + ORTHANC_OPENGL_CHECK("Entering OpenGLTextureArray::SetClampingToZero()"); + + glBindTexture(GL_TEXTURE_2D_ARRAY, texture_); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); + glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); + + GLfloat colorfv[4] = { 0, 0, 0, 0 }; + glTexParameterfv(texture_, GL_TEXTURE_BORDER_COLOR, colorfv); + + ORTHANC_OPENGL_CHECK("Exiting OpenGLTextureArray::SetClampingToZero()"); +#endif + } + + + void OpenGLTextureArray::Bind(GLint location) const + { + glActiveTexture(GL_TEXTURE0); + glBindTexture(GL_TEXTURE_2D_ARRAY, texture_); + glUniform1i(location, 0 /* texture unit */); + } + + + void OpenGLTextureArray::BindAsTextureUnit(GLint location, + unsigned int unit) const + { + if (unit >= 32) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + + assert(GL_TEXTURE0 + 1 == GL_TEXTURE1 && + GL_TEXTURE0 + 31 == GL_TEXTURE31); + + glActiveTexture(GL_TEXTURE0 + unit); + glBindTexture(GL_TEXTURE_2D_ARRAY, texture_); + glUniform1i(location, unit /* texture unit */); + } + + + void OpenGLTextureArray::Upload(const Orthanc::ImageAccessor& image, + unsigned int layer) + { + if (image.GetWidth() != width_ || + image.GetHeight() != height_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + else if (layer >= depth_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else if (image.GetPitch() != Orthanc::GetBytesPerPixel(format_) * width_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, "Minimal pitch is required for upload"); + } + else if (width_ != 0 && + height_ != 0) + { + GLenum sourceFormat, internalFormat, pixelType; + OpenGLTexture::ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, image.GetFormat()); + + glBindTexture(GL_TEXTURE_2D_ARRAY, texture_); + +#if defined(__EMSCRIPTEN__) && (ORTHANC_WEBGL2_HEAP_COMPAT == 1) + // Check out "OpenGLTexture.cpp" for an explanation + EM_ASM({ + var ptr = emscriptenWebGLGetTexPixelData($5, $4, $2, $3, $0, $1); + GLctx.texSubImage3D(GLctx.TEXTURE_2D_ARRAY, 0, 0 /* x offset */, 0 /* y offset */, + $6, $2, $3, 1 /* depth */, $4, $5, ptr); + }, + image.GetConstBuffer(), // $0 + internalFormat, // $1 + image.GetWidth(), // $2 + image.GetHeight(), // $3 + sourceFormat, // $4 + pixelType, // $5 + layer); // $6 +#else + glTexSubImage3D(GL_TEXTURE_2D_ARRAY, 0, 0 /* x offset */, 0 /* y offset */, layer /* z offset */, + width_, height_, 1 /* depth */, sourceFormat, pixelType, image.GetConstBuffer()); +#endif + } + } + + + size_t OpenGLTextureArray::GetMemoryBufferSize() const + { + return static_cast(Orthanc::GetBytesPerPixel(format_)) * width_ * height_ * depth_; + } + + + void OpenGLTextureArray::Download(void* targetBuffer, + size_t targetSize) const + { + if (targetSize != GetMemoryBufferSize()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + else if (targetSize == 0) + { + return; + } + else if (targetBuffer == NULL) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_NullPointer); + } + else + { +#if 1 || defined(__EMSCRIPTEN__) + /** + * The "glGetTexImage()" function is unavailable in WebGL, it + * is necessary to use a framebuffer: + * https://stackoverflow.com/a/15064957 + **/ + OpenGLFramebuffer framebuffer(context_); + + const size_t sliceSize = targetSize / depth_; + + Orthanc::Image tmp(GetFormat(), GetWidth(), GetHeight(), true); + if (sliceSize != tmp.GetPitch() * height_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + + for (unsigned int layer = 0; layer < depth_; layer++) + { + framebuffer.ReadTexture(tmp, *this, layer); + memcpy(reinterpret_cast(targetBuffer) + layer * sliceSize, tmp.GetBuffer(), sliceSize); + } + +#else + glBindTexture(GL_TEXTURE_3D, texture_); + + switch (format_) + { + case Orthanc::PixelFormat_Grayscale8: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RED, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_RGB24: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RGB, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_RGBA32: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RGBA, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_Float32: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RED, GL_FLOAT, targetBuffer); + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + } +#endif + } + } + + + void OpenGLTextureArray::Download(std::string& target) const + { + target.resize(GetMemoryBufferSize()); + Download(target); + } + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTextureArray.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureArray.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,117 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include "OpenGLIncludes.h" +#include "IOpenGLContext.h" + +#include + +namespace OrthancStone +{ + namespace OpenGL + { + class OpenGLTextureArray : public boost::noncopyable + { + friend class OpenGLFramebuffer; + + private: + OpenGL::IOpenGLContext& context_; + GLuint texture_; + unsigned int width_; + unsigned int height_; + unsigned int depth_; + Orthanc::PixelFormat format_; + bool isLinearInterpolation_; + + /** + * Returns the low-level OpenGL handle of the texture + * array. Beware to never change the size of the texture using + * this handle! + **/ + GLuint GetId() const + { + return texture_; + } + + public: + OpenGLTextureArray(IOpenGLContext& context); + + ~OpenGLTextureArray(); + + unsigned int GetWidth() const + { + return width_; + } + + unsigned int GetHeight() const + { + return height_; + } + + unsigned int GetDepth() const + { + return depth_; + } + + Orthanc::PixelFormat GetFormat() const + { + return format_; + } + + bool IsLinearInterpolation() const + { + return isLinearInterpolation_; + } + + void Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + unsigned int depth, + bool isLinearInterpolation); + + /** + * By default, textures are mirrored at the borders. This + * function will set out-of-image access to zero. + **/ + void SetClampingToZero(); + + void Bind(GLint location) const; + + void BindAsTextureUnit(GLint location, + unsigned int unit) const; + + void Upload(const Orthanc::ImageAccessor& image, + unsigned int layer); + + size_t GetMemoryBufferSize() const; + + // "targetSize" must be equal to "GetMemoryBufferSize()" + void Download(void* targetBuffer, + size_t targetSize) const; + + void Download(std::string& target) const; + }; + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTextureVolume.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureVolume.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,304 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "OpenGLTextureVolume.h" + +#include "OpenGLFramebuffer.h" +#include "OpenGLTexture.h" + +#include +#include + +#include +#include + + +namespace OrthancStone +{ + namespace OpenGL + { + OpenGLTextureVolume::OpenGLTextureVolume(IOpenGLContext& context) : + context_(context), + texture_(0), + width_(0), + height_(0), + depth_(0), + format_(Orthanc::PixelFormat_Float32), + isLinearInterpolation_(false) + { + if (context.IsContextLost()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGL context has been lost"); + } + + glGenTextures(1, &texture_); + ORTHANC_OPENGL_CHECK("glGenTextures()"); + + if (texture_ == 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Cannot create an OpenGL texture array"); + } + } + + + OpenGLTextureVolume::~OpenGLTextureVolume() + { + assert(texture_ != 0); + glDeleteTextures(1, &texture_); + } + + + void OpenGLTextureVolume::Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + unsigned int depth, + bool isLinearInterpolation) + { + glActiveTexture(GL_TEXTURE0); + glBindTexture(GL_TEXTURE_3D, texture_); + ORTHANC_OPENGL_CHECK("glBindTexture(GL_TEXTURE_3D)"); + + GLenum sourceFormat, internalFormat, pixelType; + OpenGLTexture::ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, format); + + glTexImage3D(GL_TEXTURE_3D, 0, internalFormat, width, height, depth, + 0, sourceFormat, pixelType, NULL); + ORTHANC_OPENGL_CHECK("glTexImage3D()"); + +#if !defined(__EMSCRIPTEN__) + /** + * glGetTexLevelParameteriv() was introduced in OpenGL ES 3.1, + * but WebGL 2 only supports OpenGL ES 3.0, so it is not + * available in WebAssembly: + * https://registry.khronos.org/OpenGL-Refpages/es3.1/html/glGetTexLevelParameter.xhtml + **/ + GLint w, h, d; + glGetTexLevelParameteriv(GL_TEXTURE_3D, 0, GL_TEXTURE_WIDTH, &w); + glGetTexLevelParameteriv(GL_TEXTURE_3D, 0, GL_TEXTURE_HEIGHT, &h); + glGetTexLevelParameteriv(GL_TEXTURE_3D, 0, GL_TEXTURE_DEPTH, &d); + if (width != static_cast(w) || + height != static_cast(h) || + depth != static_cast(d)) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "Your GPU cannot create a 3D texture of size " + + boost::lexical_cast(width) + " x " + + boost::lexical_cast(height) + " x " + + boost::lexical_cast(depth)); + } +#endif + + format_ = format; + width_ = width; + height_ = height; + depth_ = depth; + isLinearInterpolation_ = isLinearInterpolation; + + GLint interpolation = (isLinearInterpolation ? GL_LINEAR : GL_NEAREST); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_MIN_FILTER, interpolation); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_MAG_FILTER, interpolation); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_WRAP_R, GL_CLAMP_TO_EDGE); + } + + + void OpenGLTextureVolume::SetClampingToZero() + { +#if defined(__EMSCRIPTEN__) + /** + * This is because WebGL 2 derives from OpenGL ES 3.0, which + * doesn't support GL_CLAMP_TO_BORDER, as can be seen here: + * https://registry.khronos.org/OpenGL-Refpages/es3.0/html/glTexParameter.xhtml + **/ + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, + "OpenGLTextureArray::SetClampingToZero() is not available in WebGL 2"); +#else + ORTHANC_OPENGL_CHECK("Entering OpenGLTextureArray::SetClampingToZero()"); + + glBindTexture(GL_TEXTURE_3D, texture_); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); + glTexParameteri(GL_TEXTURE_3D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); + + GLfloat colorfv[4] = { 0, 0, 0, 0 }; + glTexParameterfv(texture_, GL_TEXTURE_BORDER_COLOR, colorfv); + + ORTHANC_OPENGL_CHECK("Exiting OpenGLTextureArray::SetClampingToZero()"); +#endif + } + + + void OpenGLTextureVolume::Bind(GLint location) const + { + glActiveTexture(GL_TEXTURE0); + glBindTexture(GL_TEXTURE_3D, texture_); + glUniform1i(location, 0 /* texture unit */); + } + + + void OpenGLTextureVolume::BindAsTextureUnit(GLint location, + unsigned int unit) const + { + if (unit >= 32) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + + assert(GL_TEXTURE0 + 1 == GL_TEXTURE1 && + GL_TEXTURE0 + 31 == GL_TEXTURE31); + + glActiveTexture(GL_TEXTURE0 + unit); + glBindTexture(GL_TEXTURE_3D, texture_); + glUniform1i(location, unit /* texture unit */); + } + + + void OpenGLTextureVolume::Upload(const Orthanc::ImageAccessor& image, + unsigned int z) + { + if (image.GetWidth() != width_ || + image.GetHeight() != height_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + else if (z >= depth_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + else if (image.GetPitch() != Orthanc::GetBytesPerPixel(format_) * width_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError, "Minimal pitch is required for upload"); + } + else if (width_ != 0 && + height_ != 0) + { + GLenum sourceFormat, internalFormat, pixelType; + OpenGLTexture::ConvertToOpenGLFormats(sourceFormat, internalFormat, pixelType, image.GetFormat()); + + glBindTexture(GL_TEXTURE_3D, texture_); + +#if defined(__EMSCRIPTEN__) && (ORTHANC_WEBGL2_HEAP_COMPAT == 1) + // Check out "OpenGLTexture.cpp" for an explanation + EM_ASM({ + var ptr = emscriptenWebGLGetTexPixelData($5, $4, $2, $3, $0, $1); + GLctx.texSubImage3D(GLctx.TEXTURE_3D, 0, 0 /* x offset */, 0 /* y offset */, + $6, $2, $3, 1 /* depth */, $4, $5, ptr); + }, + image.GetConstBuffer(), // $0 + internalFormat, // $1 + image.GetWidth(), // $2 + image.GetHeight(), // $3 + sourceFormat, // $4 + pixelType, // $5 + z); // $6 +#else + glTexSubImage3D(GL_TEXTURE_3D, 0, 0 /* x offset */, 0 /* y offset */, z /* z offset */, + width_, height_, 1 /* depth */, sourceFormat, pixelType, image.GetConstBuffer()); +#endif + } + } + + + size_t OpenGLTextureVolume::GetMemoryBufferSize() const + { + return static_cast(Orthanc::GetBytesPerPixel(format_)) * width_ * height_ * depth_; + } + + + void OpenGLTextureVolume::Download(void* targetBuffer, + size_t targetSize) const + { + if (targetSize != GetMemoryBufferSize()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + else if (targetSize == 0) + { + return; + } + else if (targetBuffer == NULL) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_NullPointer); + } + else + { +#if 1 || defined(__EMSCRIPTEN__) + /** + * The "glGetTexImage()" function is unavailable in WebGL, it + * is necessary to use a framebuffer: + * https://stackoverflow.com/a/15064957 + **/ + OpenGLFramebuffer framebuffer(context_); + + const size_t sliceSize = targetSize / depth_; + + Orthanc::Image tmp(GetFormat(), GetWidth(), GetHeight(), true); + if (sliceSize != tmp.GetPitch() * height_) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError); + } + + for (unsigned int z = 0; z < depth_; z++) + { + framebuffer.ReadTexture(tmp, *this, z); + memcpy(reinterpret_cast(targetBuffer) + z * sliceSize, tmp.GetBuffer(), sliceSize); + } + +#else + glBindTexture(GL_TEXTURE_3D, texture_); + + switch (format_) + { + case Orthanc::PixelFormat_Grayscale8: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RED, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_RGB24: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RGB, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_RGBA32: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RGBA, GL_UNSIGNED_BYTE, targetBuffer); + break; + + case Orthanc::PixelFormat_Float32: + glGetTexImage(GL_TEXTURE_3D, 0 /* base level */, GL_RED, GL_FLOAT, targetBuffer); + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotImplemented); + } +#endif + } + } + + + void OpenGLTextureVolume::Download(std::string& target) const + { + target.resize(GetMemoryBufferSize()); + Download(target); + } + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/OpenGL/OpenGLTextureVolume.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureVolume.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,118 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include "OpenGLIncludes.h" +#include "IOpenGLContext.h" + +#include + + +namespace OrthancStone +{ + namespace OpenGL + { + class OpenGLTextureVolume : public boost::noncopyable + { + friend class OpenGLFramebuffer; + + private: + OpenGL::IOpenGLContext& context_; + GLuint texture_; + unsigned int width_; + unsigned int height_; + unsigned int depth_; + Orthanc::PixelFormat format_; + bool isLinearInterpolation_; + + /** + * Returns the low-level OpenGL handle of the texture + * array. Beware to never change the size of the texture using + * this handle! + **/ + GLuint GetId() const + { + return texture_; + } + + public: + OpenGLTextureVolume(IOpenGLContext& context); + + ~OpenGLTextureVolume(); + + unsigned int GetWidth() const + { + return width_; + } + + unsigned int GetHeight() const + { + return height_; + } + + unsigned int GetDepth() const + { + return depth_; + } + + Orthanc::PixelFormat GetFormat() const + { + return format_; + } + + bool IsLinearInterpolation() const + { + return isLinearInterpolation_; + } + + void Setup(Orthanc::PixelFormat format, + unsigned int width, + unsigned int height, + unsigned int depth, + bool isLinearInterpolation); + + /** + * By default, textures are mirrored at the borders. This + * function will set out-of-image access to zero. + **/ + void SetClampingToZero(); + + void Bind(GLint location) const; + + void BindAsTextureUnit(GLint location, + unsigned int unit) const; + + void Upload(const Orthanc::ImageAccessor& image, + unsigned int z); + + size_t GetMemoryBufferSize() const; + + // "targetSize" must be equal to "GetMemoryBufferSize()" + void Download(void* targetBuffer, + size_t targetSize) const; + + void Download(std::string& target) const; + }; + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.cpp --- a/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.cpp Thu May 30 21:31:37 2024 +0200 @@ -45,7 +45,8 @@ bool isContextLost_; public: - explicit PImpl(const std::string& canvasSelector) : + explicit PImpl(const std::string& canvasSelector, + Version version) : canvasSelector_(canvasSelector), isContextLost_(false) { @@ -53,6 +54,20 @@ EmscriptenWebGLContextAttributes attr; emscripten_webgl_init_context_attributes(&attr); + switch (version) + { + case Version_WebGL1: + break; + + case Version_WebGL2: + attr.majorVersion = 2; + attr.minorVersion = 0; + break; + + default: + throw Orthanc::OrthancException(Orthanc::ErrorCode_ParameterOutOfRange); + } + // The next line might be necessary to print using // WebGL. Sometimes, if set to "false" (the default value), // the canvas was rendered as a fully white or black @@ -162,7 +177,13 @@ WebAssemblyOpenGLContext::WebAssemblyOpenGLContext(const std::string& canvasSelector) : - pimpl_(new PImpl(canvasSelector)) + pimpl_(new PImpl(canvasSelector, Version_WebGL1)) + { + } + + WebAssemblyOpenGLContext::WebAssemblyOpenGLContext(const std::string& canvasSelector, + Version version) : + pimpl_(new PImpl(canvasSelector, version)) { } diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.h --- a/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.h Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.h Thu May 30 21:31:37 2024 +0200 @@ -51,6 +51,13 @@ { class WebAssemblyOpenGLContext : public OpenGL::IOpenGLContext { + public: + enum Version + { + Version_WebGL1, + Version_WebGL2 + }; + private: class PImpl; boost::shared_ptr pimpl_; @@ -58,6 +65,9 @@ public: explicit WebAssemblyOpenGLContext(const std::string& canvasSelector); + explicit WebAssemblyOpenGLContext(const std::string& canvasSelector, + Version version); + virtual bool IsContextLost() ORTHANC_OVERRIDE; virtual void MakeCurrent() ORTHANC_OVERRIDE; diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Scene2D/AnnotationsSceneLayer.h --- a/OrthancStone/Sources/Scene2D/AnnotationsSceneLayer.h Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Scene2D/AnnotationsSceneLayer.h Thu May 30 21:31:37 2024 +0200 @@ -21,6 +21,8 @@ **/ +#pragma once + #include "../Messages/IObservable.h" #include "Scene2D.h" #include "../Scene2DViewport/IFlexiblePointerTracker.h" diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Scene2D/Internals/OpenGLLinesProgram.cpp --- a/OrthancStone/Sources/Scene2D/Internals/OpenGLLinesProgram.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLLinesProgram.cpp Thu May 30 21:31:37 2024 +0200 @@ -426,7 +426,7 @@ if (!context_.IsContextLost() && !data.IsEmpty()) { context_.MakeCurrent(); - program_->Use(); + program_->Use(true); GLint locationPosition = program_->GetAttributeLocation("a_position"); GLint locationMiterDirection = program_->GetAttributeLocation("a_miter_direction"); diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Scene2D/Internals/OpenGLTextProgram.cpp --- a/OrthancStone/Sources/Scene2D/Internals/OpenGLTextProgram.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLTextProgram.cpp Thu May 30 21:31:37 2024 +0200 @@ -191,7 +191,7 @@ if (!context_.IsContextLost() && !data.IsEmpty()) { context_.MakeCurrent(); - program_->Use(); + program_->Use(true); double dx, dy; // In pixels ComputeAnchorTranslation(dx, dy, data.GetAnchor(), diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Scene2D/Internals/OpenGLTextureProgram.cpp --- a/OrthancStone/Sources/Scene2D/Internals/OpenGLTextureProgram.cpp Thu May 30 21:31:07 2024 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLTextureProgram.cpp Thu May 30 21:31:37 2024 +0200 @@ -54,7 +54,7 @@ if (!context_.IsContextLost()) { context_.MakeCurrent(); - program_->Use(); + program_->Use(true); AffineTransform2D scale = AffineTransform2D::CreateScaling (texture.GetWidth(), texture.GetHeight()); diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Toolbox/AlignedMatrix.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/AlignedMatrix.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,276 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "AlignedMatrix.h" + +#include + +#include + +namespace OrthancStone +{ + static unsigned int Ceiling(unsigned int a, + unsigned int b) + { + if (a % b == 0) + { + return a / b; + } + else + { + return a / b + 1; + } + } + + + void AlignedMatrix::Setup(unsigned int rows, + unsigned int cols) + { + assert(sizeof(float) == 4); + + if (rows == 0 || + cols == 0) + { + rows_ = 0; + cols_ = 0; + pitch_ = 0; + pitchFloatPointer_ = 0; + content_ = NULL; + } + else + { + rows_ = rows; + cols_ = cols; + pitch_ = Ceiling(cols * sizeof(float), ORTHANC_MEMORY_ALIGNMENT) * ORTHANC_MEMORY_ALIGNMENT; + pitchFloatPointer_ = pitch_ / sizeof(float); + + void* tmp = NULL; + if (posix_memalign(&tmp, ORTHANC_MEMORY_ALIGNMENT, rows_ * pitch_) != 0) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_NotEnoughMemory); + } + + assert(reinterpret_cast(tmp) % ORTHANC_MEMORY_ALIGNMENT == 0); + assert(pitch_ % ORTHANC_MEMORY_ALIGNMENT == 0); + assert(pitch_ % sizeof(float) == 0); + assert((rows_ * pitch_) % ORTHANC_MEMORY_ALIGNMENT == 0); + + content_ = static_cast(tmp); + } + } + + + AlignedMatrix::~AlignedMatrix() + { + if (content_ != NULL) + { + free(content_); + } + } + + + void AlignedMatrix::FillZeros() + { + memset(content_, 0, rows_ * pitch_); + } + + + void AlignedMatrix::ProductPlain(AlignedMatrix& c, + const AlignedMatrix& a, + const AlignedMatrix& b) + { + if (c.GetRows() != a.GetRows() || + c.GetColumns() != b.GetColumns() || + a.GetColumns() != b.GetRows()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + + const unsigned int M = c.GetRows(); + const unsigned int N = c.GetColumns(); + const unsigned int K = a.GetColumns(); + + c.FillZeros(); + + for (unsigned int i = 0; i < M; i++) + { + // Loop over "k" to be more cache-friendly + // https://sahnimanas.github.io/post/anatomy-of-a-high-performance-convolution/ + for (unsigned int k = 0; k < K; k++) + { + for (unsigned int j = 0; j < N; j++) + { + c.AddValue(i, j, a.GetValue(i, k) * b.GetValue(k, j)); + } + } + } + } + + +#if ORTHANC_HAS_MATRIX_PRODUCT_TRANSPOSED_VECTORIZED == 1 + // Computes "C = A*B^T" + class AlignedMatrix::ProductTransposedVectorizedContext : public boost::noncopyable + { + private: + unsigned int vectorizedSteps_; + uint8_t finalSteps_; + + public: + ORTHANC_FORCE_INLINE + ProductTransposedVectorizedContext(const AlignedMatrix& a) + { +#if ORTHANC_HAS_AVX2 == 1 + const unsigned int blockSize = 8; +#elif ORTHANC_HAS_SSE2 == 1 || ORTHANC_HAS_WASM_SIMD == 1 + const unsigned int blockSize = 4; +#else +# error No supported SIMD instruction set +#endif + + vectorizedSteps_ = a.GetColumns() / blockSize; + finalSteps_ = a.GetColumns() - vectorizedSteps_ * blockSize; + } + + ORTHANC_FORCE_INLINE + float Apply(const float* ap, + const float* btp) const noexcept + { + float result; + +#if ORTHANC_HAS_AVX2 == 1 + __m256 accumulator = _mm256_set1_ps(0); + + for (unsigned int k = 0; k < vectorizedSteps_; k++) + { + __m256 a = _mm256_load_ps(ap); + __m256 b = _mm256_load_ps(btp); + //accumulator = _mm256_add_ps(accumulator, _mm256_mul_ps(a, b)); + accumulator = _mm256_fmadd_ps(a, b, accumulator); // Requires the "-mfma" compiler flag + + ap += 8; + btp += 8; + } + + float tmp[8] __attribute__ ((aligned (ORTHANC_MEMORY_ALIGNMENT))); + _mm256_store_ps(tmp, accumulator); + result = tmp[0] + tmp[1] + tmp[2] + tmp[3] + tmp[4] + tmp[5] + tmp[6] + tmp[7]; + +#elif ORTHANC_HAS_SSE2 == 1 + __m128 accumulator = _mm_set1_ps(0); + + for (unsigned int k = 0; k < vectorizedSteps_; k++) + { + __m128 a = _mm_load_ps(ap); + __m128 b = _mm_load_ps(btp); + accumulator = _mm_add_ps(accumulator, _mm_mul_ps(a, b)); + ap += 4; + btp += 4; + } + +#if 1 + float tmp[4] __attribute__ ((aligned (ORTHANC_MEMORY_ALIGNMENT))); + _mm_storeu_ps(tmp, accumulator); + result = tmp[0] + tmp[1] + tmp[2] + tmp[3]; +#else + // This trickier version is theoretically faster, but no much difference in practice + const __m128 sum2 = _mm_add_ps(accumulator, _mm_shuffle_ps(accumulator, accumulator, _MM_SHUFFLE(2, 3, 0, 1))); + const __m128 sum1 = _mm_add_ps(sum2, _mm_shuffle_ps(sum2, sum2, _MM_SHUFFLE(0, 1, 2, 3))); + result = _mm_cvtss_f32(sum1); +#endif + +#elif ORTHANC_HAS_WASM_SIMD == 1 + v128_t accumulator = wasm_f32x4_splat(0); + + for (unsigned int k = 0; k < vectorizedSteps_; k++) + { + v128_t a = wasm_v128_load(ap); + v128_t b = wasm_v128_load(btp); + accumulator = wasm_f32x4_add(accumulator, wasm_f32x4_mul(a, b)); + ap += 4; + btp += 4; + } + +#if 1 + float tmp[4]; + wasm_v128_store(tmp, accumulator); + result = tmp[0] + tmp[1] + tmp[2] + tmp[3]; +#else + const v128_t sum2 = wasm_f32x4_add(accumulator, wasm_i32x4_shuffle(accumulator, accumulator, 2, 3, 0, 0)); + const v128_t sum1 = wasm_f32x4_add(sum2, wasm_i32x4_shuffle(sum2, sum2, 1, 0, 0, 0)); + result = wasm_f32x4_extract_lane(sum1, 0); +#endif + +#else +# error No supported SIMD instruction set +#endif + + for (uint8_t k = 0; k < finalSteps_; k++) + { + result += (*ap) * (*btp); + ap++; + btp++; + } + + return result; + } + }; +#endif + + +#if ORTHANC_HAS_MATRIX_PRODUCT_TRANSPOSED_VECTORIZED == 1 + void AlignedMatrix::ProductTransposedVectorized(AlignedMatrix& c, + const AlignedMatrix& a, + const AlignedMatrix& bt) + { + if (c.GetRows() != a.GetRows() || + c.GetColumns() != bt.GetRows() || + a.GetColumns() != bt.GetColumns()) + { + throw Orthanc::OrthancException(Orthanc::ErrorCode_IncompatibleImageSize); + } + + AlignedMatrix::ProductTransposedVectorizedContext context(a); + + const unsigned int M = a.GetRows(); + const unsigned int N = bt.GetRows(); + + const size_t rowSizeA = a.GetPitch() / sizeof(float); + const size_t rowSizeB = bt.GetPitch() / sizeof(float); + + const float* ap = a.GetRowPointer(0); + for (unsigned int i = 0; i < M; i++) + { + float* cp = c.GetRowPointer(i); + + const float* btp = bt.GetRowPointer(0); + for (unsigned int j = 0; j < N; j++, cp++) + { + *cp = context.Apply(ap, btp); + btp += rowSizeB; + } + + ap += rowSizeA; + } + } +#endif +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Toolbox/AlignedMatrix.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/AlignedMatrix.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,135 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include "SimdIncludes.h" + +#if (ORTHANC_HAS_AVX2 == 1 || ORTHANC_HAS_SSE2 == 1 || ORTHANC_HAS_WASM_SIMD == 1) +# define ORTHANC_HAS_MATRIX_PRODUCT_TRANSPOSED_VECTORIZED 1 +#else +# define ORTHANC_HAS_MATRIX_PRODUCT_TRANSPOSED_VECTORIZED 0 +#endif + + +#include +#include + +namespace OrthancStone +{ + /** + * 2D matrix whose rows are aligned for the largest SIMD + * instructions that are available. + **/ + class AlignedMatrix : public boost::noncopyable + { + private: + class ProductTransposedVectorizedContext; + + unsigned int rows_; + unsigned int cols_; + size_t pitch_; + size_t pitchFloatPointer_; + float* content_; + + void Setup(unsigned int rows, + unsigned int cols); + + public: + AlignedMatrix(unsigned int rows, + unsigned int cols) + { + Setup(rows, cols); + } + + ~AlignedMatrix(); + + unsigned int GetRows() const + { + return rows_; + } + + unsigned int GetColumns() const + { + return cols_; + } + + unsigned int GetPitch() const + { + return pitch_; + } + + float* GetRowPointer(unsigned int row) + { + assert(row < rows_); + return content_ + row * pitchFloatPointer_; + } + + const float* GetRowPointer(unsigned int row) const + { + assert(row < rows_); + return content_ + row * pitchFloatPointer_; + } + + size_t GetIndex(unsigned int row, + unsigned int col) const + { + assert(row < rows_ && col < cols_); + return row * pitchFloatPointer_ + col; + } + + float GetValue(unsigned int row, + unsigned int col) const + { + return content_[GetIndex(row, col)]; + } + + void SetValue(unsigned int row, + unsigned int col, + float value) const + { + content_[GetIndex(row, col)] = value; + } + + void AddValue(unsigned int row, + unsigned int col, + float value) + { + content_[GetIndex(row, col)] += value; + } + + void FillZeros(); + + // Computes "C = A * B" without SIMD operations + static void ProductPlain(AlignedMatrix& c, + const AlignedMatrix& a, + const AlignedMatrix& b); + +#if ORTHANC_HAS_MATRIX_PRODUCT_TRANSPOSED_VECTORIZED == 1 + // Computes "C = A * B^T" using SIMD operations + static void ProductTransposedVectorized(AlignedMatrix& c, + const AlignedMatrix& a, + const AlignedMatrix& bt); +#endif + }; +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Toolbox/SimdIncludes.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/SimdIncludes.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,66 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#if defined(__EMSCRIPTEN__) +# include +# include +#else +# include // portable to all x86 compilers +#endif + +#if __AVX2__ == 1 +# define ORTHANC_HAS_AVX2 1 +# define ORTHANC_HAS_SSE2 1 +# define ORTHANC_HAS_WASM_SIMD 0 +# define ORTHANC_MEMORY_ALIGNMENT 32 +#elif __SSE2__ == 1 +# define ORTHANC_HAS_AVX2 0 +# define ORTHANC_HAS_SSE2 1 +# define ORTHANC_HAS_WASM_SIMD 0 +# define ORTHANC_MEMORY_ALIGNMENT 16 +#elif defined(__EMSCRIPTEN__) +# if !defined(ORTHANC_HAS_WASM_SIMD) +# error ORTHANC_HAS_WASM_SIMD must be defined to use this file +# endif +# define ORTHANC_HAS_AVX2 0 +# define ORTHANC_HAS_SSE2 0 +# if ORTHANC_HAS_WASM_SIMD == 1 +// Setting macro "ORTHANC_HAS_WASM_SIMD" to "1" means that +// "-msimd128" has been provided to Emscripten (there doesn't seem +// to exist a predefined macro to automatically check this) +# define ORTHANC_MEMORY_ALIGNMENT 16 +# else +# define ORTHANC_MEMORY_ALIGNMENT 8 +# endif +#elif defined(_MSC_VER) +# if _M_IX86_FP >= 2 // https://stackoverflow.com/a/18563988 +# define ORTHANC_HAS_AVX2 0 +# define ORTHANC_HAS_SSE2 0 +# define ORTHANC_HAS_WASM_SIMD 1 +# define ORTHANC_MEMORY_ALIGNMENT 16 +# endif +#else +# define ORTHANC_MEMORY_ALIGNMENT 8 +#endif diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Toolbox/TimerLogger.cpp --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/TimerLogger.cpp Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,53 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#include "TimerLogger.h" + +#include + + +namespace OrthancStone +{ + TimerLogger::TimerLogger(const std::string& name) : + name_(name) + { +#if defined(__EMSCRIPTEN__) + start_ = emscripten_get_now(); +#else + start_ = boost::posix_time::microsec_clock::universal_time(); +#endif + } + + + TimerLogger::~TimerLogger() + { +#if defined(__EMSCRIPTEN__) + int elapsed = static_cast(round(emscripten_get_now() - start_)); +#else + const boost::posix_time::ptime end = boost::posix_time::microsec_clock::universal_time(); + int elapsed = (end - start_).total_milliseconds(); +#endif + + LOG(WARNING) << name_ << " - Elapsed time: " << elapsed << "ms"; + } +} diff -r 2965172977a4 -r 93258466311b OrthancStone/Sources/Toolbox/TimerLogger.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/TimerLogger.h Thu May 30 21:31:37 2024 +0200 @@ -0,0 +1,54 @@ +/** + * Stone of Orthanc + * Copyright (C) 2012-2016 Sebastien Jodogne, Medical Physics + * Department, University Hospital of Liege, Belgium + * Copyright (C) 2017-2022 Osimis S.A., Belgium + * Copyright (C) 2021-2022 Sebastien Jodogne, ICTEAM UCLouvain, Belgium + * + * This program is free software: you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public License + * as published by the Free Software Foundation, either version 3 of + * the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, but + * WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with this program. If not, see + * . + **/ + + +#pragma once + +#include +#include + +#if defined(__EMSCRIPTEN__) +# include +#else +# include +#endif + + +namespace OrthancStone +{ + class TimerLogger : public boost::noncopyable + { + private: + std::string name_; + +#if defined(__EMSCRIPTEN__) + double start_; +#else + boost::posix_time::ptime start_; +#endif + + public: + TimerLogger(const std::string& name); + + ~TimerLogger(); + }; +} diff -r 2965172977a4 -r 93258466311b TODO --- a/TODO Thu May 30 21:31:07 2024 +0200 +++ b/TODO Thu May 30 21:31:37 2024 +0200 @@ -107,6 +107,9 @@ SOP Class UIDs for per-modality setting): https://groups.google.com/g/orthanc-users/c/tHPMBHe-Z-A/m/XXlRBdhoAgAJ +* Sort the studies by date in the study picker: + https://discourse.orthanc-server.org/t/the-study-pickers-list-does-not-sort-properly/3489 + -------- Won't do --------