Commit 9efbc6b11e5caff3e72ed01e98838f591ffabda2
0 parents
Initialise the repository
Showing
44 changed files
with
2354 additions
and
0 deletions
.gitignore
0 → 100644
1 | +++ a/.gitignore | ||
1 | +# Prerequisites | ||
2 | +*.d | ||
3 | + | ||
4 | +# Compiled Object files | ||
5 | +*.slo | ||
6 | +*.lo | ||
7 | +*.o | ||
8 | +*.obj | ||
9 | + | ||
10 | +# Precompiled Headers | ||
11 | +*.gch | ||
12 | +*.pch | ||
13 | + | ||
14 | +# Compiled Dynamic libraries | ||
15 | +*.so | ||
16 | +*.dylib | ||
17 | +*.dll | ||
18 | + | ||
19 | +# Fortran module files | ||
20 | +*.mod | ||
21 | +*.smod | ||
22 | + | ||
23 | +# Compiled Static libraries | ||
24 | +*.lai | ||
25 | +*.la | ||
26 | +*.a | ||
27 | +*.lib | ||
28 | + | ||
29 | +# Executables | ||
30 | +*.exe | ||
31 | +*.out | ||
32 | +*.app | ||
33 | + | ||
34 | +.metadata | ||
35 | +bin/ | ||
36 | +tmp/ | ||
37 | +*.tmp | ||
38 | +*.bak | ||
39 | +*.swp | ||
40 | +*~.nib | ||
41 | +local.properties | ||
42 | +.settings/ | ||
43 | +.loadpath | ||
44 | +.recommenders | ||
45 | + | ||
46 | +# External tool builders | ||
47 | +.externalToolBuilders/ | ||
48 | + | ||
49 | +# Locally stored "Eclipse launch configurations" | ||
50 | +*.launch | ||
51 | + | ||
52 | +# PyDev specific (Python IDE for Eclipse) | ||
53 | +*.pydevproject | ||
54 | + | ||
55 | +# CDT-specific (C/C++ Development Tooling) | ||
56 | +.cproject | ||
57 | + | ||
58 | +# CDT- autotools | ||
59 | +.autotools | ||
60 | + | ||
61 | +# Java annotation processor (APT) | ||
62 | +.factorypath | ||
63 | + | ||
64 | +# PDT-specific (PHP Development Tools) | ||
65 | +.buildpath | ||
66 | + | ||
67 | +# sbteclipse plugin | ||
68 | +.target | ||
69 | + | ||
70 | +# Tern plugin | ||
71 | +.tern-project | ||
72 | + | ||
73 | +# TeXlipse plugin | ||
74 | +.texlipse | ||
75 | + | ||
76 | +# STS (Spring Tool Suite) | ||
77 | +.springBeans | ||
78 | + | ||
79 | +# Code Recommenders | ||
80 | +.recommenders/ | ||
81 | + | ||
82 | +# Annotation Processing | ||
83 | +.apt_generated/ | ||
84 | +.apt_generated_test/ | ||
85 | + | ||
86 | +# Scala IDE specific (Scala & Java development for Eclipse) | ||
87 | +.cache-main | ||
88 | +.scala_dependencies | ||
89 | +.worksheet | ||
90 | + | ||
91 | +# Uncomment this line if you wish to ignore the project description file. | ||
92 | +# Typically, this file would be tracked if it contains build/dependency configurations: | ||
93 | +#.project | ||
94 | + | ||
95 | +build/ | ||
96 | +bin/ | ||
97 | +.**/ | ||
98 | +.ptp* |
CMakeLists.txt
0 → 100644
1 | +++ a/CMakeLists.txt | ||
1 | +cmake_minimum_required(VERSION 3.18) | ||
2 | +project(plasma-processing LANGUAGES CUDA CXX) | ||
3 | + | ||
4 | +# Link this 'library' to set the c++ standard / compile-time options requested | ||
5 | +add_library(project_options INTERFACE) | ||
6 | +target_compile_features(project_options INTERFACE cxx_std_17) | ||
7 | + | ||
8 | +set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${PROJECT_SOURCE_DIR}/bin) | ||
9 | + | ||
10 | +add_subdirectory(src) | ||
0 | \ No newline at end of file | 11 | \ No newline at end of file |
README.md
0 → 100644
1 | +++ a/README.md | ||
1 | +# Implementation of Thermal Event Image Processing Algorithms on Nvidia Tegra Embedded System on a Chip | ||
2 | + | ||
3 | +For details see Implementation of Thermal Event Image Processing Algorithms on Nvidia Tegra Embedded System on a Chip by Bartlomiej Jablonski, Dariusz Makowski, and Piotr Perek. | ||
4 | + | ||
5 | +## Dependencies | ||
6 | +- `C++17` | ||
7 | +- `CMake 3.18` | ||
8 | +- `CUDA 10.2` | ||
9 | +- `CUDAToolkit 10.2.89` | ||
10 | +- `HDF5 1.10.0.1` | ||
11 | +- `OpenCV 4.4.0` (requires compilation with `CUDA` support) |
debug_build.sh
0 → 100755
release_build.sh
0 → 100755
src/CMakeLists.txt
0 → 100644
1 | +++ a/src/CMakeLists.txt | ||
1 | +# OpenCV | ||
2 | +find_package(OpenCV REQUIRED) | ||
3 | + | ||
4 | +# NPP | ||
5 | +find_package(CUDAToolkit REQUIRED COMPONENTS nppif cudart) | ||
6 | + | ||
7 | +# HDF5 | ||
8 | +find_package(HDF5 REQUIRED COMPONENTS CXX) | ||
9 | + | ||
10 | +# Threads | ||
11 | +find_package(Threads REQUIRED) | ||
12 | + | ||
13 | +# Project files | ||
14 | +file(GLOB_RECURSE source_files CONFIGURE_DEPENDS *.cpp) | ||
15 | +file(GLOB_RECURSE header_files CONFIGURE_DEPENDS *.hpp) | ||
16 | +file(GLOB_RECURSE cuda_files CONFIGURE_DEPENDS *.cu) | ||
17 | + | ||
18 | +add_executable(executable ${source_files} ${header_files} ${cuda_files}) | ||
19 | + | ||
20 | +# CUDA compute architectures | ||
21 | +set_property(TARGET executable PROPERTY CUDA_ARCHITECTURES 62) # 62 - Jetson TX2 | ||
22 | +set_property(TARGET executable PROPERTY CUDA_STANDARD 14) | ||
23 | + | ||
24 | +# Allows to use absolute includes in regard to source directory | ||
25 | +target_include_directories(executable PUBLIC ${PROJECT_SOURCE_DIR}/src ${HDF5_INCLUDE_DIRS}) | ||
26 | + | ||
27 | +target_link_libraries(executable PUBLIC | ||
28 | + project_options | ||
29 | + CUDA::cudart | ||
30 | + CUDA::nppif | ||
31 | + ${OpenCV_LIBS} | ||
32 | + PRIVATE | ||
33 | + ${HDF5_CXX_LIBRARIES} | ||
34 | + Threads::Threads) | ||
0 | \ No newline at end of file | 35 | \ No newline at end of file |
src/algorithm/Algorithm.hpp
0 → 100644
1 | +++ a/src/algorithm/Algorithm.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/core/mat.hpp> | ||
4 | + | ||
5 | +#include "data/hdf5/CVMatLoader.hpp" | ||
6 | + | ||
7 | +class Algorithm { | ||
8 | + public: | ||
9 | + Algorithm() = default; | ||
10 | + Algorithm(const Algorithm &) = default; | ||
11 | + Algorithm(Algorithm &&) noexcept = default; | ||
12 | + Algorithm &operator=(Algorithm &&) noexcept = default; | ||
13 | + Algorithm &operator=(const Algorithm &) = default; | ||
14 | + | ||
15 | + virtual void setup(const cv::Size &frameSize, | ||
16 | + const CVMatLoader &loader) = 0; | ||
17 | + | ||
18 | + virtual void handleFrame(const cv::Mat &hostFrame, | ||
19 | + unsigned long timestamp) = 0; | ||
20 | + | ||
21 | + virtual ~Algorithm() = default; | ||
22 | +}; |
src/algorithm/common/PfcSegmenter.cpp
0 → 100644
1 | +++ a/src/algorithm/common/PfcSegmenter.cpp | ||
1 | +#include "PfcSegmenter.hpp" | ||
2 | + | ||
3 | +const std::array<PfcSegmenter::Mapping, 5> PfcSegmenter::mappings{ | ||
4 | + Mapping(400, cv::Range(400'000'000, 600'000'000), "Baffles"), | ||
5 | + Mapping(1800, cv::Range(1'000'000'000, 1'200'000'000), "Divertors"), | ||
6 | + Mapping(200, cv::Range(1'300'000'000, 1'400'000'000), | ||
7 | + "Wall Panels & Poloidal Closure"), | ||
8 | + Mapping(200, cv::Range(1'500'000'000, 1'700'000'000), "Pumping Gap"), | ||
9 | + Mapping(400, cv::Range(2'000'000'000, 2'147'483'647), "Wall Heat Shields")}; | ||
0 | \ No newline at end of file | 10 | \ No newline at end of file |
src/algorithm/common/PfcSegmenter.hpp
0 → 100644
1 | +++ a/src/algorithm/common/PfcSegmenter.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <array> | ||
4 | +#include <functional> | ||
5 | + | ||
6 | +#include <opencv2/imgproc.hpp> | ||
7 | + | ||
8 | +#include "data/hdf5/CVMatLoader.hpp" | ||
9 | + | ||
10 | +class PfcSegmenter { | ||
11 | + public: | ||
12 | + PfcSegmenter(const CVMatLoader &loader, const cv::Size &size) : size(size) { | ||
13 | + hostPfc = loader.asMat("/scene_model/PFC", CV_32F); | ||
14 | + } | ||
15 | + | ||
16 | + cv::Mat segment() const { | ||
17 | + return process( | ||
18 | + CV_16UC1, [](const Mapping &mapping, cv::Mat &component) { | ||
19 | + component.convertTo(component, CV_16UC1, (1L << 15) - 1); | ||
20 | + return component & mapping.value; | ||
21 | + }); | ||
22 | + } | ||
23 | + | ||
24 | + void contour(std::vector<std::vector<std::vector<cv::Point>>> &contours, | ||
25 | + cv::Mat &components) const { | ||
26 | + int i = 1; | ||
27 | + components = | ||
28 | + process(CV_8UC1, [&i, &contours](const Mapping & /*mapping*/, | ||
29 | + cv::Mat &component) { | ||
30 | + cv::threshold(component, component, 0, i++, cv::THRESH_BINARY); | ||
31 | + | ||
32 | + std::vector<std::vector<cv::Point>> contour; | ||
33 | + cv::findContours(component, contour, cv::RETR_TREE, | ||
34 | + cv::CHAIN_APPROX_SIMPLE); | ||
35 | + | ||
36 | + contours.emplace_back(std::move(contour)); | ||
37 | + | ||
38 | + return component; | ||
39 | + }); | ||
40 | + } | ||
41 | + | ||
42 | + cv::Mat mask(const std::string &label) const { | ||
43 | + return process(CV_8UC1, [&label](const Mapping &mapping, | ||
44 | + cv::Mat &component) { | ||
45 | + if (label == mapping.label) { | ||
46 | + return component; | ||
47 | + } | ||
48 | + | ||
49 | + return cv::Mat{cv::Mat::zeros(component.size(), component.type())}; | ||
50 | + }); | ||
51 | + } | ||
52 | + | ||
53 | + private: | ||
54 | + const cv::Size size; | ||
55 | + cv::Mat hostPfc; | ||
56 | + | ||
57 | + class Mapping { | ||
58 | + public: | ||
59 | + Mapping(int value, const cv::Range &range, | ||
60 | + const std::string &label = "") | ||
61 | + : value(value), range(range), label(label) {} | ||
62 | + | ||
63 | + const int value; | ||
64 | + const cv::Range range; | ||
65 | + const std::string label; | ||
66 | + }; | ||
67 | + | ||
68 | + static const std::array<Mapping, 5> mappings; | ||
69 | + | ||
70 | + cv::Mat process(int outputType, | ||
71 | + const std::function<cv::Mat(const Mapping &, cv::Mat &)> | ||
72 | + &modifyMask) const { | ||
73 | + cv::Mat output(size, outputType, cv::Scalar(0)); | ||
74 | + | ||
75 | + const cv::Mat kernel = | ||
76 | + cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 1)); | ||
77 | + | ||
78 | + cv::Mat component; | ||
79 | + for (const auto &mapping : mappings) { | ||
80 | + /* Workaround as the image is unsigned 32-bit thus it allows values | ||
81 | + * above the signed limit. It is assumed that values greater than | ||
82 | + * the signed maximum belongs to component, which end boundary is | ||
83 | + * at the limit | ||
84 | + */ | ||
85 | + /* Threshold PFC identifier range */ | ||
86 | + const auto upperLimit = 2'147'483'647; | ||
87 | + if (mapping.range.end >= upperLimit) { | ||
88 | + component = hostPfc >= upperLimit; | ||
89 | + } else { | ||
90 | + cv::inRange(hostPfc, mapping.range.start, mapping.range.end, | ||
91 | + component); | ||
92 | + } | ||
93 | + /* Apply close morphology */ | ||
94 | + cv::morphologyEx(component, component, cv::MORPH_CLOSE, kernel); | ||
95 | + /* Apply median filter */ | ||
96 | + cv::medianBlur(component, component, 3); | ||
97 | + | ||
98 | + /* Set pixels to operational limit value */ | ||
99 | + output = output | modifyMask(mapping, component); | ||
100 | + } | ||
101 | + | ||
102 | + /* Segmented PFC limits mask */ | ||
103 | + return output; | ||
104 | + } | ||
105 | +}; |
src/algorithm/common/SurfaceMap.hpp
0 → 100644
1 | +++ a/src/algorithm/common/SurfaceMap.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include "data/ResourceLocator.hpp" | ||
4 | +#include "data/hdf5/CVMatLoader.hpp" | ||
5 | + | ||
6 | +class SurfaceMap { | ||
7 | + public: | ||
8 | + explicit SurfaceMap(const cv::Size &size) : surface(buildSurface(size)) {} | ||
9 | + | ||
10 | + cv::Point3f at(int column, int row) const { | ||
11 | + return surface.at<cv::Point3f>(row, column); | ||
12 | + } | ||
13 | + | ||
14 | + cv::Point3f at(const cv::Point &point) const { | ||
15 | + return at(point.x, point.y); | ||
16 | + } | ||
17 | + | ||
18 | + std::vector<cv::Point3f> at(const std::vector<cv::Point> &points) const { | ||
19 | + std::vector<cv::Point3f> surfacePoints(points.size()); | ||
20 | + | ||
21 | + std::transform(points.begin(), points.end(), surfacePoints.begin(), | ||
22 | + [&](const cv::Point &point) { return at(point); }); | ||
23 | + | ||
24 | + return surfacePoints; | ||
25 | + } | ||
26 | + | ||
27 | + private: | ||
28 | + const cv::Mat surface; | ||
29 | + | ||
30 | + static cv::Mat buildSurface(const cv::Size &size) { | ||
31 | + cv::Mat result(size, CV_32FC3); | ||
32 | + | ||
33 | + const CVMatLoader loader( | ||
34 | + ResourceLocator::getPathProvider().path("20171114.053_AEF20.h5")); | ||
35 | + const std::vector<cv::Mat> coordinates = { | ||
36 | + loader.asMat("/scene_model/x", CV_32F), | ||
37 | + loader.asMat("/scene_model/y", CV_32F), | ||
38 | + loader.asMat("/scene_model/z", CV_32F), | ||
39 | + }; | ||
40 | + | ||
41 | + cv::merge(coordinates, result); | ||
42 | + return result; | ||
43 | + } | ||
44 | +}; |
src/algorithm/cpu/common/ClusterCorrespondence.hpp
0 → 100644
1 | +++ a/src/algorithm/cpu/common/ClusterCorrespondence.hpp | ||
1 | +/* Analysis of the outer divertor hot spot activity in the protection video | ||
2 | +camera recordings at JET */ | ||
3 | +#pragma once | ||
4 | + | ||
5 | +#include <opencv2/imgproc.hpp> | ||
6 | + | ||
7 | +namespace Cpu { | ||
8 | + | ||
9 | +class ClusterCorrespondence { | ||
10 | + public: | ||
11 | + ClusterCorrespondence(float minOverlap, float maxOversize) | ||
12 | + : minOverlap(minOverlap), maxOversize(maxOversize) {} | ||
13 | + | ||
14 | + bool corresponds(const cv::Mat &first, const cv::Mat &second) { | ||
15 | + cv::bitwise_and(first, second, overlap); | ||
16 | + const auto firstArea = cv::countNonZero(first), | ||
17 | + secondArea = cv::countNonZero(second), | ||
18 | + overlapArea = cv::countNonZero(overlap); | ||
19 | + | ||
20 | + return isOverlapping(firstArea, secondArea, overlapArea) && | ||
21 | + isSimilar(firstArea, secondArea); | ||
22 | + } | ||
23 | + | ||
24 | + private: | ||
25 | + float minOverlap; | ||
26 | + float maxOversize; | ||
27 | + | ||
28 | + bool isOverlapping(int firstArea, int secondArea, int overlapArea) const { | ||
29 | + return overlapArea >= minOverlap * std::min(firstArea, secondArea); | ||
30 | + } | ||
31 | + | ||
32 | + bool isSimilar(int firstArea, int secondArea) const { | ||
33 | + return (firstArea / static_cast<float>(secondArea) >= | ||
34 | + 1.F / maxOversize) && | ||
35 | + (firstArea / static_cast<float>(secondArea) <= maxOversize); | ||
36 | + } | ||
37 | + | ||
38 | + /* Buffors */ | ||
39 | + cv::Mat overlap; | ||
40 | +}; | ||
41 | + | ||
42 | +} // namespace Cpu | ||
0 | \ No newline at end of file | 43 | \ No newline at end of file |
src/algorithm/cpu/hotspot/Hotspot.cpp
0 → 100644
1 | +++ a/src/algorithm/cpu/hotspot/Hotspot.cpp | ||
1 | +#include "Hotspot.hpp" | ||
2 | + | ||
3 | +#include <opencv2/imgproc.hpp> | ||
4 | + | ||
5 | +Cpu::Hotspot::Hotspot(const cv::Mat &temperature, unsigned long timestamp, | ||
6 | + int component, const Contour &contour, | ||
7 | + const Surface &surface) | ||
8 | + : Hotspot(temperature, timestamp, component, contour, surface, | ||
9 | + createMask(temperature.size(), contour)) {} | ||
10 | + | ||
11 | +Cpu::Hotspot::Hotspot(const cv::Mat &temperature, unsigned long timestamp, | ||
12 | + int component, const Contour &contour, | ||
13 | + const Surface &surface, const cv::Mat &mask) | ||
14 | + : timestamp(timestamp), component(component), contour(contour), mask(mask), | ||
15 | + /* Compute hotspot attributes */ | ||
16 | + meanTemperature(calculateMean(temperature)), | ||
17 | + maxTemperature(calculateMax(temperature)), size(calculateSize()), | ||
18 | + surface(surface) {} | ||
19 | + | ||
20 | +double Cpu::Hotspot::perimeter() const { | ||
21 | + double perimeter = 0; | ||
22 | + | ||
23 | + for (std::size_t i = 0, size = surface.size(); i < size; ++i) { | ||
24 | + perimeter += cv::norm(surface[i] - surface[(i + 1) % size]); | ||
25 | + } | ||
26 | + | ||
27 | + return perimeter; | ||
28 | +} | ||
29 | + | ||
30 | +int Cpu::Hotspot::calculateSize() const { return cv::countNonZero(mask); } | ||
31 | + | ||
32 | +double Cpu::Hotspot::calculateMean(const cv::Mat &temperature) const { | ||
33 | + return cv::mean(temperature, mask)[0]; | ||
34 | +} | ||
35 | + | ||
36 | +double Cpu::Hotspot::calculateMax(const cv::Mat &temperature) const { | ||
37 | + double maxValue{-1}; | ||
38 | + cv::minMaxIdx(temperature, nullptr, &maxValue, nullptr, nullptr, mask); | ||
39 | + return maxValue; | ||
40 | +} | ||
41 | + | ||
42 | +cv::Mat Cpu::Hotspot::createMask(const cv::Size &size, const Contour &contour) { | ||
43 | + cv::Mat result(cv::Mat::zeros(size, CV_8UC1)); | ||
44 | + cv::drawContours(result, std::vector<Contour>{contour}, -1, 255, | ||
45 | + cv::FILLED); | ||
46 | + | ||
47 | + return result; | ||
48 | +} | ||
49 | + | ||
50 | +Cpu::PersistentHotspot::PersistentHotspot(const Hotspot &initial) | ||
51 | + : mask(initial.mask), component(initial.component) { | ||
52 | + append(initial); | ||
53 | +} | ||
54 | + | ||
55 | +void Cpu::PersistentHotspot::merge(const Hotspot &other) { | ||
56 | + cv::bitwise_or(mask, other.mask, mask); | ||
57 | + append(other); | ||
58 | +} | ||
59 | + | ||
60 | +int Cpu::PersistentHotspot::area() const { return cv::countNonZero(mask); } | ||
61 | + | ||
62 | +void Cpu::PersistentHotspot::append(const Hotspot &other) { | ||
63 | + timestamps.push_back(other.timestamp); | ||
64 | + meanTemperatures.push_back(other.meanTemperature); | ||
65 | + maxTemperatures.push_back(other.maxTemperature); | ||
66 | +} | ||
0 | \ No newline at end of file | 67 | \ No newline at end of file |
src/algorithm/cpu/hotspot/Hotspot.hpp
0 → 100644
1 | +++ a/src/algorithm/cpu/hotspot/Hotspot.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <vector> | ||
4 | + | ||
5 | +#include <opencv2/core/mat.hpp> | ||
6 | +#include <opencv2/imgproc.hpp> | ||
7 | + | ||
8 | +namespace Cpu { | ||
9 | + | ||
10 | +class Hotspot { | ||
11 | + public: | ||
12 | + using Contour = std::vector<cv::Point>; | ||
13 | + using Surface = std::vector<cv::Point3f>; | ||
14 | + | ||
15 | + const unsigned long timestamp; | ||
16 | + const int component; | ||
17 | + const Contour contour; | ||
18 | + | ||
19 | + const cv::Mat mask; | ||
20 | + | ||
21 | + const double meanTemperature, maxTemperature; | ||
22 | + const int size; | ||
23 | + | ||
24 | + Hotspot(const cv::Mat &temperature, unsigned long timestamp, int component, | ||
25 | + const Contour &contour, const Surface &surface); | ||
26 | + | ||
27 | + double perimeter() const; | ||
28 | + | ||
29 | + private: | ||
30 | + const Surface surface; | ||
31 | + | ||
32 | + Hotspot(const cv::Mat &temperature, unsigned long timestamp, int component, | ||
33 | + const Contour &contour, const Surface &surface, | ||
34 | + const cv::Mat &mask); | ||
35 | + | ||
36 | + int calculateSize() const; | ||
37 | + double calculateMean(const cv::Mat &temperature) const; | ||
38 | + double calculateMax(const cv::Mat &temperature) const; | ||
39 | + | ||
40 | + static cv::Mat createMask(const cv::Size &size, const Contour &contour); | ||
41 | +}; | ||
42 | + | ||
43 | +class PersistentHotspot { | ||
44 | + public: | ||
45 | + explicit PersistentHotspot(const Hotspot &initial); | ||
46 | + | ||
47 | + void merge(const Hotspot &other); | ||
48 | + int area() const; | ||
49 | + | ||
50 | + cv::Mat mask; | ||
51 | + int component; | ||
52 | + std::vector<unsigned long> timestamps; | ||
53 | + std::vector<double> meanTemperatures; | ||
54 | + std::vector<double> maxTemperatures; | ||
55 | + | ||
56 | + private: | ||
57 | + void append(const Hotspot &other); | ||
58 | +}; | ||
59 | + | ||
60 | +} // namespace Cpu | ||
0 | \ No newline at end of file | 61 | \ No newline at end of file |
src/algorithm/cpu/hotspot/OverloadHotspotDetectionAlgorithm.cpp
0 → 100644
1 | +++ a/src/algorithm/cpu/hotspot/OverloadHotspotDetectionAlgorithm.cpp | ||
1 | +#include "OverloadHotspotDetectionAlgorithm.hpp" | ||
2 | + | ||
3 | +#include <algorithm> | ||
4 | +#include <opencv2/imgproc.hpp> | ||
5 | + | ||
6 | +void Cpu::OverloadHotspotDetectionAlgorithm::setup(const cv::Size &frameSize, | ||
7 | + const CVMatLoader &loader) { | ||
8 | + currentFrameSize = frameSize; | ||
9 | + | ||
10 | + hostFov = loader.asMat("/scene_model/FOV", CV_8UC1); | ||
11 | + | ||
12 | + PfcSegmenter segmenter(loader, frameSize); | ||
13 | + hostPfc = segmenter.segment(); | ||
14 | + | ||
15 | + cv::Mat zerosMask; | ||
16 | + cv::compare(hostPfc, 0, zerosMask, cv::CMP_EQ); | ||
17 | + hostPfc.setTo(MAX_TEMPERATURE, zerosMask); | ||
18 | + | ||
19 | + cv::Mat model; | ||
20 | + cv::cvtColor(loader.asMat("/scene_model/CAD", CV_8UC1), model, | ||
21 | + cv::COLOR_GRAY2BGR); | ||
22 | + | ||
23 | + std::vector<Contours> pfcContours; | ||
24 | + segmenter.contour(pfcContours, hostPfcMask); | ||
25 | + | ||
26 | + surfaceMap = std::make_unique<const SurfaceMap>(frameSize); | ||
27 | +} | ||
28 | + | ||
29 | +void Cpu::OverloadHotspotDetectionAlgorithm::handleFrame( | ||
30 | + const cv::Mat &sourceFrame, unsigned long timestamp) { | ||
31 | + hostFrame = sourceFrame; | ||
32 | + | ||
33 | + /* Apply field-of-view mask */ | ||
34 | + cv::bitwise_and(hostFrame, hostFrame, hostInput, hostFov); | ||
35 | + /* Convert Kelvin to Celsius */ | ||
36 | + cv::subtract(hostInput, KELVINS, hostTemperature); | ||
37 | + | ||
38 | + /* Apply median filter */ | ||
39 | + medianBlur(hostTemperature, hostInput, 3); | ||
40 | + | ||
41 | + /* Threshold overheating pixels */ | ||
42 | + cv::compare(hostInput, hostPfc, overheating, cv::CMP_GT); | ||
43 | + | ||
44 | + currentHotspots.clear(); | ||
45 | + identifyHotspots(timestamp); | ||
46 | +} | ||
47 | + | ||
48 | +void Cpu::OverloadHotspotDetectionAlgorithm::identifyHotspots( | ||
49 | + unsigned long timestamp) { | ||
50 | + /* Analyse topologial structure */ | ||
51 | + Contours hotspotContours; | ||
52 | + cv::findContours(overheating, hotspotContours, cv::RETR_TREE, | ||
53 | + cv::CHAIN_APPROX_NONE); | ||
54 | + | ||
55 | + for (size_t i = 0, size = hotspotContours.size(); | ||
56 | + i < size && i < blobsLimit; ++i) { | ||
57 | + const auto &contour = hotspotContours[i]; | ||
58 | + /* Discard blobs below minimum area */ | ||
59 | + if (contour.size() >= pixelsThreshold) { | ||
60 | + const int component = hostPfcMask.at<uchar>(contour[0]); | ||
61 | + assert(component > 0); | ||
62 | + | ||
63 | + Cpu::Hotspot hotspot(hostTemperature, timestamp, component, contour, | ||
64 | + surfaceMap->at(contour)); | ||
65 | + | ||
66 | + /* Find corresponding blobs */ | ||
67 | + matchHotspot(hotspot); | ||
68 | + currentHotspots.emplace_back(std::move(hotspot)); | ||
69 | + } | ||
70 | + } | ||
71 | +} | ||
72 | + | ||
73 | +void Cpu::OverloadHotspotDetectionAlgorithm::matchHotspot( | ||
74 | + const Cpu::Hotspot &hotspot) { | ||
75 | + for (auto &persistentHotspot : uniqueHotspots) { | ||
76 | + if (clusterCorrespondence.corresponds(persistentHotspot.mask, | ||
77 | + hotspot.mask)) { | ||
78 | + /* Merge matching blobs */ | ||
79 | + persistentHotspot.merge(hotspot); | ||
80 | + return; | ||
81 | + } | ||
82 | + } | ||
83 | + uniqueHotspots.emplace_back(hotspot); | ||
84 | +} |
src/algorithm/cpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp
0 → 100644
1 | +++ a/src/algorithm/cpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/core/mat.hpp> | ||
4 | + | ||
5 | +#include "algorithm/Algorithm.hpp" | ||
6 | + | ||
7 | +#include "algorithm/cpu/common/ClusterCorrespondence.hpp" | ||
8 | +#include "algorithm/cpu/hotspot/Hotspot.hpp" | ||
9 | + | ||
10 | +#include "algorithm/common/PfcSegmenter.hpp" | ||
11 | +#include "algorithm/common/SurfaceMap.hpp" | ||
12 | + | ||
13 | +namespace Cpu { | ||
14 | + | ||
15 | +class OverloadHotspotDetectionAlgorithm : public Algorithm { | ||
16 | + public: | ||
17 | + /* Configurable parameters */ | ||
18 | + std::size_t blobsLimit{5}; | ||
19 | + std::size_t pixelsThreshold{3}; | ||
20 | + /* Configurable parameters */ | ||
21 | + | ||
22 | + OverloadHotspotDetectionAlgorithm() = default; | ||
23 | + | ||
24 | + void setup(const cv::Size &frameSize, const CVMatLoader &loader) override; | ||
25 | + void handleFrame(const cv::Mat &sourceFrame, | ||
26 | + unsigned long timestamp) override; | ||
27 | + | ||
28 | + std::vector<Cpu::PersistentHotspot> uniqueHotspots; /* Set of persistent hotspots */ | ||
29 | + | ||
30 | + private: | ||
31 | + using Contours = std::vector<std::vector<cv::Point>>; | ||
32 | + | ||
33 | + static constexpr int KELVINS{273}; | ||
34 | + static constexpr int MAX_TEMPERATURE{(1L << 15) - 1}; | ||
35 | + | ||
36 | + cv::Size currentFrameSize; | ||
37 | + | ||
38 | + Cpu::ClusterCorrespondence clusterCorrespondence{.8, 2.}; | ||
39 | + std::unique_ptr<const SurfaceMap> surfaceMap; | ||
40 | + | ||
41 | + cv::Mat hostFov, hostModel, hostPfc; | ||
42 | + cv::Mat hostPfcMask; | ||
43 | + | ||
44 | + std::vector<Cpu::Hotspot> currentHotspots; | ||
45 | + | ||
46 | + void identifyHotspots(unsigned long timestamp); | ||
47 | + void matchHotspot(const Cpu::Hotspot &hotspot); | ||
48 | + | ||
49 | + /* Buffors */ | ||
50 | + cv::Mat overheating; | ||
51 | + cv::Mat hostFrame, hostInput, hostTemperature; | ||
52 | + /* Buffors */ | ||
53 | +}; | ||
54 | + | ||
55 | +} // namespace Cpu | ||
0 | \ No newline at end of file | 56 | \ No newline at end of file |
src/algorithm/cpu/surfacelayer/SurfaceLayerDetectionAlgorithm.cpp
0 → 100644
1 | +++ a/src/algorithm/cpu/surfacelayer/SurfaceLayerDetectionAlgorithm.cpp | ||
1 | +#include "SurfaceLayerDetectionAlgorithm.hpp" | ||
2 | + | ||
3 | +#include <opencv2/imgproc.hpp> | ||
4 | + | ||
5 | +#include "algorithm/common/PfcSegmenter.hpp" | ||
6 | + | ||
7 | +cv::Mat | ||
8 | +Cpu::SurfaceLayerDetectionAlgorithm::setMinimumTemperature(ushort temperature) { | ||
9 | + hostMinTemperature.setTo(temperature); | ||
10 | + return hostMinTemperature; | ||
11 | +} | ||
12 | + | ||
13 | +void Cpu::SurfaceLayerDetectionAlgorithm::setup(const cv::Size &frameSize, | ||
14 | + const CVMatLoader &loader) { | ||
15 | + | ||
16 | + hostRegionMask = loader.asMat("/scene_model/FOV", CV_8UC1) & | ||
17 | + PfcSegmenter(loader, frameSize).mask("Divertors"); | ||
18 | + | ||
19 | + hostMinTemperature.create(frameSize, CV_16UC1); | ||
20 | + setMinimumTemperature(434); | ||
21 | + | ||
22 | + hostFloatInput.create(frameSize, CV_32F); | ||
23 | + hostDifference.create(frameSize, CV_32F); | ||
24 | + hostDifference.setTo(cv::Scalar::all(0)); | ||
25 | + | ||
26 | + hostDifferenceAbs.create(frameSize, CV_32F); | ||
27 | + hostDifferenceAbs.setTo(cv::Scalar::all(0)); | ||
28 | + | ||
29 | + hostSurfaceLayers.create(frameSize, CV_8UC1); | ||
30 | + hostSurfaceLayers.setTo(cv::Scalar::all(0)); | ||
31 | +} | ||
32 | + | ||
33 | +void Cpu::SurfaceLayerDetectionAlgorithm::handleFrame( | ||
34 | + const cv::Mat &sourceFrame, unsigned long timestamp) { | ||
35 | + currentTimestamp = timestamp; | ||
36 | + hostFrame = sourceFrame; | ||
37 | + | ||
38 | + /* Apply median filter */ | ||
39 | + cv::medianBlur(hostFrame, hostFiltered, 3); | ||
40 | + /* Apply detection PFC mask */ | ||
41 | + cv::bitwise_and(hostFiltered, hostFiltered, hostInput, hostRegionMask); | ||
42 | + | ||
43 | + /* Trim noise at minimum temperature */ | ||
44 | + cv::max(hostInput, hostMinTemperature, hostInput); | ||
45 | + | ||
46 | + /* Convert to float */ | ||
47 | + hostInput.convertTo(hostFloatInput, CV_32F); | ||
48 | + /* Calculate normalized derivative of the temperature evolution */ | ||
49 | + normalizedDerivative(); /* Result in hostDifference & hostDifferenceAbs */ | ||
50 | + | ||
51 | + /* Threshold rapidly evolving pixels */ | ||
52 | + cv::compare(hostDifferenceAbs, derivativeThreshold, hostResult, cv::CMP_GE); | ||
53 | + /* Merge surface layers */ | ||
54 | + cv::bitwise_or(hostResult, hostSurfaceLayers, hostSurfaceLayers); | ||
55 | +} | ||
56 | + | ||
57 | +void Cpu::SurfaceLayerDetectionAlgorithm::normalizedDerivative() { | ||
58 | + if (hostPreviousPreviousFrame.empty()) { | ||
59 | + previousPreviousTimestamp = currentTimestamp; | ||
60 | + hostFloatInput.copyTo(hostPreviousPreviousFrame); | ||
61 | + return; | ||
62 | + } | ||
63 | + if (hostPreviousFrame.empty()) { | ||
64 | + previousTimestamp = currentTimestamp; | ||
65 | + hostFloatInput.copyTo(hostPreviousFrame); | ||
66 | + return; | ||
67 | + } | ||
68 | + | ||
69 | + cv::subtract(hostFloatInput, hostPreviousPreviousFrame, hostDifference); | ||
70 | + cv::multiply(hostDifference, | ||
71 | + cv::Scalar::all( | ||
72 | + (currentTimestamp - previousPreviousTimestamp) * NS_TO_S), | ||
73 | + hostDifference); | ||
74 | + cv::divide(hostDifference, hostPreviousFrame, hostDifference); | ||
75 | + hostDifferenceAbs = cv::abs(hostDifference); | ||
76 | + | ||
77 | + previousPreviousTimestamp = previousTimestamp; | ||
78 | + previousTimestamp = currentTimestamp; | ||
79 | + | ||
80 | + hostPreviousFrame.copyTo(hostPreviousPreviousFrame); | ||
81 | + hostFloatInput.copyTo(hostPreviousFrame); | ||
82 | +} |
src/algorithm/cpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp
0 → 100644
1 | +++ a/src/algorithm/cpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/core/mat.hpp> | ||
4 | + | ||
5 | +#include "algorithm/Algorithm.hpp" | ||
6 | + | ||
7 | +namespace Cpu { | ||
8 | + | ||
9 | +class SurfaceLayerDetectionAlgorithm : public Algorithm { | ||
10 | + public: | ||
11 | + /* Configurable parameters */ | ||
12 | + cv::Scalar derivativeThreshold{7}; | ||
13 | + cv::Mat hostMinTemperature; | ||
14 | + | ||
15 | + cv::Mat setMinimumTemperature(ushort temperature); | ||
16 | + /* Configurable parameters */ | ||
17 | + | ||
18 | + SurfaceLayerDetectionAlgorithm() = default; | ||
19 | + | ||
20 | + void setup(const cv::Size &frameSize, const CVMatLoader &loader) override; | ||
21 | + void handleFrame(const cv::Mat &sourceFrame, | ||
22 | + unsigned long timestamp) override; | ||
23 | + | ||
24 | + cv::Mat hostSurfaceLayers; /* Surface layers mask */ | ||
25 | + | ||
26 | + private: | ||
27 | + /* ns to s and invert to multiply instead of divide */ | ||
28 | + static constexpr double NS_TO_S{1e-5}; | ||
29 | + | ||
30 | + cv::Mat hostRegionMask; | ||
31 | + | ||
32 | + cv::Mat hostPreviousPreviousFrame, hostPreviousFrame; | ||
33 | + unsigned long previousPreviousTimestamp{}, previousTimestamp{}, | ||
34 | + currentTimestamp{}; | ||
35 | + | ||
36 | + void normalizedDerivative(); | ||
37 | + | ||
38 | + /* Buffors */ | ||
39 | + cv::Mat hostFrame, hostFiltered, hostInput, hostFloatInput, hostDifference, | ||
40 | + hostDifferenceAbs, hostResult; | ||
41 | + /* Buffors */ | ||
42 | +}; | ||
43 | + | ||
44 | +} // namespace Cpu | ||
0 | \ No newline at end of file | 45 | \ No newline at end of file |
src/algorithm/gpu/common/ClusterCorrespondence.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/ClusterCorrespondence.hpp | ||
1 | +/* Analysis of the outer divertor hot spot activity in the protection video | ||
2 | +camera recordings at JET */ | ||
3 | +#pragma once | ||
4 | + | ||
5 | +#include <opencv2/cudaarithm.hpp> | ||
6 | + | ||
7 | +class ClusterCorrespondence { | ||
8 | + public: | ||
9 | + ClusterCorrespondence(float minOverlap, float maxOversize) | ||
10 | + : minOverlap(minOverlap), maxOversize(maxOversize) {} | ||
11 | + | ||
12 | + bool corresponds(const cv::cuda::GpuMat &first, | ||
13 | + const cv::cuda::GpuMat &second) { | ||
14 | + cv::cuda::bitwise_and(first, second, overlap, cv::noArray(), | ||
15 | + overlapStream); | ||
16 | + cv::cuda::countNonZero(first, deviceFirstArea, firstStream); | ||
17 | + cv::cuda::countNonZero(second, deviceSecondArea, secondStream), | ||
18 | + cv::cuda::countNonZero(overlap, deviceOverlapArea, overlapStream); | ||
19 | + | ||
20 | + deviceFirstArea.download(hostFirstArea, firstStream); | ||
21 | + deviceSecondArea.download(hostSecondArea, secondStream); | ||
22 | + deviceOverlapArea.download(hostOverlapArea, overlapStream); | ||
23 | + | ||
24 | + firstStream.waitForCompletion(); | ||
25 | + secondStream.waitForCompletion(); | ||
26 | + overlapStream.waitForCompletion(); | ||
27 | + return meetsCriteria(hostFirstArea.at<int>(0), | ||
28 | + hostSecondArea.at<int>(0), | ||
29 | + hostOverlapArea.at<int>(0)); | ||
30 | + } | ||
31 | + | ||
32 | + private: | ||
33 | + float minOverlap; | ||
34 | + float maxOversize; | ||
35 | + | ||
36 | + bool meetsCriteria(int firstArea, int secondArea, int overlapArea) const { | ||
37 | + return isOverlapping(firstArea, secondArea, overlapArea) && | ||
38 | + isSimilar(firstArea, secondArea); | ||
39 | + } | ||
40 | + | ||
41 | + bool isOverlapping(int firstArea, int secondArea, int overlapArea) const { | ||
42 | + return overlapArea >= minOverlap * std::min(firstArea, secondArea); | ||
43 | + } | ||
44 | + | ||
45 | + bool isSimilar(int firstArea, int secondArea) const { | ||
46 | + return (firstArea / static_cast<float>(secondArea) >= | ||
47 | + 1.F / maxOversize) && | ||
48 | + (firstArea / static_cast<float>(secondArea) <= maxOversize); | ||
49 | + } | ||
50 | + | ||
51 | + /* Buffors */ | ||
52 | + cv::cuda::Stream firstStream, secondStream, overlapStream; | ||
53 | + cv::cuda::GpuMat overlap, deviceFirstArea, deviceSecondArea, | ||
54 | + deviceOverlapArea; | ||
55 | + cv::Mat hostFirstArea, hostSecondArea, hostOverlapArea; | ||
56 | +}; | ||
0 | \ No newline at end of file | 57 | \ No newline at end of file |
src/algorithm/gpu/common/ConcurrentCopyExecute.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/ConcurrentCopyExecute.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <array> | ||
4 | +#include <functional> | ||
5 | + | ||
6 | +#include <opencv2/core/cuda.hpp> | ||
7 | + | ||
8 | +template <std::size_t N> class ConcurrentCopyExecute { | ||
9 | + public: | ||
10 | + /* Width should be divisible by N without any reminder */ | ||
11 | + void run(const std::function<void(const cv::Range &, const cv::Range &, | ||
12 | + cv::cuda::Stream &)> &callback, | ||
13 | + const cv::Size &size) { | ||
14 | + assert(size.height % N == 0); | ||
15 | + | ||
16 | + const cv::Range colRange(0, size.width); | ||
17 | + const int stride = size.height / N; | ||
18 | + for (auto i = 0, j = 0; i < size.height; i += stride, ++j) { | ||
19 | + callback(cv::Range{i, i + stride}, colRange, streams[j]); | ||
20 | + } | ||
21 | + } | ||
22 | + | ||
23 | + void synchronize() { | ||
24 | + for (auto &stream : streams) { | ||
25 | + stream.waitForCompletion(); | ||
26 | + } | ||
27 | + } | ||
28 | + | ||
29 | + private: | ||
30 | + std::array<cv::cuda::Stream, N> streams; | ||
31 | +}; | ||
0 | \ No newline at end of file | 32 | \ No newline at end of file |
src/algorithm/gpu/common/ContinuousGpuMat.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/ContinuousGpuMat.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <cuda_runtime.h> | ||
4 | + | ||
5 | +#include <opencv2/core/cuda.hpp> | ||
6 | + | ||
7 | +class ContinuousGpuMat { | ||
8 | + public: | ||
9 | + ContinuousGpuMat(const ContinuousGpuMat &) = delete; | ||
10 | + ContinuousGpuMat &operator=(const ContinuousGpuMat &) = delete; | ||
11 | + ContinuousGpuMat &operator=(ContinuousGpuMat &&) noexcept = delete; | ||
12 | + ContinuousGpuMat(ContinuousGpuMat &&) noexcept = delete; | ||
13 | + | ||
14 | + ContinuousGpuMat(const int rows, const int cols, const int type) { | ||
15 | + allocate(rows, cols, type); | ||
16 | + } | ||
17 | + | ||
18 | + ContinuousGpuMat(const cv::Size &size, const int type) | ||
19 | + : ContinuousGpuMat(size.height, size.width, type) {} | ||
20 | + | ||
21 | + explicit ContinuousGpuMat(cv::InputArray host) | ||
22 | + : ContinuousGpuMat(host.rows(), host.cols(), host.type()) { | ||
23 | + continuousMat.upload(host); | ||
24 | + } | ||
25 | + | ||
26 | + void create(const cv::Size &size, const int type) { | ||
27 | + allocate(size.height, size.width, type); | ||
28 | + } | ||
29 | + | ||
30 | + cv::Size size() const { return continuousMat.size(); } | ||
31 | + | ||
32 | + int type() const { return continuousMat.type(); } | ||
33 | + | ||
34 | + cv::cuda::GpuMat &device() { return continuousMat; } | ||
35 | + | ||
36 | + ~ContinuousGpuMat() { cudaFree(deviceMemory); } | ||
37 | + | ||
38 | + private: | ||
39 | + void *deviceMemory{nullptr}; | ||
40 | + cv::cuda::GpuMat continuousMat; | ||
41 | + | ||
42 | + void allocate(const int rows, const int cols, const int type) { | ||
43 | + cudaMalloc(&deviceMemory, rows * cols * cv::getElemSize(type)); | ||
44 | + continuousMat = cv::cuda::GpuMat(rows, cols, type, deviceMemory); | ||
45 | + } | ||
46 | +}; | ||
0 | \ No newline at end of file | 47 | \ No newline at end of file |
src/algorithm/gpu/common/ManagedMat.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/ManagedMat.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <cuda_runtime.h> | ||
4 | + | ||
5 | +#include <opencv2/core/utility.hpp> | ||
6 | +#include <opencv2/cudaarithm.hpp> | ||
7 | + | ||
8 | +class ManagedMat { | ||
9 | + public: | ||
10 | + ManagedMat(const ManagedMat &) = delete; | ||
11 | + ManagedMat &operator=(const ManagedMat &) = delete; | ||
12 | + ManagedMat &operator=(ManagedMat &&) noexcept = delete; | ||
13 | + ManagedMat(ManagedMat &&) noexcept = delete; | ||
14 | + | ||
15 | + ManagedMat(int rows, int cols, int type) | ||
16 | + : memorySize(rows * cols * cv::getElemSize(type)) { | ||
17 | + cudaMallocManaged(&managedMemory, memorySize); | ||
18 | + | ||
19 | + hostMat = cv::Mat(rows, cols, type, managedMemory); | ||
20 | + deviceMat = cv::cuda::GpuMat(rows, cols, type, managedMemory); | ||
21 | + } | ||
22 | + ManagedMat(const cv::Size &size, int type) | ||
23 | + : ManagedMat(size.height, size.width, type) {} | ||
24 | + | ||
25 | + cv::Mat &host() { | ||
26 | + prefetchToCpu(); | ||
27 | + return hostMat; | ||
28 | + } | ||
29 | + | ||
30 | + cv::cuda::GpuMat &device() { | ||
31 | + prefetchToGpu(); | ||
32 | + return deviceMat; | ||
33 | + } | ||
34 | + | ||
35 | + ~ManagedMat() { cudaFree(managedMemory); } | ||
36 | + | ||
37 | + private: | ||
38 | + const std::size_t memorySize; | ||
39 | + void *managedMemory{nullptr}; | ||
40 | + | ||
41 | + cv::Mat hostMat; | ||
42 | + cv::cuda::GpuMat deviceMat; | ||
43 | + | ||
44 | + void prefetchToGpu() const { | ||
45 | + cudaStreamAttachMemAsync(nullptr, managedMemory, memorySize, | ||
46 | + cudaMemAttachGlobal); | ||
47 | + } | ||
48 | + | ||
49 | + void prefetchToCpu() const { | ||
50 | + cudaStreamAttachMemAsync(nullptr, managedMemory, memorySize, | ||
51 | + cudaMemAttachHost); | ||
52 | + cudaStreamSynchronize(nullptr); | ||
53 | + } | ||
54 | +}; | ||
0 | \ No newline at end of file | 55 | \ No newline at end of file |
src/algorithm/gpu/common/MedianFilter.cpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/MedianFilter.cpp | ||
1 | +#include "MedianFilter.hpp" | ||
2 | + | ||
3 | +std::unique_ptr<cv::cuda::Filter> | ||
4 | +CudaFilter::createMedianFilter16U(const cv::Size &size, int kernel) { | ||
5 | + return std::make_unique<MedianFilter<Filter::U16Bit>>(size, kernel); | ||
6 | +} | ||
7 | + | ||
8 | +std::unique_ptr<cv::cuda::Filter> | ||
9 | +CudaFilter::createMedianFilter8U(const cv::Size &size, int kernel) { | ||
10 | + return std::make_unique<MedianFilter<Filter::U8Bit>>(size, kernel); | ||
11 | +} | ||
0 | \ No newline at end of file | 12 | \ No newline at end of file |
src/algorithm/gpu/common/MedianFilter.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/MedianFilter.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <cuda_runtime.h> | ||
4 | + | ||
5 | +#include <npp.h> | ||
6 | + | ||
7 | +#include <opencv2/cudaarithm.hpp> | ||
8 | +#include <opencv2/cudafilters.hpp> | ||
9 | + | ||
10 | +namespace Filter { | ||
11 | +struct U16Bit { | ||
12 | + template <typename... Args> | ||
13 | + NppStatus calculateBufferSize(Args &&... args) const { | ||
14 | + return nppiFilterMedianGetBufferSize_16u_C1R( | ||
15 | + std::forward<Args>(args)...); | ||
16 | + } | ||
17 | + | ||
18 | + template <typename... Args> NppStatus filter(Args &&... args) const { | ||
19 | + return nppiFilterMedian_16u_C1R(std::forward<Args>(args)...); | ||
20 | + } | ||
21 | + | ||
22 | + static constexpr int cvType{CV_16UC1}; | ||
23 | + using Type = Npp16u; | ||
24 | +}; | ||
25 | + | ||
26 | +struct U8Bit { | ||
27 | + template <typename... Args> | ||
28 | + NppStatus calculateBufferSize(Args &&... args) const { | ||
29 | + return nppiFilterMedianGetBufferSize_8u_C1R( | ||
30 | + std::forward<Args>(args)...); | ||
31 | + } | ||
32 | + | ||
33 | + template <typename... Args> NppStatus filter(Args &&... args) const { | ||
34 | + return nppiFilterMedian_8u_C1R(std::forward<Args>(args)...); | ||
35 | + } | ||
36 | + | ||
37 | + static constexpr int cvType{CV_8UC1}; | ||
38 | + using Type = Npp8u; | ||
39 | +}; | ||
40 | +} // namespace Filter | ||
41 | + | ||
42 | +/* | ||
43 | + * Filter uses 0 value to handle borders | ||
44 | + */ | ||
45 | +template <typename FilterType> class MedianFilter : public cv::cuda::Filter { | ||
46 | + public: | ||
47 | + MedianFilter(const MedianFilter &) = delete; | ||
48 | + MedianFilter &operator=(const MedianFilter &) = delete; | ||
49 | + MedianFilter &operator=(MedianFilter &&) noexcept = delete; | ||
50 | + MedianFilter(MedianFilter &&) noexcept = delete; | ||
51 | + | ||
52 | + MedianFilter(const cv::Size &size, int kernel) | ||
53 | + : size(size), kernelSize(NppiSize{kernel, kernel}), | ||
54 | + anchor(NppiPoint{kernel / 2, kernel / 2}), | ||
55 | + nppSize(NppiSize{size.width, size.height}), | ||
56 | + nppScratchBuffer(allocateScratch()) {} | ||
57 | + | ||
58 | + void apply(cv::InputArray source, cv::OutputArray destination, | ||
59 | + [[maybe_unused]] cv::cuda::Stream &stream = | ||
60 | + cv::cuda::Stream::Null()) override { | ||
61 | + /* Stream unsupported. If needed might be applied with NPPI Ctx function | ||
62 | + * variants */ | ||
63 | + assert(cv::cuda::Stream::Null() == stream); | ||
64 | + assert(filterDelegate.cvType == source.type()); | ||
65 | + | ||
66 | + destination.create(size, filterDelegate.cvType); | ||
67 | + cv::cuda::copyMakeBorder(source, paddedSource, anchor.y, anchor.y, | ||
68 | + anchor.x, anchor.x, cv::BORDER_CONSTANT, | ||
69 | + cv::Scalar(0)); | ||
70 | + | ||
71 | + [[maybe_unused]] const auto nppStatus = filterDelegate.filter( | ||
72 | + paddedSource.ptr<typename FilterType::Type>(anchor.y) + anchor.x, | ||
73 | + paddedSource.step, | ||
74 | + destination.getGpuMatRef().ptr<typename FilterType::Type>(), | ||
75 | + destination.step(), nppSize, kernelSize, anchor, nppScratchBuffer); | ||
76 | + assert(NPP_SUCCESS == nppStatus); | ||
77 | + } | ||
78 | + | ||
79 | + ~MedianFilter() override { cudaFree(nppScratchBuffer); } | ||
80 | + | ||
81 | + private: | ||
82 | + const FilterType filterDelegate{}; | ||
83 | + | ||
84 | + const cv::Size size; | ||
85 | + const NppiSize kernelSize; | ||
86 | + const NppiPoint anchor; | ||
87 | + const NppiSize nppSize; | ||
88 | + | ||
89 | + Npp8u *const nppScratchBuffer{nullptr}; | ||
90 | + | ||
91 | + Npp8u *allocateScratch() const { | ||
92 | + Npp32u bufferSize{0}; | ||
93 | + [[maybe_unused]] const auto nppStatus = | ||
94 | + filterDelegate.calculateBufferSize(nppSize, kernelSize, | ||
95 | + &bufferSize); | ||
96 | + assert(NPP_SUCCESS == nppStatus); | ||
97 | + | ||
98 | + Npp8u *buffer{nullptr}; | ||
99 | + [[maybe_unused]] const auto cudaStatus = | ||
100 | + cudaMalloc(&buffer, bufferSize); | ||
101 | + assert(cudaSuccess == cudaStatus); | ||
102 | + | ||
103 | + return buffer; | ||
104 | + } | ||
105 | + | ||
106 | + cv::cuda::GpuMat paddedSource; | ||
107 | +}; | ||
108 | + | ||
109 | +namespace CudaFilter { | ||
110 | + | ||
111 | +std::unique_ptr<cv::cuda::Filter> createMedianFilter16U(const cv::Size &size, | ||
112 | + int kernel); | ||
113 | +std::unique_ptr<cv::cuda::Filter> createMedianFilter8U(const cv::Size &size, | ||
114 | + int kernel); | ||
115 | + | ||
116 | +} // namespace CudaFilter | ||
0 | \ No newline at end of file | 117 | \ No newline at end of file |
src/algorithm/gpu/common/MedianFilter3x3.cu
0 → 100644
1 | +++ a/src/algorithm/gpu/common/MedianFilter3x3.cu | ||
1 | +#include "MedianFilter3x3.hpp" | ||
2 | + | ||
3 | +/* Efficient exchange */ | ||
4 | +#define swap(a, b) \ | ||
5 | + { \ | ||
6 | + T tmp = a; \ | ||
7 | + a = min(a, b); \ | ||
8 | + b = max(tmp, b); \ | ||
9 | + } | ||
10 | +#define mn3(a, b, c) \ | ||
11 | + swap(a, b); \ | ||
12 | + swap(a, c); | ||
13 | +#define mx3(a, b, c) \ | ||
14 | + swap(b, c); \ | ||
15 | + swap(a, c); | ||
16 | + | ||
17 | +#define mnmx3(a, b, c) \ | ||
18 | + mx3(a, b, c); \ | ||
19 | + swap(a, b); /* 3 exchanges */ | ||
20 | +#define mnmx4(a, b, c, d) \ | ||
21 | + swap(a, b); \ | ||
22 | + swap(c, d); \ | ||
23 | + swap(a, c); \ | ||
24 | + swap(b, d); /* 4 exchanges */ | ||
25 | +#define mnmx5(a, b, c, d, e) \ | ||
26 | + swap(a, b); \ | ||
27 | + swap(c, d); \ | ||
28 | + mn3(a, c, e); \ | ||
29 | + mx3(b, d, e); /* 6 exchanges */ | ||
30 | +#define mnmx6(a, b, c, d, e, f) \ | ||
31 | + swap(a, d); \ | ||
32 | + swap(b, e); \ | ||
33 | + swap(c, f); \ | ||
34 | + mn3(a, b, c); \ | ||
35 | + mx3(d, e, f); /* 7 exchanges */ | ||
36 | +/* Efficient exchange */ | ||
37 | + | ||
38 | +#define IDX(x, y) ((y)*width + (x)) | ||
39 | +#define SMEM(x, y) sharedMemory[(x) + 1][(y) + 1] | ||
40 | +#define GMEM(x, y) source[IDX(x, y)] | ||
41 | + | ||
42 | +template <typename T, int TileWidth, int TileHeight> | ||
43 | +__device__ T findMedian(const T sharedMemory[TileWidth + 2][TileHeight + 2], | ||
44 | + const int tx, const int ty) { | ||
45 | + /* Take upper six values from shared memory */ | ||
46 | + T v[6] = {SMEM(tx - 1, ty - 1), SMEM(tx, ty - 1), SMEM(tx + 1, ty - 1), | ||
47 | + SMEM(tx - 1, ty), SMEM(tx, ty), SMEM(tx + 1, ty)}; | ||
48 | + | ||
49 | + /* With each pass, remove min and max values and add new value */ | ||
50 | + mnmx6(v[0], v[1], v[2], v[3], v[4], v[5]); | ||
51 | + v[5] = SMEM(tx - 1, ty + 1); /* Add new value */ | ||
52 | + mnmx5(v[1], v[2], v[3], v[4], v[5]); | ||
53 | + v[5] = SMEM(tx, ty + 1); /* Add new value */ | ||
54 | + mnmx4(v[2], v[3], v[4], v[5]); | ||
55 | + v[5] = SMEM(tx + 1, ty + 1); /* Add new value */ | ||
56 | + mnmx3(v[3], v[4], v[5]); | ||
57 | + | ||
58 | + /* Select the middle value */ | ||
59 | + return v[4]; | ||
60 | +} | ||
61 | + | ||
62 | +template <typename T, int TileWidth, int TileHeight> | ||
63 | +__global__ void medianFilter3x3Kernel(const T *source, T *destination, | ||
64 | + const int width, const int height) { | ||
65 | + const int tx = threadIdx.x, ty = threadIdx.y; | ||
66 | + | ||
67 | + bool isLeftBorder = (tx == 0), isRightBorder = (tx == TileWidth - 1), | ||
68 | + isTopBorder = (ty == 0), isBottomBorder = (ty == TileHeight - 1); | ||
69 | + | ||
70 | + __shared__ T sharedMemory[TileWidth + 2][TileHeight + 2]; | ||
71 | + | ||
72 | + /* 0-initialise boundaries */ | ||
73 | + if (isLeftBorder) { | ||
74 | + SMEM(tx - 1, ty) = 0; | ||
75 | + } else if (isRightBorder) { | ||
76 | + SMEM(tx + 1, ty) = 0; | ||
77 | + } | ||
78 | + if (isTopBorder) { | ||
79 | + SMEM(tx, ty - 1) = 0; | ||
80 | + if (isLeftBorder) { | ||
81 | + SMEM(tx - 1, ty - 1) = 0; | ||
82 | + } else if (isRightBorder) { | ||
83 | + SMEM(tx + 1, ty - 1) = 0; | ||
84 | + } | ||
85 | + } else if (isBottomBorder) { | ||
86 | + SMEM(tx, ty + 1) = 0; | ||
87 | + if (isLeftBorder) { | ||
88 | + SMEM(tx - 1, ty + 1) = 0; | ||
89 | + } else if (isRightBorder) { | ||
90 | + SMEM(tx + 1, ty + 1) = 0; | ||
91 | + } | ||
92 | + } | ||
93 | + | ||
94 | + const int x = blockIdx.x * blockDim.x + tx, | ||
95 | + y = blockIdx.y * blockDim.y + ty; | ||
96 | + | ||
97 | + /* Check if padded values lies inside an image */ | ||
98 | + isLeftBorder &= (x > 0); | ||
99 | + isRightBorder &= (x < width - 1); | ||
100 | + isTopBorder &= (y > 0); | ||
101 | + isBottomBorder &= (y < height - 1); | ||
102 | + | ||
103 | + /* Set pixel value */ | ||
104 | + SMEM(tx, ty) = GMEM(x, y); | ||
105 | + | ||
106 | + /* Set boundaries */ | ||
107 | + if (isLeftBorder) { | ||
108 | + SMEM(tx - 1, ty) = GMEM(x - 1, y); | ||
109 | + } else if (isRightBorder) { | ||
110 | + SMEM(tx + 1, ty) = GMEM(x + 1, y); | ||
111 | + } | ||
112 | + if (isTopBorder) { | ||
113 | + SMEM(tx, ty - 1) = GMEM(x, y - 1); | ||
114 | + if (isLeftBorder) { | ||
115 | + SMEM(tx - 1, ty - 1) = GMEM(x - 1, y - 1); | ||
116 | + } else if (isRightBorder) { | ||
117 | + SMEM(tx + 1, ty - 1) = GMEM(x + 1, y - 1); | ||
118 | + } | ||
119 | + } else if (isBottomBorder) { | ||
120 | + SMEM(tx, ty + 1) = GMEM(x, y + 1); | ||
121 | + if (isLeftBorder) { | ||
122 | + SMEM(tx - 1, ty + 1) = GMEM(x - 1, y + 1); | ||
123 | + } else if (isRightBorder) { | ||
124 | + SMEM(tx + 1, ty + 1) = GMEM(x + 1, y + 1); | ||
125 | + } | ||
126 | + } | ||
127 | + __syncthreads(); | ||
128 | + | ||
129 | + destination[IDX(x, y)] = findMedian<T, TileWidth, TileHeight>(sharedMemory, tx, ty); | ||
130 | +} | ||
131 | + | ||
132 | +template <> | ||
133 | +void MedianFilter3x3<uchar, 16, 16>::callKernel(const uchar *source, uchar *destination, | ||
134 | + cudaStream_t &stream) const { | ||
135 | + medianFilter3x3Kernel<uchar, 16, 16><<<grid, block, 0, stream>>>( | ||
136 | + source, destination, size.width, size.height); | ||
137 | +} | ||
138 | + | ||
139 | +template <> | ||
140 | +void MedianFilter3x3<ushort, 16, 16>::callKernel(const ushort *source, | ||
141 | + ushort *destination, | ||
142 | + cudaStream_t &stream) const { | ||
143 | + medianFilter3x3Kernel<ushort, 16, 16><<<grid, block, 0, stream>>>( | ||
144 | + source, destination, size.width, size.height); | ||
145 | +} | ||
0 | \ No newline at end of file | 146 | \ No newline at end of file |
src/algorithm/gpu/common/MedianFilter3x3.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/common/MedianFilter3x3.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/core/cuda_stream_accessor.hpp> | ||
4 | +#include <opencv2/cudafilters.hpp> | ||
5 | + | ||
6 | +/* | ||
7 | + * Median Filter optimised for 3x3 kerner size and continous memory that use 0 | ||
8 | + * on borders. In order to obtain correct boundary values TileWidth and TileHeight | ||
9 | + * should be a multiple of the size | ||
10 | + */ | ||
11 | +template <typename T, std::size_t TileWidth=16, std::size_t TileHeight=TileWidth> class MedianFilter3x3 : public cv::cuda::Filter { | ||
12 | + public: | ||
13 | + static std::size_t gridSize(const std::size_t a, const std::size_t b) { | ||
14 | + assert(a % b == 0); | ||
15 | + return a / b + (a % b != 0); | ||
16 | + } | ||
17 | + | ||
18 | + explicit MedianFilter3x3(const cv::Size &size) | ||
19 | + : size(size), | ||
20 | + grid(dim3( | ||
21 | + gridSize(size.width, TileWidth), | ||
22 | + gridSize(size.height, TileHeight))), | ||
23 | + block(dim3(TileWidth, TileHeight)) {} | ||
24 | + | ||
25 | + void apply(cv::InputArray source, cv::OutputArray destination, | ||
26 | + cv::cuda::Stream &stream = cv::cuda::Stream::Null()) override { | ||
27 | + destination.create(size, source.type()); | ||
28 | + | ||
29 | + /* Requires continuous memories */ | ||
30 | + assert(source.getGpuMat().isContinuous()); | ||
31 | + assert(destination.getGpuMatRef().isContinuous()); | ||
32 | + | ||
33 | + auto cudaStream = cv::cuda::StreamAccessor::getStream(stream); | ||
34 | + callKernel(source.getGpuMat().ptr<T>(), | ||
35 | + destination.getGpuMatRef().ptr<T>(), cudaStream); | ||
36 | + | ||
37 | + stream.waitForCompletion(); | ||
38 | + } | ||
39 | + | ||
40 | + private: | ||
41 | + const cv::Size size; | ||
42 | + const dim3 grid, block; | ||
43 | + | ||
44 | + void callKernel(const T *source, T *destination, | ||
45 | + cudaStream_t &stream) const; | ||
46 | +}; | ||
0 | \ No newline at end of file | 47 | \ No newline at end of file |
src/algorithm/gpu/hotspot/Hotspot.cpp
0 → 100644
1 | +++ a/src/algorithm/gpu/hotspot/Hotspot.cpp | ||
1 | +#include "Hotspot.hpp" | ||
2 | + | ||
3 | +#include <opencv2/imgproc.hpp> | ||
4 | + | ||
5 | +Hotspot::Hotspot(const cv::cuda::GpuMat &temperature, unsigned long timestamp, | ||
6 | + int component, const Contour &contour, const Surface &surface) | ||
7 | + : Hotspot(temperature, timestamp, component, contour, surface, | ||
8 | + createMask(temperature.size(), contour)) {} | ||
9 | + | ||
10 | +Hotspot::Hotspot(const cv::cuda::GpuMat &temperature, unsigned long timestamp, | ||
11 | + int component, const Contour &contour, const Surface &surface, | ||
12 | + const cv::cuda::GpuMat &mask) | ||
13 | + : timestamp(timestamp), component(component), contour(contour), mask(mask), | ||
14 | + surface(surface) { | ||
15 | + /* Compute hotspot attributes */ | ||
16 | + cv::cuda::countNonZero(mask, nonZeroDevice, nonZeroStream); | ||
17 | + nonZeroDevice.download(nonZeroHost, nonZeroStream); | ||
18 | + | ||
19 | + cv::cuda::calcSum(temperature, sumDevice, mask, sumStream); | ||
20 | + sumDevice.download(sumHost, sumStream); | ||
21 | + | ||
22 | + cv::cuda::findMinMax(temperature, minMaxDevice, mask, minMaxStream); | ||
23 | + minMaxDevice.download(minMaxHost, minMaxStream); | ||
24 | + | ||
25 | + nonZeroStream.waitForCompletion(); | ||
26 | + size = nonZeroHost.at<int>(0); | ||
27 | + sumStream.waitForCompletion(); | ||
28 | + meanTemperature = sumHost.at<double>(0) / size; | ||
29 | + minMaxStream.waitForCompletion(); | ||
30 | + maxTemperature = minMaxHost.at<int>(1); | ||
31 | +} | ||
32 | + | ||
33 | +double Hotspot::perimeter() const { | ||
34 | + double perimeter = 0; | ||
35 | + | ||
36 | + for (std::size_t i = 0, size = surface.size(); i < size; ++i) { | ||
37 | + perimeter += cv::norm(surface[i] - surface[(i + 1) % size]); | ||
38 | + } | ||
39 | + | ||
40 | + return perimeter; | ||
41 | +} | ||
42 | + | ||
43 | +cv::cuda::GpuMat Hotspot::createMask(const cv::Size &size, | ||
44 | + const Contour &contour) { | ||
45 | + cv::Mat result(cv::Mat::zeros(size, CV_8UC1)); | ||
46 | + cv::drawContours(result, std::vector<Contour>{contour}, -1, 255, | ||
47 | + cv::FILLED); | ||
48 | + | ||
49 | + return cv::cuda::GpuMat{result}; | ||
50 | +} | ||
51 | + | ||
52 | +PersistentHotspot::PersistentHotspot(const Hotspot &initial) | ||
53 | + : mask(initial.mask), component(initial.component) { | ||
54 | + append(initial); | ||
55 | +} | ||
56 | + | ||
57 | +void PersistentHotspot::merge(const Hotspot &other) { | ||
58 | + cv::cuda::bitwise_or(mask, other.mask, mask, cv::noArray(), stream); | ||
59 | + append(other); | ||
60 | + stream.waitForCompletion(); | ||
61 | +} | ||
62 | + | ||
63 | +int PersistentHotspot::area() const { return cv::cuda::countNonZero(mask); } | ||
64 | + | ||
65 | +void PersistentHotspot::append(const Hotspot &other) { | ||
66 | + timestamps.push_back(other.timestamp); | ||
67 | + meanTemperatures.push_back(other.meanTemperature); | ||
68 | + maxTemperatures.push_back(other.maxTemperature); | ||
69 | +} | ||
0 | \ No newline at end of file | 70 | \ No newline at end of file |
src/algorithm/gpu/hotspot/Hotspot.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/hotspot/Hotspot.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <vector> | ||
4 | + | ||
5 | +#include <opencv2/core/mat.hpp> | ||
6 | +#include <opencv2/cudaarithm.hpp> | ||
7 | + | ||
8 | +class Hotspot { | ||
9 | + public: | ||
10 | + using Contour = std::vector<cv::Point>; | ||
11 | + using Surface = std::vector<cv::Point3f>; | ||
12 | + | ||
13 | + const unsigned long timestamp; | ||
14 | + const int component; | ||
15 | + const Contour contour; | ||
16 | + | ||
17 | + const cv::cuda::GpuMat mask; | ||
18 | + | ||
19 | + int size; | ||
20 | + double meanTemperature, maxTemperature; | ||
21 | + | ||
22 | + Hotspot(const cv::cuda::GpuMat &temperature, unsigned long timestamp, | ||
23 | + int component, const Contour &contour, const Surface &surface); | ||
24 | + | ||
25 | + double perimeter() const; | ||
26 | + | ||
27 | + private: | ||
28 | + const Surface surface; | ||
29 | + | ||
30 | + Hotspot(const cv::cuda::GpuMat &temperature, unsigned long timestamp, | ||
31 | + int component, const Contour &contour, const Surface &surface, | ||
32 | + const cv::cuda::GpuMat &mask); | ||
33 | + | ||
34 | + static cv::cuda::GpuMat createMask(const cv::Size &size, | ||
35 | + const Contour &contour); | ||
36 | + | ||
37 | + /* Buffors */ | ||
38 | + cv::cuda::Stream nonZeroStream, sumStream, minMaxStream; | ||
39 | + cv::cuda::GpuMat nonZeroDevice{1, 1, CV_32SC1}, sumDevice{1, 1, CV_64FC1}, | ||
40 | + minMaxDevice{1, 2, CV_32SC1}; | ||
41 | + cv::Mat nonZeroHost{1, 1, CV_32SC1}, sumHost{1, 1, CV_64FC1}, | ||
42 | + minMaxHost{1, 2, CV_32SC1}; | ||
43 | +}; | ||
44 | + | ||
45 | +class PersistentHotspot { | ||
46 | + public: | ||
47 | + explicit PersistentHotspot(const Hotspot &initial); | ||
48 | + | ||
49 | + void merge(const Hotspot &other); | ||
50 | + int area() const; | ||
51 | + | ||
52 | + cv::cuda::GpuMat mask; | ||
53 | + int component; | ||
54 | + std::vector<unsigned long> timestamps; | ||
55 | + std::vector<double> meanTemperatures; | ||
56 | + std::vector<double> maxTemperatures; | ||
57 | + | ||
58 | + private: | ||
59 | + void append(const Hotspot &other); | ||
60 | + | ||
61 | + cv::cuda::Stream stream; | ||
62 | +}; |
src/algorithm/gpu/hotspot/OverloadHotspotDetectionAlgorithm.cpp
0 → 100644
1 | +++ a/src/algorithm/gpu/hotspot/OverloadHotspotDetectionAlgorithm.cpp | ||
1 | +#include "OverloadHotspotDetectionAlgorithm.hpp" | ||
2 | + | ||
3 | +#include "algorithm/gpu/common/MedianFilter.hpp" | ||
4 | +#include "algorithm/gpu/common/MedianFilter3x3.hpp" | ||
5 | + | ||
6 | +void OverloadHotspotDetectionAlgorithm::setup(const cv::Size &frameSize, | ||
7 | + const CVMatLoader &loader) { | ||
8 | + currentFrameSize = frameSize; | ||
9 | + | ||
10 | + deviceFov.upload(loader.asMat("/scene_model/FOV", CV_8UC1)); | ||
11 | + | ||
12 | + PfcSegmenter segmenter(loader, frameSize); | ||
13 | + devicePfc.upload(segmenter.segment()); | ||
14 | + | ||
15 | + cv::cuda::GpuMat zerosMask; | ||
16 | + cv::cuda::compare(devicePfc, 0, zerosMask, cv::CMP_EQ); | ||
17 | + devicePfc.setTo(MAX_TEMPERATURE, zerosMask); | ||
18 | + | ||
19 | + cv::Mat model; | ||
20 | + cv::cvtColor(loader.asMat("/scene_model/CAD", CV_8UC1), model, | ||
21 | + cv::COLOR_GRAY2BGR); | ||
22 | + | ||
23 | + std::vector<Contours> pfcContours; | ||
24 | + segmenter.contour(pfcContours, hostPfcMask); | ||
25 | + | ||
26 | + surfaceMap = std::make_unique<const SurfaceMap>(frameSize); | ||
27 | + medianFilter = createMedianFilter(3); | ||
28 | + | ||
29 | + overheating = std::make_unique<ManagedMat>(frameSize, CV_8UC1); | ||
30 | + | ||
31 | + deviceFrame.create(frameSize, CV_16UC1); | ||
32 | + | ||
33 | + medianFilterInput = std::make_unique<ContinuousGpuMat>(frameSize, CV_16UC1); | ||
34 | + medianFilterOutput = | ||
35 | + std::make_unique<ContinuousGpuMat>(frameSize, CV_16UC1); | ||
36 | + | ||
37 | + deviceTemperature = medianFilterInput->device(); | ||
38 | + deviceInput = medianFilterOutput->device(); | ||
39 | +} | ||
40 | + | ||
41 | +void OverloadHotspotDetectionAlgorithm::handleFrame(const cv::Mat &hostFrame, | ||
42 | + unsigned long timestamp) { | ||
43 | + | ||
44 | + /* Overlap data transfer and computations */ | ||
45 | + concurrent.run( | ||
46 | + [this, &hostFrame](const auto &rows, const auto &cols, auto &stream) { | ||
47 | + deviceFrame(rows, cols).upload(hostFrame(rows, cols), stream); | ||
48 | + }, | ||
49 | + currentFrameSize); | ||
50 | + | ||
51 | + concurrent.run( | ||
52 | + [this](const auto &rows, const auto &cols, auto &stream) { | ||
53 | + /* Apply field-of-view mask */ | ||
54 | + cv::cuda::bitwise_and( | ||
55 | + deviceFrame(rows, cols), deviceFrame(rows, cols), | ||
56 | + deviceInput(rows, cols), deviceFov(rows, cols), stream); | ||
57 | + /* Convert Kelvin to Celsius */ | ||
58 | + cv::cuda::subtract(deviceInput(rows, cols), KELVINS, | ||
59 | + deviceTemperature(rows, cols), cv::noArray(), -1, | ||
60 | + stream); | ||
61 | + }, | ||
62 | + currentFrameSize); | ||
63 | + | ||
64 | + concurrent.synchronize(); | ||
65 | + | ||
66 | + /* Apply median filter */ | ||
67 | + medianFilter->apply(deviceTemperature, deviceInput); | ||
68 | + | ||
69 | + /* Threshold overheating pixels */ | ||
70 | + cv::cuda::compare(deviceInput, devicePfc, overheating->device(), | ||
71 | + cv::CMP_GT); | ||
72 | + currentHotspots.clear(); | ||
73 | + identifyHotspots(timestamp); | ||
74 | +} | ||
75 | + | ||
76 | +void OverloadHotspotDetectionAlgorithm::identifyHotspots( | ||
77 | + unsigned long timestamp) { | ||
78 | + /* Analyse topologial structure */ | ||
79 | + Contours hotspotContours; | ||
80 | + cv::findContours(overheating->host(), hotspotContours, cv::RETR_TREE, | ||
81 | + cv::CHAIN_APPROX_SIMPLE); | ||
82 | + | ||
83 | + for (std::size_t i = 0, count = 0, size = hotspotContours.size(); | ||
84 | + i < size && count < blobsLimit; ++i) { | ||
85 | + const auto &contour = hotspotContours[i]; | ||
86 | + /* Discard blobs below minimum area */ | ||
87 | + if (contour.size() >= pixelsThreshold) { | ||
88 | + count += 1; | ||
89 | + | ||
90 | + const int component = hostPfcMask.at<uchar>(contour[0]); | ||
91 | + assert(component > 0); | ||
92 | + | ||
93 | + Hotspot hotspot(deviceTemperature, timestamp, component, contour, | ||
94 | + surfaceMap->at(contour)); | ||
95 | + | ||
96 | + /* Find corresponding blobs */ | ||
97 | + matchHotspot(hotspot); | ||
98 | + currentHotspots.emplace_back(std::move(hotspot)); | ||
99 | + } | ||
100 | + } | ||
101 | +} | ||
102 | + | ||
103 | +void OverloadHotspotDetectionAlgorithm::matchHotspot(const Hotspot &hotspot) { | ||
104 | + for (auto &persistentHotspot : uniqueHotspots) { | ||
105 | + if (clusterCorrespondence.corresponds(persistentHotspot.mask, | ||
106 | + hotspot.mask)) { | ||
107 | + /* Merge matching blobs */ | ||
108 | + persistentHotspot.merge(hotspot); | ||
109 | + return; | ||
110 | + } | ||
111 | + } | ||
112 | + uniqueHotspots.emplace_back(hotspot); | ||
113 | +} | ||
114 | + | ||
115 | +std::unique_ptr<cv::cuda::Filter> | ||
116 | +OverloadHotspotDetectionAlgorithm::createMedianFilter(int kernel) const { | ||
117 | + if (kernel == 3) { | ||
118 | + /* Implementation optimised for the small kernel size */ | ||
119 | + return std::make_unique<MedianFilter3x3<ushort>>(currentFrameSize); | ||
120 | + } | ||
121 | + return CudaFilter::createMedianFilter16U(currentFrameSize, kernel); | ||
122 | +} |
src/algorithm/gpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/cudaarithm.hpp> | ||
4 | +#include <opencv2/cudafilters.hpp> | ||
5 | + | ||
6 | +#include "algorithm/Algorithm.hpp" | ||
7 | + | ||
8 | +#include "algorithm/common/PfcSegmenter.hpp" | ||
9 | +#include "algorithm/common/SurfaceMap.hpp" | ||
10 | + | ||
11 | +#include "algorithm/gpu/common/ClusterCorrespondence.hpp" | ||
12 | +#include "algorithm/gpu/common/ConcurrentCopyExecute.hpp" | ||
13 | +#include "algorithm/gpu/hotspot/Hotspot.hpp" | ||
14 | + | ||
15 | +#include "algorithm/gpu/common/ContinuousGpuMat.hpp" | ||
16 | +#include "algorithm/gpu/common/ManagedMat.hpp" | ||
17 | + | ||
18 | +class OverloadHotspotDetectionAlgorithm : public Algorithm { | ||
19 | + public: | ||
20 | + /* Configurable parameters */ | ||
21 | + std::unique_ptr<cv::cuda::Filter> medianFilter; | ||
22 | + | ||
23 | + std::size_t blobsLimit{5}; | ||
24 | + std::size_t pixelsThreshold{3}; | ||
25 | + | ||
26 | + std::unique_ptr<cv::cuda::Filter> createMedianFilter(int kernel) const; | ||
27 | + /* Configurable parameters */ | ||
28 | + | ||
29 | + OverloadHotspotDetectionAlgorithm() = default; | ||
30 | + | ||
31 | + void setup(const cv::Size &frameSize, const CVMatLoader &loader) override; | ||
32 | + void handleFrame(const cv::Mat &hostFrame, | ||
33 | + unsigned long timestamp) override; | ||
34 | + | ||
35 | + std::vector<PersistentHotspot> uniqueHotspots; /* Set of persistent hotspots */ | ||
36 | + | ||
37 | + private: | ||
38 | + using Contours = std::vector<std::vector<cv::Point>>; | ||
39 | + | ||
40 | + static constexpr int KELVINS{273}; | ||
41 | + static constexpr int MAX_TEMPERATURE{(1L << 15) - 1}; | ||
42 | + | ||
43 | + ConcurrentCopyExecute<4> concurrent; | ||
44 | + cv::Size currentFrameSize; | ||
45 | + | ||
46 | + ClusterCorrespondence clusterCorrespondence{.8, 2.}; | ||
47 | + std::unique_ptr<const SurfaceMap> surfaceMap; | ||
48 | + | ||
49 | + cv::cuda::GpuMat deviceFov, deviceModel, devicePfc; | ||
50 | + cv::Mat hostPfcMask; | ||
51 | + | ||
52 | + std::vector<Hotspot> currentHotspots; | ||
53 | + | ||
54 | + void identifyHotspots(unsigned long timestamp); | ||
55 | + void matchHotspot(const Hotspot &hotspot); | ||
56 | + | ||
57 | + /* Buffors */ | ||
58 | + std::unique_ptr<ContinuousGpuMat> medianFilterInput, medianFilterOutput; | ||
59 | + std::unique_ptr<ManagedMat> overheating; | ||
60 | + cv::cuda::GpuMat deviceFrame, deviceInput, deviceTemperature; | ||
61 | + /* Buffors */ | ||
62 | +}; |
src/algorithm/gpu/surfacelayer/SurfaceLayerDetectionAlgorithm.cpp
0 → 100644
1 | +++ a/src/algorithm/gpu/surfacelayer/SurfaceLayerDetectionAlgorithm.cpp | ||
1 | +#include "SurfaceLayerDetectionAlgorithm.hpp" | ||
2 | + | ||
3 | +#include "algorithm/common/PfcSegmenter.hpp" | ||
4 | + | ||
5 | +#include "algorithm/gpu/common/MedianFilter.hpp" | ||
6 | +#include "algorithm/gpu/common/MedianFilter3x3.hpp" | ||
7 | + | ||
8 | +cv::cuda::GpuMat | ||
9 | +SurfaceLayerDetectionAlgorithm::setMinimumTemperature(ushort temperature) { | ||
10 | + deviceMinTemperature.setTo(temperature); | ||
11 | + return deviceMinTemperature; | ||
12 | +} | ||
13 | + | ||
14 | +void SurfaceLayerDetectionAlgorithm::setup(const cv::Size &frameSize, | ||
15 | + const CVMatLoader &loader) { | ||
16 | + currentFrameSize = frameSize; | ||
17 | + medianFilter = createMedianFilter(3); | ||
18 | + | ||
19 | + deviceRegionMask.upload(loader.asMat("/scene_model/FOV", CV_8UC1) & | ||
20 | + PfcSegmenter(loader, frameSize).mask("Divertors")); | ||
21 | + | ||
22 | + medianFilterInput = std::make_unique<ContinuousGpuMat>(frameSize, CV_16UC1); | ||
23 | + medianFilterOutput = | ||
24 | + std::make_unique<ContinuousGpuMat>(frameSize, CV_16UC1); | ||
25 | + | ||
26 | + deviceFrame = medianFilterInput->device(); | ||
27 | + deviceFiltered = medianFilterOutput->device(); | ||
28 | + | ||
29 | + deviceMinTemperature.create(frameSize, CV_16UC1); | ||
30 | + setMinimumTemperature(434); | ||
31 | + | ||
32 | + deviceFloatInput.create(frameSize, CV_32F); | ||
33 | + deviceDifference.create(frameSize, CV_32F); | ||
34 | + deviceDifference.setTo(cv::Scalar::all(0)); | ||
35 | + | ||
36 | + deviceDifferenceAbs.create(frameSize, CV_32F); | ||
37 | + deviceDifferenceAbs.setTo(cv::Scalar::all(0)); | ||
38 | + | ||
39 | + deviceSurfaceLayers.create(frameSize, CV_8UC1); | ||
40 | + deviceSurfaceLayers.setTo(cv::Scalar::all(0)); | ||
41 | +} | ||
42 | + | ||
43 | +void SurfaceLayerDetectionAlgorithm::handleFrame(const cv::Mat &hostFrame, | ||
44 | + unsigned long timestamp) { | ||
45 | + currentTimestamp = timestamp; | ||
46 | + deviceFrame.upload(hostFrame); | ||
47 | + | ||
48 | + /* Apply median filter */ | ||
49 | + medianFilter->apply(deviceFrame, deviceFiltered); | ||
50 | + /* Apply detection PFC mask */ | ||
51 | + cv::cuda::bitwise_and(deviceFiltered, deviceFiltered, deviceInput, | ||
52 | + deviceRegionMask); | ||
53 | + | ||
54 | + /* Trim noise at minimum temperature */ | ||
55 | + cv::cuda::max(deviceInput, deviceMinTemperature, deviceInput); | ||
56 | + | ||
57 | + /* Convert to float */ | ||
58 | + deviceInput.convertTo(deviceFloatInput, CV_32F); | ||
59 | + /* Calculate normalized derivative of the temperature evolution */ | ||
60 | + normalizedDerivative(); /* Result in deviceDifference & deviceDifferenceAbs */ | ||
61 | + | ||
62 | + /* Threshold rapidly evolving pixels */ | ||
63 | + cv::cuda::compare(deviceDifferenceAbs, derivativeThreshold, deviceResult, | ||
64 | + cv::CMP_GE); | ||
65 | + /* Merge surface layers */ | ||
66 | + cv::cuda::bitwise_or(deviceResult, deviceSurfaceLayers, | ||
67 | + deviceSurfaceLayers); | ||
68 | +} | ||
69 | + | ||
70 | +void SurfaceLayerDetectionAlgorithm::normalizedDerivative() { | ||
71 | + if (devicePreviousPreviousFrame.empty()) { | ||
72 | + previousPreviousTimestamp = currentTimestamp; | ||
73 | + deviceFloatInput.copyTo(devicePreviousPreviousFrame); | ||
74 | + return; | ||
75 | + } | ||
76 | + if (devicePreviousFrame.empty()) { | ||
77 | + previousTimestamp = currentTimestamp; | ||
78 | + deviceFloatInput.copyTo(devicePreviousFrame); | ||
79 | + return; | ||
80 | + } | ||
81 | + | ||
82 | + cv::cuda::subtract(deviceFloatInput, devicePreviousPreviousFrame, | ||
83 | + deviceDifference); | ||
84 | + cv::cuda::multiply( | ||
85 | + deviceDifference, | ||
86 | + cv::Scalar::all((currentTimestamp - previousPreviousTimestamp) * | ||
87 | + NS_TO_S), | ||
88 | + deviceDifference); | ||
89 | + cv::cuda::divide(deviceDifference, devicePreviousFrame, deviceDifference); | ||
90 | + cv::cuda::abs(deviceDifference, deviceDifferenceAbs); | ||
91 | + | ||
92 | + previousPreviousTimestamp = previousTimestamp; | ||
93 | + previousTimestamp = currentTimestamp; | ||
94 | + | ||
95 | + devicePreviousFrame.copyTo(devicePreviousPreviousFrame); | ||
96 | + deviceFloatInput.copyTo(devicePreviousFrame); | ||
97 | +} | ||
98 | + | ||
99 | +std::unique_ptr<cv::cuda::Filter> | ||
100 | +SurfaceLayerDetectionAlgorithm::createMedianFilter(int kernel) const { | ||
101 | + if (kernel == 3) { | ||
102 | + /* Implementation optimised for the small kernel size */ | ||
103 | + return std::make_unique<MedianFilter3x3<ushort>>(currentFrameSize); | ||
104 | + } | ||
105 | + return CudaFilter::createMedianFilter16U(currentFrameSize, kernel); | ||
106 | +} | ||
0 | \ No newline at end of file | 107 | \ No newline at end of file |
src/algorithm/gpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp
0 → 100644
1 | +++ a/src/algorithm/gpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/cudaarithm.hpp> | ||
4 | +#include <opencv2/cudafilters.hpp> | ||
5 | + | ||
6 | +#include "algorithm/Algorithm.hpp" | ||
7 | +#include "algorithm/gpu/common/ContinuousGpuMat.hpp" | ||
8 | + | ||
9 | +class SurfaceLayerDetectionAlgorithm : public Algorithm { | ||
10 | + public: | ||
11 | + /* Configurable parameters */ | ||
12 | + std::unique_ptr<cv::cuda::Filter> medianFilter; | ||
13 | + cv::Scalar derivativeThreshold{7}; | ||
14 | + cv::cuda::GpuMat deviceMinTemperature; | ||
15 | + | ||
16 | + std::unique_ptr<cv::cuda::Filter> createMedianFilter(int kernel) const; | ||
17 | + cv::cuda::GpuMat setMinimumTemperature(ushort temperature); | ||
18 | + /* Configurable parameters */ | ||
19 | + | ||
20 | + SurfaceLayerDetectionAlgorithm() = default; | ||
21 | + | ||
22 | + void setup(const cv::Size &frameSize, const CVMatLoader &loader) override; | ||
23 | + void handleFrame(const cv::Mat &hostFrame, | ||
24 | + unsigned long timestamp) override; | ||
25 | + | ||
26 | + cv::cuda::GpuMat deviceSurfaceLayers; /* Surface layers mask */ | ||
27 | + | ||
28 | + private: | ||
29 | + /* ns to s and invert to multiply instead of divide */ | ||
30 | + static constexpr double NS_TO_S{1e-5}; | ||
31 | + | ||
32 | + cv::Size currentFrameSize; | ||
33 | + | ||
34 | + cv::cuda::GpuMat deviceRegionMask; | ||
35 | + | ||
36 | + cv::cuda::GpuMat devicePreviousPreviousFrame, devicePreviousFrame; | ||
37 | + unsigned long previousPreviousTimestamp{}, previousTimestamp{}, | ||
38 | + currentTimestamp{}; | ||
39 | + | ||
40 | + void normalizedDerivative(); | ||
41 | + | ||
42 | + /* Buffors */ | ||
43 | + std::unique_ptr<ContinuousGpuMat> medianFilterInput, medianFilterOutput; | ||
44 | + cv::cuda::GpuMat deviceFrame, deviceFiltered, deviceInput, deviceFloatInput, | ||
45 | + deviceDifference, deviceDifferenceAbs, deviceResult; | ||
46 | + /* Buffors */ | ||
47 | +}; |
src/benchmark/benchmark.hpp
0 → 100644
1 | +++ a/src/benchmark/benchmark.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <iostream> | ||
4 | +#include <memory> | ||
5 | + | ||
6 | +#include <opencv2/core.hpp> | ||
7 | + | ||
8 | +#include "data/hdf5/CVMatLoader.hpp" | ||
9 | +#include "data/provider/H5FrameProvider.hpp" | ||
10 | + | ||
11 | +namespace Benchmark { | ||
12 | +std::unique_ptr<H5FrameProvider> | ||
13 | +createFrameProvider(const std::string &filePath = "20171114.053_AEF20.h5") { | ||
14 | + const auto timestamps = | ||
15 | + CVMatLoader(ResourceLocator::getPathProvider().path(filePath)) | ||
16 | + .asVec<unsigned long>("/timestamps", H5SingleDataReader::UINT64); | ||
17 | + return std::make_unique<H5FrameProvider>( | ||
18 | + ResourceLocator::getPathProvider().path(filePath), "/images", | ||
19 | + timestamps); | ||
20 | +} | ||
21 | + | ||
22 | +void displayStats(const std::vector<double> &timings) { | ||
23 | + cv::Scalar mean, std; | ||
24 | + cv::meanStdDev(timings, mean, std); | ||
25 | + | ||
26 | + double min{-1}, max{-1}; | ||
27 | + cv::minMaxIdx(timings, &min, &max, nullptr, nullptr); | ||
28 | + | ||
29 | + std::cout << "Mean: " << mean[0] << " Std Dev.: " << std[0] | ||
30 | + << " Min: " << min << " Max: " << max << '\n'; | ||
31 | +} | ||
32 | +} // namespace Benchmark | ||
0 | \ No newline at end of file | 33 | \ No newline at end of file |
src/benchmark/median_filter_benchmark.hpp
0 → 100644
1 | +++ a/src/benchmark/median_filter_benchmark.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include "benchmark.hpp" | ||
4 | + | ||
5 | +#include <opencv2/cudaarithm.hpp> | ||
6 | +#include <opencv2/cudafilters.hpp> | ||
7 | +#include <opencv2/imgcodecs.hpp> | ||
8 | + | ||
9 | +#include "algorithm/gpu/common/ContinuousGpuMat.hpp" | ||
10 | +#include "algorithm/gpu/common/MedianFilter.hpp" | ||
11 | +#include "algorithm/gpu/common/MedianFilter3x3.hpp" | ||
12 | + | ||
13 | +namespace MedianFilterBenchmark { | ||
14 | + | ||
15 | +template <std::size_t iters> | ||
16 | +std::vector<double> benchmarkFilter(cv::cuda::Filter &filter, | ||
17 | + const cv::cuda::GpuMat input, | ||
18 | + const std::string &name) { | ||
19 | + std::vector<double> timings; | ||
20 | + timings.reserve(iters); | ||
21 | + | ||
22 | + cv::cuda::GpuMat output(input.size(), input.type()); | ||
23 | + /* Warm-up runs */ | ||
24 | + filter.apply(input, output); | ||
25 | + filter.apply(input, output); | ||
26 | + filter.apply(input, output); | ||
27 | + | ||
28 | + for (std::size_t i = 0; i < iters; ++i) { | ||
29 | + const auto start = static_cast<double>(cv::getTickCount()); | ||
30 | + filter.apply(input, output); | ||
31 | + const auto timeMs = 1e3 * | ||
32 | + (static_cast<double>(cv::getTickCount()) - start) / | ||
33 | + cv::getTickFrequency(); | ||
34 | + timings.push_back(timeMs); | ||
35 | + } | ||
36 | + | ||
37 | + cv::imwrite(name + ".bmp", cv::Mat(output)); | ||
38 | + return timings; | ||
39 | +} | ||
40 | + | ||
41 | +void run() { | ||
42 | + constexpr auto iterations{100}; | ||
43 | + | ||
44 | + cv::Mat frame(cv::Size(1024, 768), CV_16UC1); | ||
45 | + Benchmark::createFrameProvider()->next(frame.data); | ||
46 | + | ||
47 | + ContinuousGpuMat deviceFrame8bit(cv::Size(1024, 768), CV_16UC1); | ||
48 | + cv::cuda::normalize(cv::cuda::GpuMat(frame), deviceFrame8bit.device(), 0, | ||
49 | + 255, cv::NORM_MINMAX, CV_8UC1); | ||
50 | + | ||
51 | + { | ||
52 | + auto opencvFilter = cv::cuda::createMedianFilter(CV_8UC1, 3); | ||
53 | + std::cout << "OpenCV 3x3 8-bit\n:"; | ||
54 | + Benchmark::displayStats(benchmarkFilter<iterations>( | ||
55 | + *opencvFilter, deviceFrame8bit.device(), "opencv")); | ||
56 | + } | ||
57 | + | ||
58 | + { | ||
59 | + auto nppFilter = CudaFilter::createMedianFilter8U(frame.size(), 3); | ||
60 | + std::cout << "NPP 3x3 8-bit\n:"; | ||
61 | + Benchmark::displayStats(benchmarkFilter<iterations>( | ||
62 | + *nppFilter, deviceFrame8bit.device(), "npp")); | ||
63 | + } | ||
64 | + | ||
65 | + { | ||
66 | + auto cudaKernelFilter = | ||
67 | + std::make_unique<MedianFilter3x3<uchar>>(frame.size()); | ||
68 | + std::cout << "CUDA kernel 3x3 8-bit\n:"; | ||
69 | + Benchmark::displayStats(benchmarkFilter<iterations>( | ||
70 | + *cudaKernelFilter, deviceFrame8bit.device(), "cuda_kernel")); | ||
71 | + } | ||
72 | +} | ||
73 | +} // namespace MedianFilterBenchmark | ||
0 | \ No newline at end of file | 74 | \ No newline at end of file |
src/benchmark/thermal_event_detection_benchmark.hpp
0 → 100644
1 | +++ a/src/benchmark/thermal_event_detection_benchmark.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include "benchmark.hpp" | ||
4 | + | ||
5 | +#include "data/ResourceLocator.hpp" | ||
6 | + | ||
7 | +#include "algorithm/gpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp" | ||
8 | +#include "algorithm/gpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp" | ||
9 | +/* CPU implementations are in Cpu:: namespace */ | ||
10 | +#include "algorithm/cpu/hotspot/OverloadHotspotDetectionAlgorithm.hpp" | ||
11 | +#include "algorithm/cpu/surfacelayer/SurfaceLayerDetectionAlgorithm.hpp" | ||
12 | + | ||
13 | +namespace ThermalEventDetectionBenchmark { | ||
14 | + | ||
15 | +template <typename T, std::size_t firstFrame, std::size_t lastFrame> | ||
16 | +std::vector<double> benchmarkAlgorithm(cv::Mat frame) { | ||
17 | + std::vector<double> timings; | ||
18 | + timings.reserve(lastFrame - firstFrame + 1); | ||
19 | + | ||
20 | + const CVMatLoader loader( | ||
21 | + ResourceLocator::getPathProvider().path("20171114.053_AEF20.h5")); | ||
22 | + | ||
23 | + T algorithm{}; | ||
24 | + algorithm.setup(frame.size(), loader); | ||
25 | + | ||
26 | + auto provider = Benchmark::createFrameProvider(); | ||
27 | + for (std::size_t i = 0; provider->hasNext() && i <= lastFrame; ++i) { | ||
28 | + const auto timestamp = provider->next(frame.data); | ||
29 | + const auto start = static_cast<double>(cv::getTickCount()); | ||
30 | + algorithm.handleFrame(frame, timestamp); | ||
31 | + const auto timeMs = 1e3 * | ||
32 | + (static_cast<double>(cv::getTickCount()) - start) / | ||
33 | + cv::getTickFrequency(); | ||
34 | + if (i >= firstFrame) { | ||
35 | + timings.push_back(timeMs); | ||
36 | + } | ||
37 | + } | ||
38 | + | ||
39 | + return timings; | ||
40 | +} | ||
41 | + | ||
42 | +template <typename T, std::size_t firstFrame, std::size_t lastFrame, | ||
43 | + std::size_t iters> | ||
44 | +std::vector<double> benchmarkAlgorithmN(cv::Mat frame) { | ||
45 | + std::vector<double> timings; | ||
46 | + timings.reserve(iters); | ||
47 | + | ||
48 | + for (std::size_t i = 0; i < iters; ++i) { | ||
49 | + timings.push_back( | ||
50 | + cv::mean(benchmarkAlgorithm<T, firstFrame, lastFrame>(frame))[0]); | ||
51 | + } | ||
52 | + | ||
53 | + return timings; | ||
54 | +} | ||
55 | + | ||
56 | +template <typename T, std::size_t firstFrame, std::size_t lastFrame, | ||
57 | + std::size_t iters> | ||
58 | +std::vector<double> benchmarkGpu() { | ||
59 | + /* Page-locked memory */ | ||
60 | + void *data{nullptr}; | ||
61 | + cudaSetDeviceFlags(cudaDeviceMapHost); | ||
62 | + cudaHostAlloc(&data, 1024 * 768 * cv::getElemSize(CV_16UC1), | ||
63 | + cudaHostAllocMapped); | ||
64 | + const cv::Mat frame(cv::Size(1024, 768), CV_16UC1, data); | ||
65 | + | ||
66 | + auto timings = benchmarkAlgorithmN<T, firstFrame, lastFrame, iters>(frame); | ||
67 | + cudaFreeHost(data); | ||
68 | + | ||
69 | + return timings; | ||
70 | +} | ||
71 | + | ||
72 | +template <typename T, std::size_t firstFrame, std::size_t lastFrame, | ||
73 | + std::size_t iters> | ||
74 | +std::vector<double> benchmarkCpu() { | ||
75 | + const cv::Mat frame(cv::Size(1024, 768), CV_16UC1); | ||
76 | + return benchmarkAlgorithmN<T, firstFrame, lastFrame, iters>(frame); | ||
77 | +} | ||
78 | + | ||
79 | +void run() { | ||
80 | + constexpr auto iterations{20}; | ||
81 | + | ||
82 | + std::cout << "GPU Overload Hotspot Detection Algorithm\n"; | ||
83 | + Benchmark::displayStats(benchmarkGpu<OverloadHotspotDetectionAlgorithm, 239, | ||
84 | + 258, iterations>()); | ||
85 | + std::cout << "GPU Surface Layer Detection Algorithm\n"; | ||
86 | + Benchmark::displayStats( | ||
87 | + benchmarkGpu<SurfaceLayerDetectionAlgorithm, 16, 41, iterations>()); | ||
88 | + std::cout << "CPU Overload Hotspot Detection Algorithm\n"; | ||
89 | + Benchmark::displayStats(benchmarkCpu<Cpu::OverloadHotspotDetectionAlgorithm, | ||
90 | + 239, 258, iterations>()); | ||
91 | + std::cout << "CPU Surface Layer Detection Algorithm\n"; | ||
92 | + Benchmark::displayStats(benchmarkCpu<Cpu::SurfaceLayerDetectionAlgorithm, | ||
93 | + 16, 41, iterations>()); | ||
94 | +} | ||
95 | + | ||
96 | +} // namespace ThermalEventDetectionBenchmark | ||
0 | \ No newline at end of file | 97 | \ No newline at end of file |
src/data/DataNotFoundException.hpp
0 → 100644
1 | +++ a/src/data/DataNotFoundException.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <exception> | ||
4 | +#include <string> | ||
5 | + | ||
6 | +class DataNotFoundException : public std::runtime_error { | ||
7 | + public: | ||
8 | + explicit DataNotFoundException(const std::string &what) | ||
9 | + : std::runtime_error(what) {} | ||
10 | +}; | ||
0 | \ No newline at end of file | 11 | \ No newline at end of file |
src/data/ResourceLocator.cpp
0 → 100644
src/data/ResourceLocator.hpp
0 → 100644
1 | +++ a/src/data/ResourceLocator.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <exception> | ||
4 | +#include <memory> | ||
5 | +#include <string> | ||
6 | + | ||
7 | +class PathProvider { | ||
8 | + public: | ||
9 | + PathProvider() = default; | ||
10 | + PathProvider(const PathProvider &) = default; | ||
11 | + PathProvider(PathProvider &&) noexcept = default; | ||
12 | + PathProvider &operator=(PathProvider &&) noexcept = default; | ||
13 | + PathProvider &operator=(const PathProvider &) = default; | ||
14 | + virtual ~PathProvider() = default; | ||
15 | + | ||
16 | + virtual std::string path(const std::string &file) const = 0; | ||
17 | +}; | ||
18 | + | ||
19 | +class AbsolutePathProvider : public PathProvider { | ||
20 | + public: | ||
21 | + explicit AbsolutePathProvider(const std::string &relativePath) | ||
22 | + : relativePath(relativePath) {} | ||
23 | + | ||
24 | + std::string path(const std::string &file) const override { | ||
25 | + return relativePath + file; | ||
26 | + } | ||
27 | + | ||
28 | + private: | ||
29 | + const std::string relativePath; | ||
30 | +}; | ||
31 | + | ||
32 | +class ResourceLocator { | ||
33 | + public: | ||
34 | + static const PathProvider &getPathProvider() { | ||
35 | + if (!pathProvider) { | ||
36 | + throw std::runtime_error("Path provider not provided"); | ||
37 | + } | ||
38 | + return *pathProvider; | ||
39 | + } | ||
40 | + | ||
41 | + static void setPathProvider(std::unique_ptr<PathProvider> provider) { | ||
42 | + pathProvider = std::move(provider); | ||
43 | + } | ||
44 | + | ||
45 | + private: | ||
46 | + static std::unique_ptr<PathProvider> pathProvider; | ||
47 | +}; | ||
0 | \ No newline at end of file | 48 | \ No newline at end of file |
src/data/hdf5/CVMatLoader.hpp
0 → 100644
1 | +++ a/src/data/hdf5/CVMatLoader.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <opencv2/core/mat.hpp> | ||
4 | + | ||
5 | +#include "data/DataNotFoundException.hpp" | ||
6 | +#include "data/hdf5/H5SingleDataReader.hpp" | ||
7 | + | ||
8 | +class CVMatLoader { | ||
9 | + public: | ||
10 | + explicit CVMatLoader(const std::string &path) : path(path) {} | ||
11 | + | ||
12 | + cv::Mat asMat(const std::string &sourceDataSet, int destinationType) const { | ||
13 | + const auto reader = | ||
14 | + createReader(sourceDataSet, resolveType(destinationType), 2); | ||
15 | + | ||
16 | + cv::Mat output(cv::Size{static_cast<int>(reader.dimensions[1]), | ||
17 | + static_cast<int>(reader.dimensions[0])}, | ||
18 | + destinationType); | ||
19 | + assert(output.isContinuous()); | ||
20 | + | ||
21 | + reader.read(output.data); | ||
22 | + return output; | ||
23 | + } | ||
24 | + | ||
25 | + template <typename T> | ||
26 | + std::vector<T> asVec(const std::string &sourceDataSet, | ||
27 | + const H5::PredType &sourceType) const { | ||
28 | + const auto reader = createReader(sourceDataSet, sourceType, 1); | ||
29 | + | ||
30 | + std::vector<T> output(reader.size()); | ||
31 | + reader.read(output.data()); | ||
32 | + | ||
33 | + return output; | ||
34 | + } | ||
35 | + | ||
36 | + private: | ||
37 | + const std::string path; | ||
38 | + | ||
39 | + H5SingleDataReader createReader(const std::string &sourceDataSet, | ||
40 | + const H5::PredType &sourceType, | ||
41 | + int dimensions) const { | ||
42 | + try { | ||
43 | + return {path, sourceDataSet, sourceType, dimensions}; | ||
44 | + } catch (const H5::FileIException &exception) { | ||
45 | + throw DataNotFoundException(exception.getDetailMsg()); | ||
46 | + } | ||
47 | + } | ||
48 | + | ||
49 | + static const H5::PredType &resolveType(int cvType) { | ||
50 | + switch (CV_MAT_DEPTH(cvType)) { | ||
51 | + case CV_8U: | ||
52 | + return H5SingleDataReader::UINT8; | ||
53 | + case CV_16U: | ||
54 | + return H5SingleDataReader::UINT16; | ||
55 | + case CV_8S: | ||
56 | + return H5SingleDataReader::INT8; | ||
57 | + case CV_16S: | ||
58 | + return H5SingleDataReader::INT16; | ||
59 | + case CV_32S: | ||
60 | + return H5SingleDataReader::INT32; | ||
61 | + case CV_32F: | ||
62 | + return H5SingleDataReader::FLOAT; | ||
63 | + case CV_64F: | ||
64 | + return H5SingleDataReader::DOUBLE; | ||
65 | + default: | ||
66 | + throw std::runtime_error("Unknown OpenCV type " + | ||
67 | + std::to_string(cvType)); | ||
68 | + } | ||
69 | + } | ||
70 | +}; |
src/data/hdf5/H5SingleDataReader.cpp
0 → 100644
1 | +++ a/src/data/hdf5/H5SingleDataReader.cpp | ||
1 | +#include "H5SingleDataReader.hpp" | ||
2 | + | ||
3 | +const H5::PredType H5SingleDataReader::UINT8 = H5::PredType::NATIVE_UINT8; | ||
4 | +const H5::PredType H5SingleDataReader::UINT16 = H5::PredType::NATIVE_UINT16; | ||
5 | +const H5::PredType H5SingleDataReader::UINT32 = H5::PredType::NATIVE_UINT32; | ||
6 | +const H5::PredType H5SingleDataReader::UINT64 = H5::PredType::NATIVE_UINT64; | ||
7 | +const H5::PredType H5SingleDataReader::INT8 = H5::PredType::NATIVE_INT8; | ||
8 | +const H5::PredType H5SingleDataReader::INT16 = H5::PredType::NATIVE_INT16; | ||
9 | +const H5::PredType H5SingleDataReader::INT32 = H5::PredType::NATIVE_INT32; | ||
10 | +const H5::PredType H5SingleDataReader::INT64 = H5::PredType::NATIVE_INT64; | ||
11 | +const H5::PredType H5SingleDataReader::FLOAT = H5::PredType::NATIVE_FLOAT; | ||
12 | +const H5::PredType H5SingleDataReader::DOUBLE = H5::PredType::NATIVE_DOUBLE; | ||
13 | + | ||
14 | +std::mutex H5SingleDataReader::lock{}; | ||
0 | \ No newline at end of file | 15 | \ No newline at end of file |
src/data/hdf5/H5SingleDataReader.hpp
0 → 100644
1 | +++ a/src/data/hdf5/H5SingleDataReader.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include <algorithm> | ||
4 | +#include <array> | ||
5 | +#include <cassert> | ||
6 | +#include <memory> | ||
7 | +#include <mutex> | ||
8 | +#include <numeric> | ||
9 | +#include <string> | ||
10 | + | ||
11 | +#include <H5Cpp.h> | ||
12 | + | ||
13 | +/* | ||
14 | + * HDF5 C++ API doesn't support multiple threads, even instantiation of types | ||
15 | + * since it uses global structres. | ||
16 | + * | ||
17 | + * Thus, to avoid deafult construction of member H5 types, they are stored using | ||
18 | + * not owning pointers. | ||
19 | + */ | ||
20 | +class H5SingleDataReader { | ||
21 | + public: | ||
22 | + H5SingleDataReader(const H5SingleDataReader &) = delete; | ||
23 | + H5SingleDataReader &operator=(const H5SingleDataReader &) = delete; | ||
24 | + | ||
25 | + H5SingleDataReader(H5SingleDataReader &&) noexcept = default; | ||
26 | + H5SingleDataReader &operator=(H5SingleDataReader &&) noexcept = default; | ||
27 | + | ||
28 | + std::array<hsize_t, 3> dimensions{0, 0, 0}; | ||
29 | + | ||
30 | + H5SingleDataReader(const std::string &path, const std::string &dataSetPath, | ||
31 | + const H5::PredType &dataType, int outputDimension) { | ||
32 | + std::lock_guard<std::mutex> guard(lock); | ||
33 | + | ||
34 | + type.reset(assign(dataType)); | ||
35 | + | ||
36 | + const H5::H5File file(path, H5F_ACC_RDONLY); | ||
37 | + dataSet.reset(assign(file.openDataSet(dataSetPath))); | ||
38 | + fileSpace.reset(assign(dataSet->getSpace())); | ||
39 | + | ||
40 | + [[maybe_unused]] const auto rank = | ||
41 | + fileSpace->getSimpleExtentDims(dimensions.data(), nullptr); | ||
42 | + assert(outputDimension <= rank); | ||
43 | + | ||
44 | + const std::array<hsize_t, 1> memoryDimensions{std::accumulate( | ||
45 | + dimensions.cbegin(), dimensions.cbegin() + outputDimension, 1LLU, | ||
46 | + std::multiplies<>())}; | ||
47 | + const std::array<hsize_t, 1> memoryOffset{0}; | ||
48 | + memorySpace.reset(assign(H5::DataSpace(1, memoryDimensions.data()))); | ||
49 | + memorySpace->selectHyperslab(H5S_SELECT_SET, memoryDimensions.data(), | ||
50 | + memoryOffset.data()); | ||
51 | + | ||
52 | + totalSize = memorySpace->getSelectNpoints(); | ||
53 | + } | ||
54 | + | ||
55 | + hsize_t size() const { return totalSize; } | ||
56 | + | ||
57 | + void read(void *buffer) const { | ||
58 | + std::lock_guard<std::mutex> guard(lock); | ||
59 | + dataSet->read(buffer, *type, *memorySpace, *fileSpace); | ||
60 | + } | ||
61 | + | ||
62 | + void select(const std::array<hsize_t, 3> selectionSize, | ||
63 | + const std::array<hsize_t, 3> selectionOffset) const { | ||
64 | + std::lock_guard<std::mutex> guard(lock); | ||
65 | + fileSpace->selectHyperslab(H5S_SELECT_SET, selectionSize.data(), | ||
66 | + selectionOffset.data()); | ||
67 | + } | ||
68 | + | ||
69 | + virtual ~H5SingleDataReader() { | ||
70 | + std::lock_guard<std::mutex> guard(lock); | ||
71 | + delete memorySpace.get(); | ||
72 | + delete fileSpace.get(); | ||
73 | + delete dataSet.get(); | ||
74 | + delete type.get(); | ||
75 | + } | ||
76 | + | ||
77 | + const static H5::PredType UINT8; | ||
78 | + const static H5::PredType UINT16; | ||
79 | + const static H5::PredType UINT32; | ||
80 | + const static H5::PredType UINT64; | ||
81 | + const static H5::PredType INT8; | ||
82 | + const static H5::PredType INT16; | ||
83 | + const static H5::PredType INT32; | ||
84 | + const static H5::PredType INT64; | ||
85 | + const static H5::PredType FLOAT; | ||
86 | + const static H5::PredType DOUBLE; | ||
87 | + | ||
88 | + private: | ||
89 | + static std::mutex lock; | ||
90 | + | ||
91 | + hsize_t totalSize; | ||
92 | + | ||
93 | + struct NoOp { | ||
94 | + template <typename T> void operator()(T const &) const noexcept {} | ||
95 | + }; | ||
96 | + | ||
97 | + /** | ||
98 | + * not_owning_ptr allows to use default move constructor | ||
99 | + * and manually manage target memory | ||
100 | + */ | ||
101 | + template <typename T> using not_owning_ptr = std::unique_ptr<T, NoOp>; | ||
102 | + | ||
103 | + not_owning_ptr<const H5::PredType> type; | ||
104 | + not_owning_ptr<const H5::DataSet> dataSet; | ||
105 | + not_owning_ptr<const H5::DataSpace> fileSpace; | ||
106 | + not_owning_ptr<const H5::DataSpace> memorySpace; | ||
107 | + | ||
108 | + template <typename T> static const T *assign(const T &value) { | ||
109 | + T *const result = new T(value); | ||
110 | + *result = value; | ||
111 | + return result; | ||
112 | + } | ||
113 | +}; |
src/data/provider/FrameProvider.hpp
0 → 100644
1 | +++ a/src/data/provider/FrameProvider.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +class FrameProvider { | ||
4 | + public: | ||
5 | + FrameProvider() = default; | ||
6 | + FrameProvider(const FrameProvider &) = default; | ||
7 | + FrameProvider(FrameProvider &&) noexcept = default; | ||
8 | + FrameProvider &operator=(FrameProvider &&) noexcept = default; | ||
9 | + FrameProvider &operator=(const FrameProvider &) = default; | ||
10 | + | ||
11 | + virtual unsigned long next(void *buffer) = 0; | ||
12 | + | ||
13 | + virtual bool hasNext() const = 0; | ||
14 | + | ||
15 | + virtual size_t rows() const = 0; | ||
16 | + | ||
17 | + virtual size_t columns() const = 0; | ||
18 | + | ||
19 | + virtual ~FrameProvider() = default; | ||
20 | +}; |
src/data/provider/H5FrameProvider.hpp
0 → 100644
1 | +++ a/src/data/provider/H5FrameProvider.hpp | ||
1 | + | ||
2 | +#pragma once | ||
3 | + | ||
4 | +#include <array> | ||
5 | + | ||
6 | +#include "data/hdf5/H5SingleDataReader.hpp" | ||
7 | +#include "data/provider/FrameProvider.hpp" | ||
8 | +#include "data/provider/ProviderOpenFailureException.hpp" | ||
9 | + | ||
10 | +class H5FrameProvider : public FrameProvider, public H5SingleDataReader { | ||
11 | + public: | ||
12 | + H5FrameProvider(const std::string &path, const std::string &dataSetPath, | ||
13 | + const std::vector<unsigned long> ×tamps) try | ||
14 | + : H5SingleDataReader(path, dataSetPath, H5SingleDataReader::UINT16, 2), | ||
15 | + timestamps(timestamps) { | ||
16 | + } catch (const H5::FileIException &exception) { | ||
17 | + throw ProviderOpenFailureException(exception.getDetailMsg()); | ||
18 | + } | ||
19 | + | ||
20 | + H5FrameProvider(const std::string &path, const std::string &dataSetPath) try | ||
21 | + : H5SingleDataReader(path, dataSetPath, H5SingleDataReader::UINT16, 2), | ||
22 | + timestamps(std::vector<unsigned long>(dimensions[2], 0)) { | ||
23 | + } catch (const H5::FileIException &exception) { | ||
24 | + throw ProviderOpenFailureException(exception.getDetailMsg()); | ||
25 | + } | ||
26 | + | ||
27 | + bool hasNext() const override { | ||
28 | + return current < dimensions[2] && current <= lastFrame; | ||
29 | + } | ||
30 | + | ||
31 | + unsigned long next(void *buffer) override { | ||
32 | + selectImage(current); | ||
33 | + read(buffer); | ||
34 | + return timestamps[current++]; | ||
35 | + } | ||
36 | + | ||
37 | + void first(void *buffer) const { | ||
38 | + selectImage(0); | ||
39 | + read(buffer); | ||
40 | + } | ||
41 | + | ||
42 | + size_t rows() const override { return dimensions[0]; } | ||
43 | + | ||
44 | + size_t columns() const override { return dimensions[1]; } | ||
45 | + | ||
46 | + private: | ||
47 | + hsize_t current = 0; // 239 - 5; /* Modify to start from further frame */ | ||
48 | + | ||
49 | + static constexpr hsize_t lastFrame{258}; /* Peak */ | ||
50 | + // static constexpr hsize_t lastFrame{41}; /* Rise */ | ||
51 | + // static constexpr hsize_t lastFrame{9999}; /* Entire */ | ||
52 | + | ||
53 | + const std::vector<unsigned long> timestamps; | ||
54 | + | ||
55 | + void selectImage(hsize_t image) const { | ||
56 | + select({dimensions[0], dimensions[1], 1}, {0, 0, image}); | ||
57 | + } | ||
58 | +}; | ||
0 | \ No newline at end of file | 59 | \ No newline at end of file |
src/data/provider/ProviderOpenFailureException.hpp
0 → 100644
1 | +++ a/src/data/provider/ProviderOpenFailureException.hpp | ||
1 | +#pragma once | ||
2 | + | ||
3 | +#include "data/DataNotFoundException.hpp" | ||
4 | + | ||
5 | +class ProviderOpenFailureException : public DataNotFoundException { | ||
6 | + public: | ||
7 | + explicit ProviderOpenFailureException(const std::string &what) | ||
8 | + : DataNotFoundException(what) {} | ||
9 | +}; | ||
0 | \ No newline at end of file | 10 | \ No newline at end of file |
src/main.cpp
0 → 100644
1 | +++ a/src/main.cpp | ||
1 | +#include "data/ResourceLocator.hpp" | ||
2 | + | ||
3 | +#include "benchmark/median_filter_benchmark.hpp" | ||
4 | +#include "benchmark/thermal_event_detection_benchmark.hpp" | ||
5 | + | ||
6 | +int main(int, char **) { | ||
7 | + ResourceLocator::setPathProvider( | ||
8 | + std::make_unique<AbsolutePathProvider>("resource/")); | ||
9 | + | ||
10 | + ThermalEventDetectionBenchmark::run(); | ||
11 | + MedianFilterBenchmark::run(); | ||
12 | + | ||
13 | + return 0; | ||
14 | +} | ||
0 | \ No newline at end of file | 15 | \ No newline at end of file |