# HG changeset patch
# User Sebastien Jodogne <s.jodogne@gmail.com>
# 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 @@
 
             </div>        
             <div class="wvLayoutLeft__contentBottom">
+              <div v-if="deepLearningReady">
+                <div style="width:100%;padding:10px;text-align:center;">
+                  <button class="btn btn-primary" @click="ApplyDeepLearning()">Apply deep learning</button>
+                </div>
+                
+                <div style="padding: 10px; position: relative; width:100%;">
+                  <div style="background-color: #007000; position: relative; height: 10px;">    
+                    <div v-bind:style="{ 'background-color': '#00ff00', position: 'absolute', height: '100%', width: (deepLearningProgress*100) + '%' }"></div>
+                  </div>
+                </div>
+              </div>
+                
               <div style="width:100%;padding:10px;text-align:center;"
                    v-if="globalConfiguration.InstitutionLogo != ''">
                 <img style="max-width:100%" v-bind:src="globalConfiguration.InstitutionLogo" />
               </div>
-            </div>        
+            </div>
           </div>
         </div>
         <div class="wvLayout__main"
diff -r 2965172977a4 -r 93258466311b Applications/StoneWebViewer/WebAssembly/CMakeLists.txt
--- a/Applications/StoneWebViewer/WebAssembly/CMakeLists.txt	Thu May 30 21:31:07 2024 +0200
+++ b/Applications/StoneWebViewer/WebAssembly/CMakeLists.txt	Thu May 30 21:31:37 2024 +0200
@@ -34,14 +34,18 @@
 set(WASM_FLAGS "${WASM_FLAGS} -s WASM=1 -s FETCH=1 -s ASSERTIONS=1 -s DISABLE_EXCEPTION_CATCHING=0")
 if (CMAKE_BUILD_TYPE STREQUAL "Debug")
   set(WASM_FLAGS "${WASM_FLAGS} -s SAFE_HEAP=1")
+  set(WASM_LINKER_FLAGS "${WASM_LINKER_FLAGS} -s GL_DEBUG=1 -s GL_ASSERTIONS=1 -s TRACE_WEBGL_CALLS=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
 )
 
 
@@ -59,6 +63,8 @@
 SET(ENABLE_MODULE_JOBS OFF)
 SET(ENABLE_PUGIXML ON) # Necessary for OsiriX annotations
 SET(ORTHANC_SANDBOXED ON)
+SET(ENABLE_PROTOBUF ON)           # Necessary for deep learning
+SET(ENABLE_PROTOBUF_COMPILER ON)  # Necessary for deep learning
 
 # this will set up the build system for Stone of Orthanc and will
 # populate the ORTHANC_STONE_SOURCES CMake variable
@@ -116,9 +122,21 @@
   ${STONE_WRAPPER}
   )
 
-add_custom_target(StoneWrapper
+add_custom_command(
+  OUTPUT
+  ${AUTOGENERATED_DIR}/DeepLearningWorker.pb.h
+  ${AUTOGENERATED_DIR}/DeepLearningWorker.pb.cc
+  COMMAND ${PROTOC_EXECUTABLE} --cpp_out=${AUTOGENERATED_DIR} -I${CMAKE_SOURCE_DIR}
+  ${CMAKE_SOURCE_DIR}/DeepLearningWorker.proto
+  DEPENDS
+  ${CMAKE_SOURCE_DIR}/DeepLearningWorker.proto
+  ProtobufCompiler
+  )
+
+add_custom_target(AutogeneratedFiles
   DEPENDS
   ${STONE_WRAPPER}
+  ${AUTOGENERATED_DIR}/DeepLearningWorker.pb.cc
   )  
 
 
@@ -131,7 +149,8 @@
 
 add_executable(StoneWebViewer
   ${ORTHANC_STONE_SOURCES}
-  ${AUTOGENERATED_SOURCES}
+  ${AUTOGENERATED_SOURCES}  # Populated by "EmbedResources()"
+  ${AUTOGENERATED_DIR}/DeepLearningWorker.pb.cc
   StoneWebViewer.cpp
   )
 
@@ -149,7 +168,7 @@
   )
 
 # Make sure to have the wrapper generated
-add_dependencies(StoneWebViewer StoneWrapper)
+add_dependencies(StoneWebViewer AutogeneratedFiles)
 
 
 # Declare installation files for the module
diff -r 2965172977a4 -r 93258466311b Applications/StoneWebViewer/WebAssembly/DeepLearningWorker.proto
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/Applications/StoneWebViewer/WebAssembly/DeepLearningWorker.proto	Thu May 30 21:31:37 2024 +0200
@@ -0,0 +1,61 @@
+syntax = "proto2";
+option optimize_for = LITE_RUNTIME;
+
+package OrthancStone.Messages;
+
+enum RequestType {
+  PARSE_MODEL = 1;
+  LOAD_IMAGE = 2;
+  EXECUTE_STEP = 3;
+}
+
+enum ResponseType {
+  INITIALIZED = 1;
+  PARSED_MODEL = 2;
+  LOADED_IMAGE = 3;
+  STEP_DONE = 4;
+}
+
+message ParseModelRequest {
+  required bytes content = 1;
+}
+
+message ParseModelResponse {
+  required uint32 number_of_steps = 1;
+}
+
+message LoadImageRequest {
+  required string sop_instance_uid = 1;
+  required uint32 frame_number = 2;
+  required uint32 height = 3;
+  required uint32 width = 4;
+  repeated float values = 5 [packed=true];
+}
+
+message SegmentationMask {
+  required string sop_instance_uid = 1;
+  required uint32 frame_number = 2;
+  required uint32 height = 3;
+  required uint32 width = 4;
+  repeated bool values = 5 [packed=true];
+}
+
+message StepResponse {
+  required bool done = 1;
+  required float progress = 2;
+  optional SegmentationMask mask = 3;
+}
+
+message Request {
+  required RequestType type = 1;
+
+  optional ParseModelRequest parse_model = 2;
+  optional LoadImageRequest load_image = 3;
+}
+
+message Response {
+  required ResponseType type = 1;
+
+  optional ParseModelResponse parse_model = 2;
+  optional StepResponse step = 3;
+}
diff -r 2965172977a4 -r 93258466311b Applications/StoneWebViewer/WebAssembly/StoneModule/CMakeLists.txt
--- a/Applications/StoneWebViewer/WebAssembly/StoneModule/CMakeLists.txt	Thu May 30 21:31:07 2024 +0200
+++ b/Applications/StoneWebViewer/WebAssembly/StoneModule/CMakeLists.txt	Thu May 30 21:31:37 2024 +0200
@@ -46,6 +46,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 ASSERTIONS=1 -s DISABLE_EXCEPTION_CATCHING=0")
@@ -53,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 Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp
--- a/Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp	Thu May 30 21:31:07 2024 +0200
+++ b/Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp	Thu May 30 21:31:37 2024 +0200
@@ -1684,10 +1684,11 @@
 private:
   static const int LAYER_TEXTURE = 0;
   static const int LAYER_OVERLAY = 1;
-  static const int LAYER_ORIENTATION_MARKERS = 2;
-  static const int LAYER_REFERENCE_LINES = 3;
-  static const int LAYER_ANNOTATIONS_OSIRIX = 4;
-  static const int LAYER_ANNOTATIONS_STONE = 5;
+  static const int LAYER_DEEP_LEARNING = 2;
+  static const int LAYER_ORIENTATION_MARKERS = 3;
+  static const int LAYER_REFERENCE_LINES = 4;
+  static const int LAYER_ANNOTATIONS_OSIRIX = 5;
+  static const int LAYER_ANNOTATIONS_STONE = 6;
 
   
   class ICommand : public Orthanc::IDynamicObject
@@ -2109,9 +2110,12 @@
   // coordinates of the current texture, with (0,0) corresponding to
   // the center of the top-left pixel
   boost::shared_ptr<OrthancStone::AnnotationsSceneLayer>  stoneAnnotations_;
-
+  
   bool linearInterpolation_;
 
+  boost::shared_ptr<Orthanc::ImageAccessor>  deepLearningMask_;
+  std::string deepLearningSopInstanceUid_;
+  unsigned int deepLearningFrameNumber_;
 
   void ScheduleNextPrefetch()
   {
@@ -2330,6 +2334,26 @@
       }
     }
 
+    std::unique_ptr<OrthancStone::LookupTableTextureSceneLayer> deepLearningLayer;
+
+    if (deepLearningMask_.get() != NULL &&
+        deepLearningSopInstanceUid_ == instance.GetSopInstanceUid() &&
+        deepLearningFrameNumber_ == frameIndex)
+    {
+      std::vector<uint8_t> 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 <emscripten/fetch.h>
+#include <DeepLearningWorker.pb.h>
+
+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<const float*>(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<uint8_t*>(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<ViewerViewport> 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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#include "ImageProcessingProgram.h"
+
+#include "OpenGLFramebuffer.h"
+
+#include <OrthancException.h>
+
+
+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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#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 <OrthancException.h>
+
+
+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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#pragma once
+
+#include "OpenGLIncludes.h"
+#include "IOpenGLContext.h"
+
+#include <Images/ImageAccessor.h>
+
+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 <Images/Image.h>
 #include <Logging.h>
 #include <OrthancException.h>
 
@@ -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<unsigned int>(w) ||
+            height != static_cast<unsigned int>(h))
+        {
+          throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError,
+                                          "Your GPU cannot create a texture of size " +
+                                          boost::lexical_cast<std::string>(width) + " x " +
+                                          boost::lexical_cast<std::string>(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<Orthanc::ImageAccessor> 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 <Images/ImageAccessor.h>
 
@@ -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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#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 <Images/Image.h>
+#include <OrthancException.h>
+
+#include <boost/lexical_cast.hpp>
+#include <cassert>
+
+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<unsigned int>(w) ||
+          height != static_cast<unsigned int>(h) ||
+          depth != static_cast<unsigned int>(d))
+      {
+        throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError,
+                                        "Your GPU cannot create an array of textures of size " +
+                                        boost::lexical_cast<std::string>(width) + " x " +
+                                        boost::lexical_cast<std::string>(height) + " x " +
+                                        boost::lexical_cast<std::string>(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<size_t>(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<uint8_t*>(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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#pragma once
+
+#include "OpenGLIncludes.h"
+#include "IOpenGLContext.h"
+
+#include <Images/ImageAccessor.h>
+
+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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#include "OpenGLTextureVolume.h"
+
+#include "OpenGLFramebuffer.h"
+#include "OpenGLTexture.h"
+
+#include <Images/Image.h>
+#include <OrthancException.h>
+
+#include <boost/lexical_cast.hpp>
+#include <cassert>
+
+
+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<unsigned int>(w) ||
+          height != static_cast<unsigned int>(h) ||
+          depth != static_cast<unsigned int>(d))
+      {
+        throw Orthanc::OrthancException(Orthanc::ErrorCode_InternalError,
+                                        "Your GPU cannot create a 3D texture of size " +
+                                        boost::lexical_cast<std::string>(width) + " x " +
+                                        boost::lexical_cast<std::string>(height) + " x " +
+                                        boost::lexical_cast<std::string>(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<size_t>(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<uint8_t*>(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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#pragma once
+
+#include "OpenGLIncludes.h"
+#include "IOpenGLContext.h"
+
+#include <Images/ImageAccessor.h>
+
+
+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>  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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#include "AlignedMatrix.h"
+
+#include <OrthancException.h>
+
+#include <string.h>
+
+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<intptr_t>(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<float*>(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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#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 <boost/noncopyable.hpp>
+#include <cassert>
+
+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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#pragma once
+
+#if defined(__EMSCRIPTEN__)
+#  include <emscripten.h>
+#  include <wasm_simd128.h>
+#else
+#  include <immintrin.h>  // 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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#include "TimerLogger.h"
+
+#include <Logging.h>
+
+
+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<int>(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
+ * <http://www.gnu.org/licenses/>.
+ **/
+
+
+#pragma once
+
+#include <boost/noncopyable.hpp>
+#include <string>
+
+#if defined(__EMSCRIPTEN__)
+#  include <emscripten.h>
+#else
+#  include <boost/date_time/posix_time/posix_time_types.hpp>
+#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
 --------