Mercurial > hg > orthanc-stone
changeset 2204:e1613509a939 deep-learning
integration default->deep-learning, turning deep learning into plugin
| author | Sebastien Jodogne <s.jodogne@gmail.com> |
|---|---|
| date | Sat, 19 Apr 2025 14:46:27 +0200 (2 weeks ago) |
| parents | 2795f1ee4a1a (diff) dcfabb36dc21 (current diff) |
| children | a8066ce6bacc |
| files | Applications/StoneWebViewer/WebAssembly/CMakeLists.txt Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp |
| diffstat | 35 files changed, 2630 insertions(+), 84 deletions(-) [+] |
line wrap: on
line diff
--- a/Applications/Samples/WebAssembly/CMakeLists.txt Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/Samples/WebAssembly/CMakeLists.txt Sat Apr 19 14:46:27 2025 +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 )
--- a/Applications/StoneWebViewer/WebApplication/app.js Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/StoneWebViewer/WebApplication/app.js Sat Apr 19 14:46:27 2025 +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: { @@ -1377,6 +1381,11 @@ }); }, + ApplyDeepLearning: function() { + stone.ApplyDeepLearningModel(this.GetActiveCanvas()); + app.deepLearningStartTime = performance.now(); + }, + ChangeActiveSeries: function(offset) { var seriesTags = this.GetActiveViewportSeriesTags(); if (seriesTags !== null) { @@ -1832,3 +1841,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'); +});
--- a/Applications/StoneWebViewer/WebApplication/index.html Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/StoneWebViewer/WebApplication/index.html Sat Apr 19 14:46:27 2025 +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"
--- a/Applications/StoneWebViewer/WebAssembly/CMakeLists.txt Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/StoneWebViewer/WebAssembly/CMakeLists.txt Sat Apr 19 14:46:27 2025 +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\", \"lengthBytesUTF8\", \"stringToUTF8\"]'") 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 @@ -108,7 +114,7 @@ add_custom_command( COMMAND - ${CLANG_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/ParseWebAssemblyExports.py --libclang=${LIBCLANG} ${CMAKE_SOURCE_DIR}/StoneWebViewer.cpp > ${STONE_WRAPPER} + ${CLANG_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/ParseWebAssemblyExports.py --libclang=${LIBCLANG} ${CMAKE_SOURCE_DIR}/StoneWebViewer.cpp ${CMAKE_SOURCE_DIR}/DeepLearning.cpp > ${STONE_WRAPPER} DEPENDS ${CMAKE_SOURCE_DIR}/StoneWebViewer.cpp ${CMAKE_SOURCE_DIR}/ParseWebAssemblyExports.py @@ -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,9 @@ add_executable(StoneWebViewer ${ORTHANC_STONE_SOURCES} - ${AUTOGENERATED_SOURCES} + ${AUTOGENERATED_SOURCES} # Populated by "EmbedResources()" + ${AUTOGENERATED_DIR}/DeepLearningWorker.pb.cc + DeepLearning.cpp StoneWebViewer.cpp ) @@ -149,7 +169,7 @@ ) # Make sure to have the wrapper generated -add_dependencies(StoneWebViewer StoneWrapper) +add_dependencies(StoneWebViewer AutogeneratedFiles) # Declare installation files for the module
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/Applications/StoneWebViewer/WebAssembly/DeepLearningWorker.proto Sat Apr 19 14:46:27 2025 +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; +}
--- a/Applications/StoneWebViewer/WebAssembly/StoneModule/CMakeLists.txt Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/StoneWebViewer/WebAssembly/StoneModule/CMakeLists.txt Sat Apr 19 14:46:27 2025 +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 )
--- a/Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/Applications/StoneWebViewer/WebAssembly/StoneWebViewer.cpp Sat Apr 19 14:46:27 2025 +0200 @@ -2262,10 +2262,8 @@ // the center of the top-left pixel boost::shared_ptr<OrthancStone::AnnotationsSceneLayer> stoneAnnotations_; - bool linearInterpolation_; - // WARNING: The ownership is not transferred std::list<ILayerSource*> layerSources_; @@ -3890,7 +3888,6 @@ // Orientation markers, new in Stone Web viewer 2.4 static std::unique_ptr<OrientationMarkersSource> orientationMarkersSource_; - static void UpdateReferenceLines() { if (showReferenceLines_) @@ -4209,9 +4206,12 @@ } +IStoneWebViewerPlugin* DeepLearningInitialization(IStoneWebViewerContext& context); + typedef IStoneWebViewerPlugin* (*PluginInitializer) (IStoneWebViewerContext&); static const PluginInitializer pluginsInitializers_[] = { + DeepLearningInitialization, NULL };
--- a/OrthancStone/Resources/CMake/OrthancStoneConfiguration.cmake Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Resources/CMake/OrthancStoneConfiguration.cmake Sat Apr 19 14:46:27 2025 +0200 @@ -320,6 +320,7 @@ ${ORTHANC_STONE_ROOT}/StoneInitialization.cpp ${ORTHANC_STONE_ROOT}/Toolbox/AffineTransform2D.cpp + ${ORTHANC_STONE_ROOT}/Toolbox/AlignedMatrix.cpp ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator1D.cpp ${ORTHANC_STONE_ROOT}/Toolbox/BucketAccumulator2D.cpp ${ORTHANC_STONE_ROOT}/Toolbox/CoordinateSystem3D.cpp @@ -343,6 +344,7 @@ ${ORTHANC_STONE_ROOT}/Toolbox/SortedFrames.cpp ${ORTHANC_STONE_ROOT}/Toolbox/StoneToolbox.cpp ${ORTHANC_STONE_ROOT}/Toolbox/TextRenderer.cpp + ${ORTHANC_STONE_ROOT}/Toolbox/TimerLogger.cpp ${ORTHANC_STONE_ROOT}/Toolbox/UndoRedoStack.cpp ${ORTHANC_STONE_ROOT}/Toolbox/UnionOfRectangles.cpp ${ORTHANC_STONE_ROOT}/Toolbox/Windowing.cpp @@ -385,6 +387,10 @@ ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLProgram.cpp ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLShader.cpp ${ORTHANC_STONE_ROOT}/OpenGL/OpenGLTexture.cpp + ${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
--- a/OrthancStone/Resources/WebAssemblySharedLibrary/CMakeLists.txt Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Resources/WebAssemblySharedLibrary/CMakeLists.txt Sat Apr 19 14:46:27 2025 +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 )
--- a/OrthancStone/Resources/WebAssemblyUnitTests/CMakeLists.txt Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Resources/WebAssemblyUnitTests/CMakeLists.txt Sat Apr 19 14:46:27 2025 +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 )
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/ImageProcessingProgram.cpp Sat Apr 19 14:46:27 2025 +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()"); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/ImageProcessingProgram.h Sat Apr 19 14:46:27 2025 +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); + } + }; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLFramebuffer.cpp Sat Apr 19 14:46:27 2025 +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); + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLFramebuffer.h Sat Apr 19 14:46:27 2025 +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); + }; + } +}
--- a/OrthancStone/Sources/OpenGL/OpenGLProgram.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLProgram.cpp Sat Apr 19 14:46:27 2025 +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"); }
--- a/OrthancStone/Sources/OpenGL/OpenGLProgram.h Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLProgram.h Sat Apr 19 14:46:27 2025 +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,
--- a/OrthancStone/Sources/OpenGL/OpenGLShader.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLShader.cpp Sat Apr 19 14:46:27 2025 +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; } }
--- a/OrthancStone/Sources/OpenGL/OpenGLTexture.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLTexture.cpp Sat Apr 19 14:46:27 2025 +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))); + } + } } }
--- a/OrthancStone/Sources/OpenGL/OpenGLTexture.h Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/OpenGL/OpenGLTexture.h Sat Apr 19 14:46:27 2025 +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); }; } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureArray.cpp Sat Apr 19 14:46:27 2025 +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); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureArray.h Sat Apr 19 14:46:27 2025 +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; + }; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureVolume.cpp Sat Apr 19 14:46:27 2025 +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); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/OpenGL/OpenGLTextureVolume.h Sat Apr 19 14:46:27 2025 +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; + }; + } +}
--- a/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.cpp Sat Apr 19 14:46:27 2025 +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)) { }
--- a/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.h Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Platforms/WebAssembly/WebAssemblyOpenGLContext.h Sat Apr 19 14:46:27 2025 +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;
--- a/OrthancStone/Sources/Scene2D/AnnotationsSceneLayer.h Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Scene2D/AnnotationsSceneLayer.h Sat Apr 19 14:46:27 2025 +0200 @@ -21,6 +21,8 @@ **/ +#pragma once + #include "../Messages/IObservable.h" #include "Scene2D.h" #include "../Scene2DViewport/IFlexiblePointerTracker.h"
--- a/OrthancStone/Sources/Scene2D/Internals/OpenGLLinesProgram.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLLinesProgram.cpp Sat Apr 19 14:46:27 2025 +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");
--- a/OrthancStone/Sources/Scene2D/Internals/OpenGLTextProgram.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLTextProgram.cpp Sat Apr 19 14:46:27 2025 +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(),
--- a/OrthancStone/Sources/Scene2D/Internals/OpenGLTextureProgram.cpp Sat Apr 19 14:40:38 2025 +0200 +++ b/OrthancStone/Sources/Scene2D/Internals/OpenGLTextureProgram.cpp Sat Apr 19 14:46:27 2025 +0200 @@ -54,7 +54,7 @@ if (!context_.IsContextLost()) { context_.MakeCurrent(); - program_->Use(); + program_->Use(true); AffineTransform2D scale = AffineTransform2D::CreateScaling (texture.GetWidth(), texture.GetHeight());
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/AlignedMatrix.cpp Sat Apr 19 14:46:27 2025 +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 +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/AlignedMatrix.h Sat Apr 19 14:46:27 2025 +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 + }; +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/SimdIncludes.h Sat Apr 19 14:46:27 2025 +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
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/TimerLogger.cpp Sat Apr 19 14:46:27 2025 +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"; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/OrthancStone/Sources/Toolbox/TimerLogger.h Sat Apr 19 14:46:27 2025 +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(); + }; +}
--- a/TODO Sat Apr 19 14:40:38 2025 +0200 +++ b/TODO Sat Apr 19 14:46:27 2025 +0200 @@ -116,6 +116,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 --------