+
+
+
+
+
+
+
+  |
+
+
+
+ $projectname
+ $projectnumber
+
+ $projectbrief
+ |
+
+
+
+
+ $projectbrief
+ |
+
+
+
+
+ $searchbox |
+
+
+
+
+
+
+
+
diff --git a/docs/doxygen-awesome.css b/docs/doxygen-awesome.css
new file mode 100644
index 00000000..121eb0fd
--- /dev/null
+++ b/docs/doxygen-awesome.css
@@ -0,0 +1,2405 @@
+/**
+
+Doxygen Awesome
+https://github.com/jothepro/doxygen-awesome-css
+
+MIT License
+
+Copyright (c) 2021 - 2022 jothepro
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
+
+*/
+
+html {
+ /* primary theme color. This will affect the entire websites color scheme: links, arrows, labels, ... */
+ --primary-color: #1779c4;
+ --primary-dark-color: #335c80;
+ --primary-light-color: #70b1e9;
+
+ /* page base colors */
+ --page-background-color: #ffffff;
+ --page-foreground-color: #2f4153;
+ --page-secondary-foreground-color: #6f7e8e;
+
+ /* color for all separators on the website: hr, borders, ... */
+ --separator-color: #dedede;
+
+ /* border radius for all rounded components. Will affect many components, like dropdowns, memitems, codeblocks, ... */
+ --border-radius-large: 8px;
+ --border-radius-small: 4px;
+ --border-radius-medium: 6px;
+
+ /* default spacings. Most components reference these values for spacing, to provide uniform spacing on the page. */
+ --spacing-small: 5px;
+ --spacing-medium: 10px;
+ --spacing-large: 16px;
+
+ /* default box shadow used for raising an element above the normal content. Used in dropdowns, search result, ... */
+ --box-shadow: 0 2px 8px 0 rgba(0,0,0,.075);
+
+ --odd-color: rgba(0,0,0,.028);
+
+ /* font-families. will affect all text on the website
+ * font-family: the normal font for text, headlines, menus
+ * font-family-monospace: used for preformatted text in memtitle, code, fragments
+ */
+ --font-family: -apple-system,BlinkMacSystemFont,Segoe UI,Roboto,Oxygen,Ubuntu,Cantarell,Fira Sans,Droid Sans,Helvetica Neue,sans-serif;
+ --font-family-monospace: ui-monospace,SFMono-Regular,SF Mono,Menlo,Consolas,Liberation Mono,monospace;
+
+ /* font sizes */
+ --page-font-size: 15.6px;
+ --navigation-font-size: 14.4px;
+ --toc-font-size: 13.4px;
+ --code-font-size: 14px; /* affects code, fragment */
+ --title-font-size: 22px;
+
+ /* content text properties. These only affect the page content, not the navigation or any other ui elements */
+ --content-line-height: 27px;
+ /* The content is centered and constraint in it's width. To make the content fill the whole page, set the variable to auto.*/
+ --content-maxwidth: 1050px;
+ --table-line-height: 24px;
+ --toc-sticky-top: var(--spacing-medium);
+ --toc-width: 200px;
+ --toc-max-height: calc(100vh - 2 * var(--spacing-medium) - 85px);
+
+ /* colors for various content boxes: @warning, @note, @deprecated @bug */
+ --warning-color: #f8d1cc;
+ --warning-color-dark: #b61825;
+ --warning-color-darker: #75070f;
+ --note-color: #faf3d8;
+ --note-color-dark: #f3a600;
+ --note-color-darker: #5f4204;
+ --todo-color: #e4f3ff;
+ --todo-color-dark: #1879C4;
+ --todo-color-darker: #274a5c;
+ --deprecated-color: #ecf0f3;
+ --deprecated-color-dark: #5b6269;
+ --deprecated-color-darker: #43454a;
+ --bug-color: #e4dafd;
+ --bug-color-dark: #5b2bdd;
+ --bug-color-darker: #2a0d72;
+ --invariant-color: #d8f1e3;
+ --invariant-color-dark: #44b86f;
+ --invariant-color-darker: #265532;
+
+ /* blockquote colors */
+ --blockquote-background: #f8f9fa;
+ --blockquote-foreground: #636568;
+
+ /* table colors */
+ --tablehead-background: #f1f1f1;
+ --tablehead-foreground: var(--page-foreground-color);
+
+ /* menu-display: block | none
+ * Visibility of the top navigation on screens >= 768px. On smaller screen the menu is always visible.
+ * `GENERATE_TREEVIEW` MUST be enabled!
+ */
+ --menu-display: block;
+
+ --menu-focus-foreground: var(--page-background-color);
+ --menu-focus-background: var(--primary-color);
+ --menu-selected-background: rgba(0,0,0,.05);
+
+
+ --header-background: var(--page-background-color);
+ --header-foreground: var(--page-foreground-color);
+
+ /* searchbar colors */
+ --searchbar-background: var(--side-nav-background);
+ --searchbar-foreground: var(--page-foreground-color);
+
+ /* searchbar size
+ * (`searchbar-width` is only applied on screens >= 768px.
+ * on smaller screens the searchbar will always fill the entire screen width) */
+ --searchbar-height: 33px;
+ --searchbar-width: 210px;
+ --searchbar-border-radius: var(--searchbar-height);
+
+ /* code block colors */
+ --code-background: #f5f5f5;
+ --code-foreground: var(--page-foreground-color);
+
+ /* fragment colors */
+ --fragment-background: #F8F9FA;
+ --fragment-foreground: #37474F;
+ --fragment-keyword: #bb6bb2;
+ --fragment-keywordtype: #8258b3;
+ --fragment-keywordflow: #d67c3b;
+ --fragment-token: #438a59;
+ --fragment-comment: #969696;
+ --fragment-link: #5383d6;
+ --fragment-preprocessor: #46aaa5;
+ --fragment-linenumber-color: #797979;
+ --fragment-linenumber-background: #f4f4f5;
+ --fragment-linenumber-border: #e3e5e7;
+ --fragment-lineheight: 20px;
+
+ /* sidebar navigation (treeview) colors */
+ --side-nav-background: #fbfbfb;
+ --side-nav-foreground: var(--page-foreground-color);
+ --side-nav-arrow-opacity: 0;
+ --side-nav-arrow-hover-opacity: 0.9;
+
+ --toc-background: var(--side-nav-background);
+ --toc-foreground: var(--side-nav-foreground);
+
+ /* height of an item in any tree / collapsable table */
+ --tree-item-height: 30px;
+
+ --memname-font-size: var(--code-font-size);
+ --memtitle-font-size: 18px;
+
+ --webkit-scrollbar-size: 7px;
+ --webkit-scrollbar-padding: 4px;
+ --webkit-scrollbar-color: var(--separator-color);
+}
+
+@media screen and (max-width: 767px) {
+ html {
+ --page-font-size: 16px;
+ --navigation-font-size: 16px;
+ --toc-font-size: 15px;
+ --code-font-size: 15px; /* affects code, fragment */
+ --title-font-size: 22px;
+ }
+}
+
+@media (prefers-color-scheme: dark) {
+ html:not(.light-mode) {
+ color-scheme: dark;
+
+ --primary-color: #1982d2;
+ --primary-dark-color: #86a9c4;
+ --primary-light-color: #4779ac;
+
+ --box-shadow: 0 2px 8px 0 rgba(0,0,0,.35);
+
+ --odd-color: rgba(100,100,100,.06);
+
+ --menu-selected-background: rgba(0,0,0,.4);
+
+ --page-background-color: #1C1D1F;
+ --page-foreground-color: #d2dbde;
+ --page-secondary-foreground-color: #859399;
+ --separator-color: #38393b;
+ --side-nav-background: #252628;
+
+ --code-background: #2a2c2f;
+
+ --tablehead-background: #2a2c2f;
+
+ --blockquote-background: #222325;
+ --blockquote-foreground: #7e8c92;
+
+ --warning-color: #2e1917;
+ --warning-color-dark: #ad2617;
+ --warning-color-darker: #f5b1aa;
+ --note-color: #3b2e04;
+ --note-color-dark: #f1b602;
+ --note-color-darker: #ceb670;
+ --todo-color: #163750;
+ --todo-color-dark: #1982D2;
+ --todo-color-darker: #dcf0fa;
+ --deprecated-color: #2e323b;
+ --deprecated-color-dark: #738396;
+ --deprecated-color-darker: #abb0bd;
+ --bug-color: #2a2536;
+ --bug-color-dark: #7661b3;
+ --bug-color-darker: #ae9ed6;
+ --invariant-color: #303a35;
+ --invariant-color-dark: #76ce96;
+ --invariant-color-darker: #cceed5;
+
+ --fragment-background: #282c34;
+ --fragment-foreground: #dbe4eb;
+ --fragment-keyword: #cc99cd;
+ --fragment-keywordtype: #ab99cd;
+ --fragment-keywordflow: #e08000;
+ --fragment-token: #7ec699;
+ --fragment-comment: #999999;
+ --fragment-link: #98c0e3;
+ --fragment-preprocessor: #65cabe;
+ --fragment-linenumber-color: #cccccc;
+ --fragment-linenumber-background: #35393c;
+ --fragment-linenumber-border: #1f1f1f;
+ }
+}
+
+/* dark mode variables are defined twice, to support both the dark-mode without and with doxygen-awesome-darkmode-toggle.js */
+html.dark-mode {
+ color-scheme: dark;
+
+ --primary-color: #1982d2;
+ --primary-dark-color: #86a9c4;
+ --primary-light-color: #4779ac;
+
+ --box-shadow: 0 2px 8px 0 rgba(0,0,0,.30);
+
+ --odd-color: rgba(100,100,100,.06);
+
+ --menu-selected-background: rgba(0,0,0,.4);
+
+ --page-background-color: #1C1D1F;
+ --page-foreground-color: #d2dbde;
+ --page-secondary-foreground-color: #859399;
+ --separator-color: #38393b;
+ --side-nav-background: #252628;
+
+ --code-background: #2a2c2f;
+
+ --tablehead-background: #2a2c2f;
+
+ --blockquote-background: #222325;
+ --blockquote-foreground: #7e8c92;
+
+ --warning-color: #2e1917;
+ --warning-color-dark: #ad2617;
+ --warning-color-darker: #f5b1aa;
+ --note-color: #3b2e04;
+ --note-color-dark: #f1b602;
+ --note-color-darker: #ceb670;
+ --todo-color: #163750;
+ --todo-color-dark: #1982D2;
+ --todo-color-darker: #dcf0fa;
+ --deprecated-color: #2e323b;
+ --deprecated-color-dark: #738396;
+ --deprecated-color-darker: #abb0bd;
+ --bug-color: #2a2536;
+ --bug-color-dark: #7661b3;
+ --bug-color-darker: #ae9ed6;
+ --invariant-color: #303a35;
+ --invariant-color-dark: #76ce96;
+ --invariant-color-darker: #cceed5;
+
+ --fragment-background: #282c34;
+ --fragment-foreground: #dbe4eb;
+ --fragment-keyword: #cc99cd;
+ --fragment-keywordtype: #ab99cd;
+ --fragment-keywordflow: #e08000;
+ --fragment-token: #7ec699;
+ --fragment-comment: #999999;
+ --fragment-link: #98c0e3;
+ --fragment-preprocessor: #65cabe;
+ --fragment-linenumber-color: #cccccc;
+ --fragment-linenumber-background: #35393c;
+ --fragment-linenumber-border: #1f1f1f;
+}
+
+body {
+ color: var(--page-foreground-color);
+ background-color: var(--page-background-color);
+ font-size: var(--page-font-size);
+}
+
+body, table, div, p, dl, #nav-tree .label, .title,
+.sm-dox a, .sm-dox a:hover, .sm-dox a:focus, #projectname,
+.SelectItem, #MSearchField, .navpath li.navelem a,
+.navpath li.navelem a:hover, p.reference, p.definition {
+ font-family: var(--font-family);
+}
+
+h1, h2, h3, h4, h5 {
+ margin-top: .9em;
+ font-weight: 600;
+ line-height: initial;
+}
+
+p, div, table, dl, p.reference, p.definition {
+ font-size: var(--page-font-size);
+}
+
+p.reference, p.definition {
+ color: var(--page-secondary-foreground-color);
+}
+
+a:link, a:visited, a:hover, a:focus, a:active {
+ color: var(--primary-color) !important;
+ font-weight: 500;
+}
+
+a.anchor {
+ scroll-margin-top: var(--spacing-large);
+ display: block;
+}
+
+/*
+Title and top navigation
+*/
+
+#top {
+ background: var(--header-background);
+ border-bottom: 1px solid var(--separator-color);
+}
+
+@media screen and (min-width: 768px) {
+ #top {
+ display: flex;
+ flex-wrap: wrap;
+ justify-content: space-between;
+ align-items: center;
+ }
+}
+
+#main-nav {
+ flex-grow: 5;
+ padding: var(--spacing-small) var(--spacing-medium);
+}
+
+#titlearea {
+ width: auto;
+ padding: var(--spacing-medium) var(--spacing-large);
+ background: none;
+ color: var(--header-foreground);
+ border-bottom: none;
+}
+
+@media screen and (max-width: 767px) {
+ #titlearea {
+ padding-bottom: var(--spacing-small);
+ }
+}
+
+#titlearea table tbody tr {
+ height: auto !important;
+}
+
+#projectname {
+ font-size: var(--title-font-size);
+ font-weight: 600;
+}
+
+#projectnumber {
+ font-family: inherit;
+ font-size: 60%;
+}
+
+#projectbrief {
+ font-family: inherit;
+ font-size: 80%;
+}
+
+#projectlogo {
+ vertical-align: middle;
+}
+
+#projectlogo img {
+ max-height: 64px;
+ margin-right: var(--spacing-small);
+}
+
+.sm-dox, .tabs, .tabs2, .tabs3 {
+ background: none;
+ padding: 0;
+}
+
+.tabs, .tabs2, .tabs3 {
+ border-bottom: 1px solid var(--separator-color);
+ margin-bottom: -1px;
+}
+
+.main-menu-btn-icon, .main-menu-btn-icon:before, .main-menu-btn-icon:after {
+ background: var(--page-secondary-foreground-color);
+}
+
+@media screen and (max-width: 767px) {
+ .sm-dox a span.sub-arrow {
+ background: var(--code-background);
+ }
+
+ #main-menu a.has-submenu span.sub-arrow {
+ color: var(--page-secondary-foreground-color);
+ border-radius: var(--border-radius-medium);
+ }
+
+ #main-menu a.has-submenu:hover span.sub-arrow {
+ color: var(--page-foreground-color);
+ }
+}
+
+@media screen and (min-width: 768px) {
+ .sm-dox li, .tablist li {
+ display: var(--menu-display);
+ }
+
+ .sm-dox a span.sub-arrow {
+ border-color: var(--header-foreground) transparent transparent transparent;
+ }
+
+ .sm-dox a:hover span.sub-arrow {
+ border-color: var(--menu-focus-foreground) transparent transparent transparent;
+ }
+
+ .sm-dox ul a span.sub-arrow {
+ border-color: transparent transparent transparent var(--page-foreground-color);
+ }
+
+ .sm-dox ul a:hover span.sub-arrow {
+ border-color: transparent transparent transparent var(--menu-focus-foreground);
+ }
+}
+
+.sm-dox ul {
+ background: var(--page-background-color);
+ box-shadow: var(--box-shadow);
+ border: 1px solid var(--separator-color);
+ border-radius: var(--border-radius-medium) !important;
+ padding: var(--spacing-small);
+ animation: ease-out 150ms slideInMenu;
+}
+
+@keyframes slideInMenu {
+ from {
+ opacity: 0;
+ transform: translate(0px, -2px);
+ }
+
+ to {
+ opacity: 1;
+ transform: translate(0px, 0px);
+ }
+}
+
+.sm-dox ul a {
+ color: var(--page-foreground-color) !important;
+ background: var(--page-background-color);
+ font-size: var(--navigation-font-size);
+}
+
+.sm-dox>li>ul:after {
+ border-bottom-color: var(--page-background-color) !important;
+}
+
+.sm-dox>li>ul:before {
+ border-bottom-color: var(--separator-color) !important;
+}
+
+.sm-dox ul a:hover, .sm-dox ul a:active, .sm-dox ul a:focus {
+ font-size: var(--navigation-font-size) !important;
+ color: var(--menu-focus-foreground) !important;
+ text-shadow: none;
+ background-color: var(--menu-focus-background);
+ border-radius: var(--border-radius-small) !important;
+}
+
+.sm-dox a, .sm-dox a:focus, .tablist li, .tablist li a, .tablist li.current a {
+ text-shadow: none;
+ background: transparent;
+ background-image: none !important;
+ color: var(--header-foreground) !important;
+ font-weight: normal;
+ font-size: var(--navigation-font-size);
+ border-radius: var(--border-radius-small) !important;
+}
+
+.sm-dox a:focus {
+ outline: auto;
+}
+
+.sm-dox a:hover, .sm-dox a:active, .tablist li a:hover {
+ text-shadow: none;
+ font-weight: normal;
+ background: var(--menu-focus-background);
+ color: var(--menu-focus-foreground) !important;
+ border-radius: var(--border-radius-small) !important;
+ font-size: var(--navigation-font-size);
+}
+
+.tablist li.current {
+ border-radius: var(--border-radius-small);
+ background: var(--menu-selected-background);
+}
+
+.tablist li {
+ margin: var(--spacing-small) 0 var(--spacing-small) var(--spacing-small);
+}
+
+.tablist a {
+ padding: 0 var(--spacing-large);
+}
+
+
+/*
+Search box
+*/
+
+#MSearchBox {
+ height: var(--searchbar-height);
+ background: var(--searchbar-background);
+ border-radius: var(--searchbar-border-radius);
+ border: 1px solid var(--separator-color);
+ overflow: hidden;
+ width: var(--searchbar-width);
+ position: relative;
+ box-shadow: none;
+ display: block;
+ margin-top: 0;
+}
+
+/* until Doxygen 1.9.4 */
+.left img#MSearchSelect {
+ left: 0;
+ user-select: none;
+ padding-left: 8px;
+}
+
+/* Doxygen 1.9.5 */
+.left span#MSearchSelect {
+ left: 0;
+ user-select: none;
+ margin-left: 8px;
+ padding: 0;
+}
+
+.left #MSearchSelect[src$=".png"] {
+ padding-left: 0
+}
+
+.SelectionMark {
+ user-select: none;
+}
+
+.tabs .left #MSearchSelect {
+ padding-left: 0;
+}
+
+.tabs #MSearchBox {
+ position: absolute;
+ right: var(--spacing-medium);
+}
+
+@media screen and (max-width: 767px) {
+ .tabs #MSearchBox {
+ position: relative;
+ right: 0;
+ margin-left: var(--spacing-medium);
+ margin-top: 0;
+ }
+}
+
+#MSearchSelectWindow, #MSearchResultsWindow {
+ z-index: 9999;
+}
+
+#MSearchBox.MSearchBoxActive {
+ border-color: var(--primary-color);
+ box-shadow: inset 0 0 0 1px var(--primary-color);
+}
+
+#main-menu > li:last-child {
+ margin-right: 0;
+}
+
+@media screen and (max-width: 767px) {
+ #main-menu > li:last-child {
+ height: 50px;
+ }
+}
+
+#MSearchField {
+ font-size: var(--navigation-font-size);
+ height: calc(var(--searchbar-height) - 2px);
+ background: transparent;
+ width: calc(var(--searchbar-width) - 64px);
+}
+
+.MSearchBoxActive #MSearchField {
+ color: var(--searchbar-foreground);
+}
+
+#MSearchSelect {
+ top: calc(calc(var(--searchbar-height) / 2) - 11px);
+}
+
+#MSearchBox span.left, #MSearchBox span.right {
+ background: none;
+ background-image: none;
+}
+
+#MSearchBox span.right {
+ padding-top: calc(calc(var(--searchbar-height) / 2) - 12px);
+ position: absolute;
+ right: var(--spacing-small);
+}
+
+.tabs #MSearchBox span.right {
+ top: calc(calc(var(--searchbar-height) / 2) - 12px);
+}
+
+@keyframes slideInSearchResults {
+ from {
+ opacity: 0;
+ transform: translate(0, 15px);
+ }
+
+ to {
+ opacity: 1;
+ transform: translate(0, 20px);
+ }
+}
+
+#MSearchResultsWindow {
+ left: auto !important;
+ right: var(--spacing-medium);
+ border-radius: var(--border-radius-large);
+ border: 1px solid var(--separator-color);
+ transform: translate(0, 20px);
+ box-shadow: var(--box-shadow);
+ animation: ease-out 280ms slideInSearchResults;
+ background: var(--page-background-color);
+}
+
+iframe#MSearchResults {
+ margin: 4px;
+}
+
+iframe {
+ color-scheme: normal;
+}
+
+@media (prefers-color-scheme: dark) {
+ html:not(.light-mode) iframe#MSearchResults {
+ filter: invert() hue-rotate(180deg);
+ }
+}
+
+html.dark-mode iframe#MSearchResults {
+ filter: invert() hue-rotate(180deg);
+}
+
+#MSearchResults .SRPage {
+ background-color: transparent;
+}
+
+#MSearchResults .SRPage .SREntry {
+ font-size: 10pt;
+ padding: var(--spacing-small) var(--spacing-medium);
+}
+
+#MSearchSelectWindow {
+ border: 1px solid var(--separator-color);
+ border-radius: var(--border-radius-medium);
+ box-shadow: var(--box-shadow);
+ background: var(--page-background-color);
+ padding-top: var(--spacing-small);
+ padding-bottom: var(--spacing-small);
+}
+
+#MSearchSelectWindow a.SelectItem {
+ font-size: var(--navigation-font-size);
+ line-height: var(--content-line-height);
+ margin: 0 var(--spacing-small);
+ border-radius: var(--border-radius-small);
+ color: var(--page-foreground-color) !important;
+ font-weight: normal;
+}
+
+#MSearchSelectWindow a.SelectItem:hover {
+ background: var(--menu-focus-background);
+ color: var(--menu-focus-foreground) !important;
+}
+
+@media screen and (max-width: 767px) {
+ #MSearchBox {
+ margin-top: var(--spacing-medium);
+ margin-bottom: var(--spacing-medium);
+ width: calc(100vw - 30px);
+ }
+
+ #main-menu > li:last-child {
+ float: none !important;
+ }
+
+ #MSearchField {
+ width: calc(100vw - 110px);
+ }
+
+ @keyframes slideInSearchResultsMobile {
+ from {
+ opacity: 0;
+ transform: translate(0, 15px);
+ }
+
+ to {
+ opacity: 1;
+ transform: translate(0, 20px);
+ }
+ }
+
+ #MSearchResultsWindow {
+ left: var(--spacing-medium) !important;
+ right: var(--spacing-medium);
+ overflow: auto;
+ transform: translate(0, 20px);
+ animation: ease-out 280ms slideInSearchResultsMobile;
+ width: auto !important;
+ }
+
+ /*
+ * Overwrites for fixing the searchbox on mobile in doxygen 1.9.2
+ */
+ label.main-menu-btn ~ #searchBoxPos1 {
+ top: 3px !important;
+ right: 6px !important;
+ left: 45px;
+ display: flex;
+ }
+
+ label.main-menu-btn ~ #searchBoxPos1 > #MSearchBox {
+ margin-top: 0;
+ margin-bottom: 0;
+ flex-grow: 2;
+ float: left;
+ }
+}
+
+/*
+Tree view
+*/
+
+#side-nav {
+ padding: 0 !important;
+ background: var(--side-nav-background);
+}
+
+@media screen and (max-width: 767px) {
+ #side-nav {
+ display: none;
+ }
+
+ #doc-content {
+ margin-left: 0 !important;
+ }
+}
+
+#nav-tree {
+ background: transparent;
+}
+
+#nav-tree .label {
+ font-size: var(--navigation-font-size);
+}
+
+#nav-tree .item {
+ height: var(--tree-item-height);
+ line-height: var(--tree-item-height);
+}
+
+#nav-sync {
+ bottom: 12px;
+ right: 12px;
+ top: auto !important;
+ user-select: none;
+}
+
+#nav-tree .selected {
+ text-shadow: none;
+ background-image: none;
+ background-color: transparent;
+ position: relative;
+}
+
+#nav-tree .selected::after {
+ content: "";
+ position: absolute;
+ top: 1px;
+ bottom: 1px;
+ left: 0;
+ width: 4px;
+ border-radius: 0 var(--border-radius-small) var(--border-radius-small) 0;
+ background: var(--primary-color);
+}
+
+
+#nav-tree a {
+ color: var(--side-nav-foreground) !important;
+ font-weight: normal;
+}
+
+#nav-tree a:focus {
+ outline-style: auto;
+}
+
+#nav-tree .arrow {
+ opacity: var(--side-nav-arrow-opacity);
+}
+
+.arrow {
+ color: inherit;
+ cursor: pointer;
+ font-size: 45%;
+ vertical-align: middle;
+ margin-right: 2px;
+ font-family: serif;
+ height: auto;
+ text-align: right;
+}
+
+#nav-tree div.item:hover .arrow, #nav-tree a:focus .arrow {
+ opacity: var(--side-nav-arrow-hover-opacity);
+}
+
+#nav-tree .selected a {
+ color: var(--primary-color) !important;
+ font-weight: bolder;
+ font-weight: 600;
+}
+
+.ui-resizable-e {
+ background: var(--separator-color);
+ width: 1px;
+}
+
+/*
+Contents
+*/
+
+div.header {
+ border-bottom: 1px solid var(--separator-color);
+ background-color: var(--page-background-color);
+ background-image: none;
+}
+
+@media screen and (min-width: 1000px) {
+ #doc-content > div > div.contents,
+ .PageDoc > div.contents {
+ display: flex;
+ flex-direction: row-reverse;
+ flex-wrap: nowrap;
+ align-items: flex-start;
+ }
+
+ div.contents .textblock {
+ min-width: 200px;
+ flex-grow: 1;
+ }
+}
+
+div.contents, div.header .title, div.header .summary {
+ max-width: var(--content-maxwidth);
+}
+
+div.contents, div.header .title {
+ line-height: initial;
+ margin: calc(var(--spacing-medium) + .2em) auto var(--spacing-medium) auto;
+}
+
+div.header .summary {
+ margin: var(--spacing-medium) auto 0 auto;
+}
+
+div.headertitle {
+ padding: 0;
+}
+
+div.header .title {
+ font-weight: 600;
+ font-size: 225%;
+ padding: var(--spacing-medium) var(--spacing-large);
+ word-break: break-word;
+}
+
+div.header .summary {
+ width: auto;
+ display: block;
+ float: none;
+ padding: 0 var(--spacing-large);
+}
+
+td.memSeparator {
+ border-color: var(--separator-color);
+}
+
+span.mlabel {
+ background: var(--primary-color);
+ border: none;
+ padding: 4px 9px;
+ border-radius: 12px;
+ margin-right: var(--spacing-medium);
+}
+
+span.mlabel:last-of-type {
+ margin-right: 2px;
+}
+
+div.contents {
+ padding: 0 var(--spacing-large);
+}
+
+div.contents p, div.contents li {
+ line-height: var(--content-line-height);
+}
+
+div.contents div.dyncontent {
+ margin: var(--spacing-medium) 0;
+}
+
+@media (prefers-color-scheme: dark) {
+ html:not(.light-mode) div.contents div.dyncontent img,
+ html:not(.light-mode) div.contents center img,
+ html:not(.light-mode) div.contents > table img,
+ html:not(.light-mode) div.contents div.dyncontent iframe,
+ html:not(.light-mode) div.contents center iframe,
+ html:not(.light-mode) div.contents table iframe {
+ filter: hue-rotate(180deg) invert();
+ }
+}
+
+html.dark-mode div.contents div.dyncontent img,
+html.dark-mode div.contents center img,
+html.dark-mode div.contents > table img,
+html.dark-mode div.contents div.dyncontent iframe,
+html.dark-mode div.contents center iframe,
+html.dark-mode div.contents table iframe {
+ filter: hue-rotate(180deg) invert();
+}
+
+h2.groupheader {
+ border-bottom: 0px;
+ color: var(--page-foreground-color);
+ box-shadow:
+ 100px 0 var(--page-background-color),
+ -100px 0 var(--page-background-color),
+ 100px 0.75px var(--separator-color),
+ -100px 0.75px var(--separator-color),
+ 500px 0 var(--page-background-color),
+ -500px 0 var(--page-background-color),
+ 500px 0.75px var(--separator-color),
+ -500px 0.75px var(--separator-color),
+ 900px 0 var(--page-background-color),
+ -900px 0 var(--page-background-color),
+ 900px 0.75px var(--separator-color),
+ -900px 0.75px var(--separator-color),
+ 1400px 0 var(--page-background-color),
+ -1400px 0 var(--page-background-color),
+ 1400px 0.75px var(--separator-color),
+ -1400px 0.75px var(--separator-color),
+ 1900px 0 var(--page-background-color),
+ -1900px 0 var(--page-background-color),
+ 1900px 0.75px var(--separator-color),
+ -1900px 0.75px var(--separator-color);
+}
+
+blockquote {
+ margin: 0 var(--spacing-medium) 0 var(--spacing-medium);
+ padding: var(--spacing-small) var(--spacing-large);
+ background: var(--blockquote-background);
+ color: var(--blockquote-foreground);
+ border-left: 0;
+ overflow: visible;
+ border-radius: var(--border-radius-medium);
+ overflow: visible;
+ position: relative;
+}
+
+blockquote::before, blockquote::after {
+ font-weight: bold;
+ font-family: serif;
+ font-size: 360%;
+ opacity: .15;
+ position: absolute;
+}
+
+blockquote::before {
+ content: "“";
+ left: -10px;
+ top: 4px;
+}
+
+blockquote::after {
+ content: "”";
+ right: -8px;
+ bottom: -25px;
+}
+
+blockquote p {
+ margin: var(--spacing-small) 0 var(--spacing-medium) 0;
+}
+.paramname {
+ font-weight: 600;
+ color: var(--primary-dark-color);
+}
+
+.paramname > code {
+ border: 0;
+}
+
+table.params .paramname {
+ font-weight: 600;
+ font-family: var(--font-family-monospace);
+ font-size: var(--code-font-size);
+ padding-right: var(--spacing-small);
+ line-height: var(--table-line-height);
+}
+
+h1.glow, h2.glow, h3.glow, h4.glow, h5.glow, h6.glow {
+ text-shadow: 0 0 15px var(--primary-light-color);
+}
+
+.alphachar a {
+ color: var(--page-foreground-color);
+}
+
+/*
+Table of Contents
+*/
+
+div.contents .toc {
+ max-height: var(--toc-max-height);
+ min-width: var(--toc-width);
+ border: 0;
+ border-left: 1px solid var(--separator-color);
+ border-radius: 0;
+ background-color: transparent;
+ box-shadow: none;
+ position: sticky;
+ top: var(--toc-sticky-top);
+ padding: 0 var(--spacing-large);
+ margin: var(--spacing-small) 0 var(--spacing-large) var(--spacing-large);
+}
+
+div.toc h3 {
+ color: var(--toc-foreground);
+ font-size: var(--navigation-font-size);
+ margin: var(--spacing-large) 0 var(--spacing-medium) 0;
+}
+
+div.toc li {
+ padding: 0;
+ background: none;
+ line-height: var(--toc-font-size);
+ margin: var(--toc-font-size) 0 0 0;
+}
+
+div.toc li::before {
+ display: none;
+}
+
+div.toc ul {
+ margin-top: 0
+}
+
+div.toc li a {
+ font-size: var(--toc-font-size);
+ color: var(--page-foreground-color) !important;
+ text-decoration: none;
+}
+
+div.toc li a:hover, div.toc li a.active {
+ color: var(--primary-color) !important;
+}
+
+div.toc li a.aboveActive {
+ color: var(--page-secondary-foreground-color) !important;
+}
+
+
+@media screen and (max-width: 999px) {
+ div.contents .toc {
+ max-height: 45vh;
+ float: none;
+ width: auto;
+ margin: 0 0 var(--spacing-medium) 0;
+ position: relative;
+ top: 0;
+ position: relative;
+ border: 1px solid var(--separator-color);
+ border-radius: var(--border-radius-medium);
+ background-color: var(--toc-background);
+ box-shadow: var(--box-shadow);
+ }
+
+ div.contents .toc.interactive {
+ max-height: calc(var(--navigation-font-size) + 2 * var(--spacing-large));
+ overflow: hidden;
+ }
+
+ div.contents .toc > h3 {
+ -webkit-tap-highlight-color: transparent;
+ cursor: pointer;
+ position: sticky;
+ top: 0;
+ background-color: var(--toc-background);
+ margin: 0;
+ padding: var(--spacing-large) 0;
+ display: block;
+ }
+
+ div.contents .toc.interactive > h3::before {
+ content: "";
+ width: 0;
+ height: 0;
+ border-left: 4px solid transparent;
+ border-right: 4px solid transparent;
+ border-top: 5px solid var(--primary-color);
+ display: inline-block;
+ margin-right: var(--spacing-small);
+ margin-bottom: calc(var(--navigation-font-size) / 4);
+ transform: rotate(-90deg);
+ transition: transform 0.25s ease-out;
+ }
+
+ div.contents .toc.interactive.open > h3::before {
+ transform: rotate(0deg);
+ }
+
+ div.contents .toc.interactive.open {
+ max-height: 45vh;
+ overflow: auto;
+ transition: max-height 0.2s ease-in-out;
+ }
+
+ div.contents .toc a, div.contents .toc a.active {
+ color: var(--primary-color) !important;
+ }
+
+ div.contents .toc a:hover {
+ text-decoration: underline;
+ }
+}
+
+/*
+Code & Fragments
+*/
+
+code, div.fragment, pre.fragment {
+ border-radius: var(--border-radius-small);
+ border: 1px solid var(--separator-color);
+ overflow: hidden;
+}
+
+code {
+ display: inline;
+ background: var(--code-background);
+ color: var(--code-foreground);
+ padding: 2px 6px;
+}
+
+div.fragment, pre.fragment {
+ margin: var(--spacing-medium) 0;
+ padding: calc(var(--spacing-large) - (var(--spacing-large) / 6)) var(--spacing-large);
+ background: var(--fragment-background);
+ color: var(--fragment-foreground);
+ overflow-x: auto;
+}
+
+@media screen and (max-width: 767px) {
+ div.fragment, pre.fragment {
+ border-top-right-radius: 0;
+ border-bottom-right-radius: 0;
+ border-right: 0;
+ }
+
+ .contents > div.fragment,
+ .textblock > div.fragment,
+ .textblock > pre.fragment,
+ .contents > .doxygen-awesome-fragment-wrapper > div.fragment,
+ .textblock > .doxygen-awesome-fragment-wrapper > div.fragment,
+ .textblock > .doxygen-awesome-fragment-wrapper > pre.fragment {
+ margin: var(--spacing-medium) calc(0px - var(--spacing-large));
+ border-radius: 0;
+ border-left: 0;
+ }
+
+ .textblock li > .fragment,
+ .textblock li > .doxygen-awesome-fragment-wrapper > .fragment {
+ margin: var(--spacing-medium) calc(0px - var(--spacing-large));
+ }
+
+ .memdoc li > .fragment,
+ .memdoc li > .doxygen-awesome-fragment-wrapper > .fragment {
+ margin: var(--spacing-medium) calc(0px - var(--spacing-medium));
+ }
+
+ .textblock ul, .memdoc ul {
+ overflow: initial;
+ }
+
+ .memdoc > div.fragment,
+ .memdoc > pre.fragment,
+ dl dd > div.fragment,
+ dl dd pre.fragment,
+ .memdoc > .doxygen-awesome-fragment-wrapper > div.fragment,
+ .memdoc > .doxygen-awesome-fragment-wrapper > pre.fragment,
+ dl dd > .doxygen-awesome-fragment-wrapper > div.fragment,
+ dl dd .doxygen-awesome-fragment-wrapper > pre.fragment {
+ margin: var(--spacing-medium) calc(0px - var(--spacing-medium));
+ border-radius: 0;
+ border-left: 0;
+ }
+}
+
+code, code a, pre.fragment, div.fragment, div.fragment .line, div.fragment span, div.fragment .line a, div.fragment .line span {
+ font-family: var(--font-family-monospace);
+ font-size: var(--code-font-size) !important;
+}
+
+div.line:after {
+ margin-right: var(--spacing-medium);
+}
+
+div.fragment .line, pre.fragment {
+ white-space: pre;
+ word-wrap: initial;
+ line-height: var(--fragment-lineheight);
+}
+
+div.fragment span.keyword {
+ color: var(--fragment-keyword);
+}
+
+div.fragment span.keywordtype {
+ color: var(--fragment-keywordtype);
+}
+
+div.fragment span.keywordflow {
+ color: var(--fragment-keywordflow);
+}
+
+div.fragment span.stringliteral {
+ color: var(--fragment-token)
+}
+
+div.fragment span.comment {
+ color: var(--fragment-comment);
+}
+
+div.fragment a.code {
+ color: var(--fragment-link) !important;
+}
+
+div.fragment span.preprocessor {
+ color: var(--fragment-preprocessor);
+}
+
+div.fragment span.lineno {
+ display: inline-block;
+ width: 27px;
+ border-right: none;
+ background: var(--fragment-linenumber-background);
+ color: var(--fragment-linenumber-color);
+}
+
+div.fragment span.lineno a {
+ background: none;
+ color: var(--fragment-link) !important;
+}
+
+div.fragment .line:first-child .lineno {
+ box-shadow: -999999px 0px 0 999999px var(--fragment-linenumber-background), -999998px 0px 0 999999px var(--fragment-linenumber-border);
+}
+
+div.line {
+ border-radius: var(--border-radius-small);
+}
+
+div.line.glow {
+ background-color: var(--primary-light-color);
+ box-shadow: none;
+}
+
+/*
+dl warning, attention, note, deprecated, bug, ...
+*/
+
+dl.bug dt a, dl.deprecated dt a, dl.todo dt a {
+ font-weight: bold !important;
+}
+
+dl.warning, dl.attention, dl.note, dl.deprecated, dl.bug, dl.invariant, dl.pre, dl.post, dl.todo, dl.remark {
+ padding: var(--spacing-medium);
+ margin: var(--spacing-medium) 0;
+ color: var(--page-background-color);
+ overflow: hidden;
+ margin-left: 0;
+ border-radius: var(--border-radius-small);
+}
+
+dl.section dd {
+ margin-bottom: 2px;
+}
+
+dl.warning, dl.attention {
+ background: var(--warning-color);
+ border-left: 8px solid var(--warning-color-dark);
+ color: var(--warning-color-darker);
+}
+
+dl.warning dt, dl.attention dt {
+ color: var(--warning-color-dark);
+}
+
+dl.note, dl.remark {
+ background: var(--note-color);
+ border-left: 8px solid var(--note-color-dark);
+ color: var(--note-color-darker);
+}
+
+dl.note dt, dl.remark dt {
+ color: var(--note-color-dark);
+}
+
+dl.todo {
+ background: var(--todo-color);
+ border-left: 8px solid var(--todo-color-dark);
+ color: var(--todo-color-darker);
+}
+
+dl.todo dt {
+ color: var(--todo-color-dark);
+}
+
+dl.bug dt a {
+ color: var(--todo-color-dark) !important;
+}
+
+dl.bug {
+ background: var(--bug-color);
+ border-left: 8px solid var(--bug-color-dark);
+ color: var(--bug-color-darker);
+}
+
+dl.bug dt a {
+ color: var(--bug-color-dark) !important;
+}
+
+dl.deprecated {
+ background: var(--deprecated-color);
+ border-left: 8px solid var(--deprecated-color-dark);
+ color: var(--deprecated-color-darker);
+}
+
+dl.deprecated dt a {
+ color: var(--deprecated-color-dark) !important;
+}
+
+dl.section dd, dl.bug dd, dl.deprecated dd, dl.todo dd {
+ margin-inline-start: 0px;
+}
+
+dl.invariant, dl.pre, dl.post {
+ background: var(--invariant-color);
+ border-left: 8px solid var(--invariant-color-dark);
+ color: var(--invariant-color-darker);
+}
+
+dl.invariant dt, dl.pre dt, dl.post dt {
+ color: var(--invariant-color-dark);
+}
+
+/*
+memitem
+*/
+
+div.memdoc, div.memproto, h2.memtitle {
+ box-shadow: none;
+ background-image: none;
+ border: none;
+}
+
+div.memdoc {
+ padding: 0 var(--spacing-medium);
+ background: var(--page-background-color);
+}
+
+h2.memtitle, div.memitem {
+ border: 1px solid var(--separator-color);
+ box-shadow: var(--box-shadow);
+}
+
+h2.memtitle {
+ box-shadow: 0px var(--spacing-medium) 0 -1px var(--fragment-background), var(--box-shadow);
+}
+
+div.memitem {
+ transition: none;
+}
+
+div.memproto, h2.memtitle {
+ background: var(--fragment-background);
+}
+
+h2.memtitle {
+ font-weight: 500;
+ font-size: var(--memtitle-font-size);
+ font-family: var(--font-family-monospace);
+ border-bottom: none;
+ border-top-left-radius: var(--border-radius-medium);
+ border-top-right-radius: var(--border-radius-medium);
+ word-break: break-all;
+ position: relative;
+}
+
+h2.memtitle:after {
+ content: "";
+ display: block;
+ background: var(--fragment-background);
+ height: var(--spacing-medium);
+ bottom: calc(0px - var(--spacing-medium));
+ left: 0;
+ right: -14px;
+ position: absolute;
+ border-top-right-radius: var(--border-radius-medium);
+}
+
+h2.memtitle > span.permalink {
+ font-size: inherit;
+}
+
+h2.memtitle > span.permalink > a {
+ text-decoration: none;
+ padding-left: 3px;
+ margin-right: -4px;
+ user-select: none;
+ display: inline-block;
+ margin-top: -6px;
+}
+
+h2.memtitle > span.permalink > a:hover {
+ color: var(--primary-dark-color) !important;
+}
+
+a:target + h2.memtitle, a:target + h2.memtitle + div.memitem {
+ border-color: var(--primary-light-color);
+}
+
+div.memitem {
+ border-top-right-radius: var(--border-radius-medium);
+ border-bottom-right-radius: var(--border-radius-medium);
+ border-bottom-left-radius: var(--border-radius-medium);
+ overflow: hidden;
+ display: block !important;
+}
+
+div.memdoc {
+ border-radius: 0;
+}
+
+div.memproto {
+ border-radius: 0 var(--border-radius-small) 0 0;
+ overflow: auto;
+ border-bottom: 1px solid var(--separator-color);
+ padding: var(--spacing-medium);
+ margin-bottom: -1px;
+}
+
+div.memtitle {
+ border-top-right-radius: var(--border-radius-medium);
+ border-top-left-radius: var(--border-radius-medium);
+}
+
+div.memproto table.memname {
+ font-family: var(--font-family-monospace);
+ color: var(--page-foreground-color);
+ font-size: var(--memname-font-size);
+ text-shadow: none;
+}
+
+div.memproto div.memtemplate {
+ font-family: var(--font-family-monospace);
+ color: var(--primary-dark-color);
+ font-size: var(--memname-font-size);
+ margin-left: 2px;
+ text-shadow: none;
+}
+
+table.mlabels, table.mlabels > tbody {
+ display: block;
+}
+
+td.mlabels-left {
+ width: auto;
+}
+
+td.mlabels-right {
+ margin-top: 3px;
+ position: sticky;
+ left: 0;
+}
+
+table.mlabels > tbody > tr:first-child {
+ display: flex;
+ justify-content: space-between;
+ flex-wrap: wrap;
+}
+
+.memname, .memitem span.mlabels {
+ margin: 0
+}
+
+/*
+reflist
+*/
+
+dl.reflist {
+ box-shadow: var(--box-shadow);
+ border-radius: var(--border-radius-medium);
+ border: 1px solid var(--separator-color);
+ overflow: hidden;
+ padding: 0;
+}
+
+
+dl.reflist dt, dl.reflist dd {
+ box-shadow: none;
+ text-shadow: none;
+ background-image: none;
+ border: none;
+ padding: 12px;
+}
+
+
+dl.reflist dt {
+ font-weight: 500;
+ border-radius: 0;
+ background: var(--code-background);
+ border-bottom: 1px solid var(--separator-color);
+ color: var(--page-foreground-color)
+}
+
+
+dl.reflist dd {
+ background: none;
+}
+
+/*
+Table
+*/
+
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname),
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody {
+ display: inline-block;
+ max-width: 100%;
+}
+
+.contents > table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname):not(.classindex) {
+ margin-left: calc(0px - var(--spacing-large));
+ margin-right: calc(0px - var(--spacing-large));
+ max-width: calc(100% + 2 * var(--spacing-large));
+}
+
+table.fieldtable,
+table.markdownTable tbody,
+table.doxtable tbody {
+ border: none;
+ margin: var(--spacing-medium) 0;
+ box-shadow: 0 0 0 1px var(--separator-color);
+ border-radius: var(--border-radius-small);
+}
+
+table.doxtable caption {
+ display: block;
+}
+
+table.fieldtable {
+ border-collapse: collapse;
+ width: 100%;
+}
+
+th.markdownTableHeadLeft,
+th.markdownTableHeadRight,
+th.markdownTableHeadCenter,
+th.markdownTableHeadNone,
+table.doxtable th {
+ background: var(--tablehead-background);
+ color: var(--tablehead-foreground);
+ font-weight: 600;
+ font-size: var(--page-font-size);
+}
+
+th.markdownTableHeadLeft:first-child,
+th.markdownTableHeadRight:first-child,
+th.markdownTableHeadCenter:first-child,
+th.markdownTableHeadNone:first-child,
+table.doxtable tr th:first-child {
+ border-top-left-radius: var(--border-radius-small);
+}
+
+th.markdownTableHeadLeft:last-child,
+th.markdownTableHeadRight:last-child,
+th.markdownTableHeadCenter:last-child,
+th.markdownTableHeadNone:last-child,
+table.doxtable tr th:last-child {
+ border-top-right-radius: var(--border-radius-small);
+}
+
+table.markdownTable td,
+table.markdownTable th,
+table.fieldtable td,
+table.fieldtable th,
+table.doxtable td,
+table.doxtable th {
+ border: 1px solid var(--separator-color);
+ padding: var(--spacing-small) var(--spacing-medium);
+}
+
+table.markdownTable td:last-child,
+table.markdownTable th:last-child,
+table.fieldtable td:last-child,
+table.fieldtable th:last-child,
+table.doxtable td:last-child,
+table.doxtable th:last-child {
+ border-right: none;
+}
+
+table.markdownTable td:first-child,
+table.markdownTable th:first-child,
+table.fieldtable td:first-child,
+table.fieldtable th:first-child,
+table.doxtable td:first-child,
+table.doxtable th:first-child {
+ border-left: none;
+}
+
+table.markdownTable tr:first-child td,
+table.markdownTable tr:first-child th,
+table.fieldtable tr:first-child td,
+table.fieldtable tr:first-child th,
+table.doxtable tr:first-child td,
+table.doxtable tr:first-child th {
+ border-top: none;
+}
+
+table.markdownTable tr:last-child td,
+table.markdownTable tr:last-child th,
+table.fieldtable tr:last-child td,
+table.fieldtable tr:last-child th,
+table.doxtable tr:last-child td,
+table.doxtable tr:last-child th {
+ border-bottom: none;
+}
+
+table.markdownTable tr, table.doxtable tr {
+ border-bottom: 1px solid var(--separator-color);
+}
+
+table.markdownTable tr:last-child, table.doxtable tr:last-child {
+ border-bottom: none;
+}
+
+table.fieldtable th {
+ font-size: var(--page-font-size);
+ font-weight: 600;
+ background-image: none;
+ background-color: var(--tablehead-background);
+ color: var(--tablehead-foreground);
+}
+
+table.fieldtable td.fieldtype, .fieldtable td.fieldname, .fieldtable td.fielddoc, .fieldtable th {
+ border-bottom: 1px solid var(--separator-color);
+ border-right: 1px solid var(--separator-color);
+}
+
+table.fieldtable tr:last-child td:first-child {
+ border-bottom-left-radius: var(--border-radius-small);
+}
+
+table.fieldtable tr:last-child td:last-child {
+ border-bottom-right-radius: var(--border-radius-small);
+}
+
+.memberdecls td.glow, .fieldtable tr.glow {
+ background-color: var(--primary-light-color);
+ box-shadow: none;
+}
+
+table.memberdecls {
+ display: block;
+ -webkit-tap-highlight-color: transparent;
+}
+
+table.memberdecls tr[class^='memitem'] {
+ font-family: var(--font-family-monospace);
+ font-size: var(--code-font-size);
+}
+
+table.memberdecls tr[class^='memitem'] .memTemplParams {
+ font-family: var(--font-family-monospace);
+ font-size: var(--code-font-size);
+ color: var(--primary-dark-color);
+ white-space: normal;
+}
+
+table.memberdecls .memItemLeft,
+table.memberdecls .memItemRight,
+table.memberdecls .memTemplItemLeft,
+table.memberdecls .memTemplItemRight,
+table.memberdecls .memTemplParams {
+ transition: none;
+ padding-top: var(--spacing-small);
+ padding-bottom: var(--spacing-small);
+ border-top: 1px solid var(--separator-color);
+ border-bottom: 1px solid var(--separator-color);
+ background-color: var(--fragment-background);
+}
+
+table.memberdecls .memTemplItemLeft,
+table.memberdecls .memTemplItemRight {
+ padding-top: 2px;
+}
+
+table.memberdecls .memTemplParams {
+ border-bottom: 0;
+ border-left: 1px solid var(--separator-color);
+ border-right: 1px solid var(--separator-color);
+ border-radius: var(--border-radius-small) var(--border-radius-small) 0 0;
+ padding-bottom: var(--spacing-small);
+}
+
+table.memberdecls .memTemplItemLeft {
+ border-radius: 0 0 0 var(--border-radius-small);
+ border-left: 1px solid var(--separator-color);
+ border-top: 0;
+}
+
+table.memberdecls .memTemplItemRight {
+ border-radius: 0 0 var(--border-radius-small) 0;
+ border-right: 1px solid var(--separator-color);
+ padding-left: 0;
+ border-top: 0;
+}
+
+table.memberdecls .memItemLeft {
+ border-radius: var(--border-radius-small) 0 0 var(--border-radius-small);
+ border-left: 1px solid var(--separator-color);
+ padding-left: var(--spacing-medium);
+ padding-right: 0;
+}
+
+table.memberdecls .memItemRight {
+ border-radius: 0 var(--border-radius-small) var(--border-radius-small) 0;
+ border-right: 1px solid var(--separator-color);
+ padding-right: var(--spacing-medium);
+ padding-left: 0;
+
+}
+
+table.memberdecls .mdescLeft, table.memberdecls .mdescRight {
+ background: none;
+ color: var(--page-foreground-color);
+ padding: var(--spacing-small) 0;
+}
+
+table.memberdecls .memItemLeft,
+table.memberdecls .memTemplItemLeft {
+ padding-right: var(--spacing-medium);
+}
+
+table.memberdecls .memSeparator {
+ background: var(--page-background-color);
+ height: var(--spacing-large);
+ border: 0;
+ transition: none;
+}
+
+table.memberdecls .groupheader {
+ margin-bottom: var(--spacing-large);
+}
+
+table.memberdecls .inherit_header td {
+ padding: 0 0 var(--spacing-medium) 0;
+ text-indent: -12px;
+ color: var(--page-secondary-foreground-color);
+}
+
+table.memberdecls img[src="closed.png"],
+table.memberdecls img[src="open.png"],
+div.dynheader img[src="open.png"],
+div.dynheader img[src="closed.png"] {
+ width: 0;
+ height: 0;
+ border-left: 4px solid transparent;
+ border-right: 4px solid transparent;
+ border-top: 5px solid var(--primary-color);
+ margin-top: 8px;
+ display: block;
+ float: left;
+ margin-left: -10px;
+ transition: transform 0.25s ease-out;
+}
+
+table.memberdecls img {
+ margin-right: 10px;
+}
+
+table.memberdecls img[src="closed.png"],
+div.dynheader img[src="closed.png"] {
+ transform: rotate(-90deg);
+
+}
+
+.compoundTemplParams {
+ font-family: var(--font-family-monospace);
+ color: var(--primary-dark-color);
+ font-size: var(--code-font-size);
+}
+
+@media screen and (max-width: 767px) {
+
+ table.memberdecls .memItemLeft,
+ table.memberdecls .memItemRight,
+ table.memberdecls .mdescLeft,
+ table.memberdecls .mdescRight,
+ table.memberdecls .memTemplItemLeft,
+ table.memberdecls .memTemplItemRight,
+ table.memberdecls .memTemplParams {
+ display: block;
+ text-align: left;
+ padding-left: var(--spacing-large);
+ margin: 0 calc(0px - var(--spacing-large)) 0 calc(0px - var(--spacing-large));
+ border-right: none;
+ border-left: none;
+ border-radius: 0;
+ white-space: normal;
+ }
+
+ table.memberdecls .memItemLeft,
+ table.memberdecls .mdescLeft,
+ table.memberdecls .memTemplItemLeft {
+ border-bottom: 0;
+ padding-bottom: 0;
+ }
+
+ table.memberdecls .memTemplItemLeft {
+ padding-top: 0;
+ }
+
+ table.memberdecls .mdescLeft {
+ margin-bottom: calc(0px - var(--page-font-size));
+ }
+
+ table.memberdecls .memItemRight,
+ table.memberdecls .mdescRight,
+ table.memberdecls .memTemplItemRight {
+ border-top: 0;
+ padding-top: 0;
+ padding-right: var(--spacing-large);
+ overflow-x: auto;
+ }
+
+ table.memberdecls tr[class^='memitem']:not(.inherit) {
+ display: block;
+ width: calc(100vw - 2 * var(--spacing-large));
+ }
+
+ table.memberdecls .mdescRight {
+ color: var(--page-foreground-color);
+ }
+
+ table.memberdecls tr.inherit {
+ visibility: hidden;
+ }
+
+ table.memberdecls tr[style="display: table-row;"] {
+ display: block !important;
+ visibility: visible;
+ width: calc(100vw - 2 * var(--spacing-large));
+ animation: fade .5s;
+ }
+
+ @keyframes fade {
+ 0% {
+ opacity: 0;
+ max-height: 0;
+ }
+
+ 100% {
+ opacity: 1;
+ max-height: 200px;
+ }
+ }
+}
+
+
+/*
+Horizontal Rule
+*/
+
+hr {
+ margin-top: var(--spacing-large);
+ margin-bottom: var(--spacing-large);
+ height: 1px;
+ background-color: var(--separator-color);
+ border: 0;
+}
+
+.contents hr {
+ box-shadow: 100px 0 0 var(--separator-color),
+ -100px 0 0 var(--separator-color),
+ 500px 0 0 var(--separator-color),
+ -500px 0 0 var(--separator-color),
+ 1500px 0 0 var(--separator-color),
+ -1500px 0 0 var(--separator-color),
+ 2000px 0 0 var(--separator-color),
+ -2000px 0 0 var(--separator-color);
+}
+
+.contents img, .contents .center, .contents center, .contents div.image object {
+ max-width: 100%;
+ overflow: auto;
+}
+
+@media screen and (max-width: 767px) {
+ .contents .dyncontent > .center, .contents > center {
+ margin-left: calc(0px - var(--spacing-large));
+ margin-right: calc(0px - var(--spacing-large));
+ max-width: calc(100% + 2 * var(--spacing-large));
+ }
+}
+
+/*
+Directories
+*/
+div.directory {
+ border-top: 1px solid var(--separator-color);
+ border-bottom: 1px solid var(--separator-color);
+ width: auto;
+}
+
+table.directory {
+ font-family: var(--font-family);
+ font-size: var(--page-font-size);
+ font-weight: normal;
+ width: 100%;
+}
+
+table.directory td.entry, table.directory td.desc {
+ padding: calc(var(--spacing-small) / 2) var(--spacing-small);
+ line-height: var(--table-line-height);
+}
+
+table.directory tr.even td:last-child {
+ border-radius: 0 var(--border-radius-small) var(--border-radius-small) 0;
+}
+
+table.directory tr.even td:first-child {
+ border-radius: var(--border-radius-small) 0 0 var(--border-radius-small);
+}
+
+table.directory tr.even:last-child td:last-child {
+ border-radius: 0 var(--border-radius-small) 0 0;
+}
+
+table.directory tr.even:last-child td:first-child {
+ border-radius: var(--border-radius-small) 0 0 0;
+}
+
+table.directory td.desc {
+ min-width: 250px;
+}
+
+table.directory tr.even {
+ background-color: var(--odd-color);
+}
+
+table.directory tr.odd {
+ background-color: transparent;
+}
+
+.icona {
+ width: auto;
+ height: auto;
+ margin: 0 var(--spacing-small);
+}
+
+.icon {
+ background: var(--primary-color);
+ border-radius: var(--border-radius-small);
+ font-size: var(--page-font-size);
+ padding: calc(var(--page-font-size) / 5);
+ line-height: var(--page-font-size);
+ transform: scale(0.8);
+ height: auto;
+ width: var(--page-font-size);
+ user-select: none;
+}
+
+.iconfopen, .icondoc, .iconfclosed {
+ background-position: center;
+ margin-bottom: 0;
+ height: var(--table-line-height);
+}
+
+.icondoc {
+ filter: saturate(0.2);
+}
+
+@media screen and (max-width: 767px) {
+ div.directory {
+ margin-left: calc(0px - var(--spacing-large));
+ margin-right: calc(0px - var(--spacing-large));
+ }
+}
+
+@media (prefers-color-scheme: dark) {
+ html:not(.light-mode) .iconfopen, html:not(.light-mode) .iconfclosed {
+ filter: hue-rotate(180deg) invert();
+ }
+}
+
+html.dark-mode .iconfopen, html.dark-mode .iconfclosed {
+ filter: hue-rotate(180deg) invert();
+}
+
+/*
+Class list
+*/
+
+.classindex dl.odd {
+ background: var(--odd-color);
+ border-radius: var(--border-radius-small);
+}
+
+.classindex dl.even {
+ background-color: transparent;
+}
+
+/*
+Class Index Doxygen 1.8
+*/
+
+table.classindex {
+ margin-left: 0;
+ margin-right: 0;
+ width: 100%;
+}
+
+table.classindex table div.ah {
+ background-image: none;
+ background-color: initial;
+ border-color: var(--separator-color);
+ color: var(--page-foreground-color);
+ box-shadow: var(--box-shadow);
+ border-radius: var(--border-radius-large);
+ padding: var(--spacing-small);
+}
+
+div.qindex {
+ background-color: var(--odd-color);
+ border-radius: var(--border-radius-small);
+ border: 1px solid var(--separator-color);
+ padding: var(--spacing-small) 0;
+}
+
+/*
+Footer and nav-path
+*/
+
+#nav-path {
+ width: 100%;
+}
+
+#nav-path ul {
+ background-image: none;
+ background: var(--page-background-color);
+ border: none;
+ border-top: 1px solid var(--separator-color);
+ border-bottom: 1px solid var(--separator-color);
+ border-bottom: 0;
+ box-shadow: 0 0.75px 0 var(--separator-color);
+ font-size: var(--navigation-font-size);
+}
+
+img.footer {
+ width: 60px;
+}
+
+.navpath li.footer {
+ color: var(--page-secondary-foreground-color);
+}
+
+address.footer {
+ color: var(--page-secondary-foreground-color);
+ margin-bottom: var(--spacing-large);
+}
+
+#nav-path li.navelem {
+ background-image: none;
+ display: flex;
+ align-items: center;
+}
+
+.navpath li.navelem a {
+ text-shadow: none;
+ display: inline-block;
+ color: var(--primary-color) !important;
+}
+
+.navpath li.navelem b {
+ color: var(--primary-dark-color);
+ font-weight: 500;
+}
+
+li.navelem {
+ padding: 0;
+ margin-left: -8px;
+}
+
+li.navelem:first-child {
+ margin-left: var(--spacing-large);
+}
+
+li.navelem:first-child:before {
+ display: none;
+}
+
+#nav-path li.navelem:after {
+ content: '';
+ border: 5px solid var(--page-background-color);
+ border-bottom-color: transparent;
+ border-right-color: transparent;
+ border-top-color: transparent;
+ transform: translateY(-1px) scaleY(4.2);
+ z-index: 10;
+ margin-left: 6px;
+}
+
+#nav-path li.navelem:before {
+ content: '';
+ border: 5px solid var(--separator-color);
+ border-bottom-color: transparent;
+ border-right-color: transparent;
+ border-top-color: transparent;
+ transform: translateY(-1px) scaleY(3.2);
+ margin-right: var(--spacing-small);
+}
+
+.navpath li.navelem a:hover {
+ color: var(--primary-color);
+}
+
+/*
+Scrollbars for Webkit
+*/
+
+#nav-tree::-webkit-scrollbar,
+div.fragment::-webkit-scrollbar,
+pre.fragment::-webkit-scrollbar,
+div.memproto::-webkit-scrollbar,
+.contents center::-webkit-scrollbar,
+.contents .center::-webkit-scrollbar,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody::-webkit-scrollbar,
+div.contents .toc::-webkit-scrollbar {
+ background: transparent;
+ width: calc(var(--webkit-scrollbar-size) + var(--webkit-scrollbar-padding) + var(--webkit-scrollbar-padding));
+ height: calc(var(--webkit-scrollbar-size) + var(--webkit-scrollbar-padding) + var(--webkit-scrollbar-padding));
+}
+
+#nav-tree::-webkit-scrollbar-thumb,
+div.fragment::-webkit-scrollbar-thumb,
+pre.fragment::-webkit-scrollbar-thumb,
+div.memproto::-webkit-scrollbar-thumb,
+.contents center::-webkit-scrollbar-thumb,
+.contents .center::-webkit-scrollbar-thumb,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody::-webkit-scrollbar-thumb,
+div.contents .toc::-webkit-scrollbar-thumb {
+ background-color: transparent;
+ border: var(--webkit-scrollbar-padding) solid transparent;
+ border-radius: calc(var(--webkit-scrollbar-padding) + var(--webkit-scrollbar-padding));
+ background-clip: padding-box;
+}
+
+#nav-tree:hover::-webkit-scrollbar-thumb,
+div.fragment:hover::-webkit-scrollbar-thumb,
+pre.fragment:hover::-webkit-scrollbar-thumb,
+div.memproto:hover::-webkit-scrollbar-thumb,
+.contents center:hover::-webkit-scrollbar-thumb,
+.contents .center:hover::-webkit-scrollbar-thumb,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody:hover::-webkit-scrollbar-thumb,
+div.contents .toc:hover::-webkit-scrollbar-thumb {
+ background-color: var(--webkit-scrollbar-color);
+}
+
+#nav-tree::-webkit-scrollbar-track,
+div.fragment::-webkit-scrollbar-track,
+pre.fragment::-webkit-scrollbar-track,
+div.memproto::-webkit-scrollbar-track,
+.contents center::-webkit-scrollbar-track,
+.contents .center::-webkit-scrollbar-track,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody::-webkit-scrollbar-track,
+div.contents .toc::-webkit-scrollbar-track {
+ background: transparent;
+}
+
+#nav-tree::-webkit-scrollbar-corner {
+ background-color: var(--side-nav-background);
+}
+
+#nav-tree,
+div.fragment,
+pre.fragment,
+div.memproto,
+.contents center,
+.contents .center,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody,
+div.contents .toc {
+ overflow-x: auto;
+ overflow-x: overlay;
+}
+
+#nav-tree {
+ overflow-x: auto;
+ overflow-y: auto;
+ overflow-y: overlay;
+}
+
+/*
+Scrollbars for Firefox
+*/
+
+#nav-tree,
+div.fragment,
+pre.fragment,
+div.memproto,
+.contents center,
+.contents .center,
+.contents table:not(.memberdecls):not(.mlabels):not(.fieldtable):not(.memname) tbody,
+div.contents .toc {
+ scrollbar-width: thin;
+}
+
+/*
+Optional Dark mode toggle button
+*/
+
+doxygen-awesome-dark-mode-toggle {
+ display: inline-block;
+ margin: 0 0 0 var(--spacing-small);
+ padding: 0;
+ width: var(--searchbar-height);
+ height: var(--searchbar-height);
+ background: none;
+ border: none;
+ border-radius: var(--searchbar-height);
+ vertical-align: middle;
+ text-align: center;
+ line-height: var(--searchbar-height);
+ font-size: 22px;
+ display: flex;
+ align-items: center;
+ justify-content: center;
+ user-select: none;
+ cursor: pointer;
+}
+
+doxygen-awesome-dark-mode-toggle > svg {
+ transition: transform .1s ease-in-out;
+}
+
+doxygen-awesome-dark-mode-toggle:active > svg {
+ transform: scale(.5);
+}
+
+doxygen-awesome-dark-mode-toggle:hover {
+ background-color: rgba(0,0,0,.03);
+}
+
+html.dark-mode doxygen-awesome-dark-mode-toggle:hover {
+ background-color: rgba(0,0,0,.18);
+}
+
+/*
+Optional fragment copy button
+*/
+.doxygen-awesome-fragment-wrapper {
+ position: relative;
+}
+
+doxygen-awesome-fragment-copy-button {
+ opacity: 0;
+ background: var(--fragment-background);
+ width: 28px;
+ height: 28px;
+ position: absolute;
+ right: calc(var(--spacing-large) - (var(--spacing-large) / 2.5));
+ top: calc(var(--spacing-large) - (var(--spacing-large) / 2.5));
+ border: 1px solid var(--fragment-foreground);
+ cursor: pointer;
+ border-radius: var(--border-radius-small);
+ display: flex;
+ justify-content: center;
+ align-items: center;
+}
+
+.doxygen-awesome-fragment-wrapper:hover doxygen-awesome-fragment-copy-button, doxygen-awesome-fragment-copy-button.success {
+ opacity: .28;
+}
+
+doxygen-awesome-fragment-copy-button:hover, doxygen-awesome-fragment-copy-button.success {
+ opacity: 1 !important;
+}
+
+doxygen-awesome-fragment-copy-button:active:not([class~=success]) svg {
+ transform: scale(.91);
+}
+
+doxygen-awesome-fragment-copy-button svg {
+ fill: var(--fragment-foreground);
+ width: 18px;
+ height: 18px;
+}
+
+doxygen-awesome-fragment-copy-button.success svg {
+ fill: rgb(14, 168, 14);
+}
+
+doxygen-awesome-fragment-copy-button.success {
+ border-color: rgb(14, 168, 14);
+}
+
+@media screen and (max-width: 767px) {
+ .textblock > .doxygen-awesome-fragment-wrapper > doxygen-awesome-fragment-copy-button,
+ .textblock li > .doxygen-awesome-fragment-wrapper > doxygen-awesome-fragment-copy-button,
+ .memdoc li > .doxygen-awesome-fragment-wrapper > doxygen-awesome-fragment-copy-button,
+ .memdoc > .doxygen-awesome-fragment-wrapper > doxygen-awesome-fragment-copy-button,
+ dl dd > .doxygen-awesome-fragment-wrapper > doxygen-awesome-fragment-copy-button {
+ right: 0;
+ }
+}
+
+/*
+Optional paragraph link button
+*/
+
+a.anchorlink {
+ font-size: 90%;
+ margin-left: var(--spacing-small);
+ color: var(--page-foreground-color) !important;
+ text-decoration: none;
+ opacity: .15;
+ display: none;
+ transition: opacity .1s ease-in-out, color .1s ease-in-out;
+}
+
+a.anchorlink svg {
+ fill: var(--page-foreground-color);
+}
+
+h3 a.anchorlink svg, h4 a.anchorlink svg {
+ margin-bottom: -3px;
+ margin-top: -4px;
+}
+
+a.anchorlink:hover {
+ opacity: .45;
+}
+
+h2:hover a.anchorlink, h1:hover a.anchorlink, h3:hover a.anchorlink, h4:hover a.anchorlink {
+ display: inline-block;
+}
\ No newline at end of file
diff --git a/docs/images/nvblox_logo.png b/docs/images/nvblox_logo.png
new file mode 100644
index 00000000..9cf09f40
Binary files /dev/null and b/docs/images/nvblox_logo.png differ
diff --git a/docs/images/nvblox_logo_128.png b/docs/images/nvblox_logo_128.png
new file mode 100644
index 00000000..c2961b13
Binary files /dev/null and b/docs/images/nvblox_logo_128.png differ
diff --git a/docs/images/nvblox_logo_32.png b/docs/images/nvblox_logo_32.png
new file mode 100644
index 00000000..738a25af
Binary files /dev/null and b/docs/images/nvblox_logo_32.png differ
diff --git a/docs/images/nvblox_logo_64.png b/docs/images/nvblox_logo_64.png
new file mode 100644
index 00000000..69b0b0e2
Binary files /dev/null and b/docs/images/nvblox_logo_64.png differ
diff --git a/docs/images/system_diagram.png b/docs/images/system_diagram.png
new file mode 100644
index 00000000..abd39567
Binary files /dev/null and b/docs/images/system_diagram.png differ
diff --git a/docs/pages/integrators.md b/docs/pages/integrators.md
new file mode 100644
index 00000000..b33d6d94
--- /dev/null
+++ b/docs/pages/integrators.md
@@ -0,0 +1,18 @@
+# Integrators
+
+An integrator is, very generally, a class that modifies the content of a layers.
+
+The integrators currently offered can be split into two types:
+
+* Those which fuse incoming sensor data into a layer. For example, the [TsdfIntegrator](@ref nvblox::ProjectiveTsdfIntegrator) and [ColorIntegrator](@ref nvblox::ProjectiveColorIntegrator) fused depth images and color images into [TsdfLayer](@ref nvblox::TsdfLayer) and [ColorLayer](@ref nvblox::ColorLayer) respectively.
+* Those which transform the contents of one layer to update the data in another layer. For example the [EsdfIntegrator](@ref nvblox::EsdfIntegrator) transforms an TSDF in a [TsdfLayer](@ref nvblox::TsdfLayer) into an ESDF in a [EsdfLayer](@ref nvblox::EsdfLayer).
+
+## API
+
+The API for the currently available integrators are here:
+
+* [TsdfIntegrator](@ref nvblox::ProjectiveTsdfIntegrator)
+* [ColorIntegrator](@ref nvblox::ProjectiveColorIntegrator)
+* [EsdfIntegrator](@ref nvblox::EsdfIntegrator)
+* [MeshIntegrator](@ref nvblox::MeshIntegrator)
+
diff --git a/docs/pages/map.md b/docs/pages/map.md
new file mode 100644
index 00000000..188b3098
--- /dev/null
+++ b/docs/pages/map.md
@@ -0,0 +1,53 @@
+# Map Overview
+We implement a hierarchical sparse voxel grid for storing data. At the top level we have the ``LayerCake``, which contains several layers, each of which contains a different type of mapped quantity (eg TSDF and ESDF). A layer is a collection of sparsely allocated blocks. Each block is in charge of mapping a cubular small region of space. Most blocks are composed of many voxels, each of which captures a single value of the mapped quantity (eg the TSDF).
+
+
+
+The API for the various classes implementing a map nvblox map:
+
+## Voxels
+
+* [TsdfVoxel](@ref nvblox::TsdfVoxel)
+* [EsdfVoxel](@ref nvblox::EsdfVoxel)
+* [ColorVoxel](@ref nvblox::ColorVoxel)
+
+## Blocks
+
+A template for a block containing voxels.
+
+* [VoxelBlock](@ref nvblox::VoxelBlock)
+
+Blocks containing specific voxel types. These are just typedefs of the above voxel blocks.
+
+* [TsdfBlock](@ref nvblox::TsdfBlock)
+* [Esdflock](@ref nvblox::TsdfBlock)
+* [ColorBlock](@ref nvblox::TsdfBlock)
+
+We store the mesh of the environment as a collection of smaller meshes. In particular, a ``MeshLayer`` contains ``MeshBlocks``, each of which contains a mesh of the surfaces in that region contained within that block. Note that we store a mesh per *block*, there are no mesh *voxels*.
+Storage of the mesh in this way allows us to perform incremental updates (we only update the mesh in blocks where the underlying TSDF has changed).
+
+* [MeshBlock](@ref nvblox::MeshBlock>)
+
+## Layers
+
+Template classes for layers containing blocks, and the child class for layers containing ``VoxelBlock`` s.
+
+* [BlockLayer](@ref nvblox::BlockLayer)
+* [VoxelBlockLayer](@ref nvblox::VoxelBlockLayer)
+
+Typedefs for layers containing specific block types and voxel block types.
+
+* [TsdfLayer](@ref nvblox::TsdfLayer)
+* [EsdfLayer](@ref nvblox::EsdfLayer)
+* [ColorLayer](@ref nvblox::ColorLayer)
+* [MeshLayer](@ref nvblox::MeshLayer)
+
+## Layer Cake
+
+The ``LayerCake`` combines several layers into a single object.
+
+* [LayerCake](@ref nvblox::LayerCake)
+
+
+
+
diff --git a/docs/pages/nvblox_node.rst b/docs/pages/nvblox_node.rst
deleted file mode 100644
index 7bb2e792..00000000
--- a/docs/pages/nvblox_node.rst
+++ /dev/null
@@ -1,158 +0,0 @@
-
-.. # Package Reference
-.. In this section we give an overview of parameters, inputs, and outputs within this repository and their purpose and settings.
-
-.. ## nvblox_ros
-.. The package centers around the `nvblox_node`, whose parameters, inputs, and outputs are described in detail here.
-
-
-.. ### Published Topics
-
-.. **mesh** ``nvblox_msgs::msg::Mesh`` \
-.. A visualization topic showing the mesh produced from the TSDF in a form that can be seen in RViz using `nvblox_rviz_plugin`. Set ``max_mesh_update_hz`` to control its update rate.
-
-.. **pointcloud** ``sensor_msgs::msg::PointCloud2`` \
-.. A pointcloud of the 2D ESDF (Euclidean Signed Distance Field), with intensity as the metric distance to the nearest obstacle. Set ``max_esdf_update_hz`` to control its update rate.
-
-.. **map_slice** ``nvblox_msgs::msg::DistanceMapSlice`` \
-.. A 2D slice of the ESDF, to be consumed by `nvblox_nav2` package for interfacing with Nav2. Set ``max_esdf_update_hz`` to control its update rate.
-
-.. ### Subscribed Topics
-
-.. **tf** ``tf2_msgs::msg::TFMessage`` \
-.. The pose of the sensor relative to the global frame is resolved through TF2. Please see the [ROS2 documentation](https://docs.ros.org/en/foxy/Tutorials/Tf2/Introduction-To-Tf2.html) for more details.
-
-.. **depth/image** ``sensor_msgs::msg::Image`` \
-.. The input depth image to be integrated. Must be paired with a `camera_info` message below. Supports both floating-point (depth in meters) and uint16 (depth in millimeters, OpenNI format).
-
-.. **depth/camera_info** ``sensor_msgs::msg::CameraInfo`` \
-.. Required topic along with the depth image; contains intrinsic calibration parameters of the depth camera.
-
-.. **color/image** ``sensor_msgs::msg::Image`` \
-.. Optional input color image to be integrated. Must be paired with a `camera_info` message below. Only used to color the mesh.
-
-.. **color/camera_info** ``sensor_msgs::msg::CameraInfo`` \
-.. Optional topic along with the color image above, contains intrinsics of the color camera.
-
-.. ### Services
-
-.. **save_ply** ``std_srvs::srv::Empty`` \
-.. This service has an empty request and response. Calling this service will write a mesh to disk at the path specified by the `output_dir` parameter.
-
-.. ### Parameters
-
-.. A summary of the user settable parameters. All parameters are listed as:
-
-.. **Parameter** `Default` \
-.. Description.
-
-.. #### General Parameters
-
-.. **voxel_size** `0.05` \
-.. Voxel size (in meters) to use for the map.
-
-.. **esdf** `true` \
-.. Whether to compute the ESDF (map of distances to the nearest obstacle).
-
-.. **esdf_2d** `true` \
-.. Whether to compute the ESDF in 2D (true) or 3D (false).
-
-.. **distance_slice** `true` \
-.. Whether to output a distance slice of the ESDF to be used for path planning.
-
-.. **mesh** `true` \
-.. Whether to output a mesh for visualization in rviz, to be used with `nvblox_rviz_plugin`.
-
-.. **global_frame** `map` \
-.. The name of the TF frame to be used as the global frame.
-
-.. **slice_height** `1.0` \
-.. The *output* slice height for the distance slice and ESDF pointcloud. Does not need to be within min and max height below. In units of meters.
-
-.. **min_height** `0.0` \
-.. The minimum height, in meters, to consider obstacles part of the 2D ESDF slice.
-
-.. **max_height** `1.0` \
-.. The maximum height, in meters, to consider obstacles part of the 2D ESDF slice.
-
-.. **max_tsdf_update_hz** `10.0` \
-.. The maximum rate (in Hz) at which to integrate depth images into the TSDF. A value of 0.0 means there is no cap.
-
-.. **max_color_update_hz** `5.0` \
-.. The maximum rate (in Hz) at which to integrate color images into the color layer. A value of 0.0 means there is no cap.
-
-.. **max_mesh_update_hz** `5.0` \
-.. The maximum rate (in Hz) at which to update and color the mesh. A value of 0.0 means there is no cap.
-
-.. **max_esdf_update_hz** `2.0` \
-.. The maximum rate (in Hz) at which to update the ESDF and output the distance slice. A value of 0.0 means there is no cap.
-
-
-.. #### Integrator Settings
-.. **tsdf_integrator_max_integration_distance_m** `10.0` \
-.. The maximum distance, in meters, to integrate the TSDF up to.
-
-.. **tsdf_integrator_truncation_distance_vox** `4.0` \
-.. The truncation distance, in units of voxels, for the TSDF.
-
-.. **tsdf_integrator_max_weight** `100.0` \
-.. Maximum weight for the TSDF. Setting this number higher will lead to higher-quality reconstructions but worse performance in dynamic scenes.
-
-.. **mesh_integrator_min_weight** `1e-4` \
-.. Minimum weight of the TSDF to consider for inclusion in the mesh.
-
-.. **mesh_integrator_weld_vertices** `false` \
-.. Whether to weld identical vertices together in the mesh. Currently reduces the number of vertices by a factor of 5x, but is quite slow so we do not recommend you use this setting.
-
-.. **color_integrator_max_integration_distance_m** `10.0` \
-.. Maximum distance, in meters, to integrate the color up to.
-
-.. **esdf_integrator_min_weight** `1e-4` \
-.. Minimum weight of the TSDF to consider for inclusion in the ESDF.
-
-.. **esdf_integrator_min_site_distance_vox** `1.0` \
-.. Minimum distance to consider a voxel within a surface for the ESDF calculation.
-
-.. **esdf_integrator_max_distance_m** `10.0` \
-.. Maximum distance to compute the ESDF up to, in meters.
-
-
-.. ## nvblox_nav2
-.. `nvblox_nav2` consists of two parts: an `nvblox_costmap_layer`, which is a Nav2 costmap plugin, and launch files for running the set-up with Nav2 in the loop (described above).
-
-.. ### `nvblox_costmap_layer` Parameters
-.. **nvblox_map_slice_topic** `/nvblox_node/map_slice` \
-.. Topic to listen for the slice (set via parameters to allow easy configuration from a parameter file).
-
-.. **max_obstacle_distance** `1.0` \
-.. Maximum distance from the surface to consider something to have a collision cost, in meters. This is *NOT* in addition to the inflation distance, but the total.
-
-.. **inflation_distance** `0.5` \
-.. Distance to inflate all obstacles by, in meters.
-
-.. # Troubleshooting
-
-.. ## Reconstruction without the planner (troubleshooting)
-.. If something isn't working, or as a quick sanity check, you can run reconstruction without the planner.
-
-.. Use the `carter_sim.launch.py` launch file parameters `run\_rviz` and `run\_nav2` to turn on and off running rviz and nav2. To run the reconstruction without the planner (but with rviz), run the following:
-.. ```
-.. ros2 launch nvblox_nav2 carter_sim.launch.py run_rviz:=true run_nav2:=false
-.. ```
-
-.. Now, command the robot to spin in place. In another terminal, source your ROS2 workspace again and enter:
-.. ```
-.. ros2 topic pub /cmd_vel geometry_msgs/Twist '{linear: {x: 0.0, y: 0.0, z: 0.0}, angular: {x: 0.0,y: 0.0,z: 0.1}}'
-.. ```
-.. You should see the robot spin in a circle and reconstruct its environment.
-
-
-.. ## Do you have a connection to Isaac Sim?
-.. A quick way to check if ROS2 is communicating correctly with Isaac Sim is to check whether the depth images are being sent.
-
-.. ```
-.. ros2 topic hz /left/depth
-.. ```
-
-.. If this does not receive any messages, then something is wrong with ROS2's connection to Isaac Sim. Either the file hasn't been set up correctly (make sure to run `nvblox_isaac_sim/omniverse_scripts/carter_warehouse.py` rather than opening an OV scene manually), or the scene is paused, or there are different ROS2 Domain IDs at play.
-
diff --git a/docs/pages/technical.md b/docs/pages/technical.md
new file mode 100644
index 00000000..c918f260
--- /dev/null
+++ b/docs/pages/technical.md
@@ -0,0 +1,26 @@
+# Technical Details
+
+## Input/Outputs
+
+Here we discuss the inputs you have to provide to nvblox, and the outputs it produces for downstream tasks. This is the default setup within ROS2 for 2D navigation, but note that other outputs are possible (such as the full 3D distance map).
+
+_Inputs_:
+* **Depth Images**: (@ref nvblox::Image) We require input from a sensor supplying depth per pixel. Examples of such sensors are the Intel Realsense series and Kinect cameras.
+* **Camera Intrinsics**: (@ref nvblox::Camera) The intrinsics associated with the camera that generated the depth (or color) image.
+* **Sensor Pose**: (@ref nvblox::Transform) We require localization of the depth sensor as input to nvblox
+* [Optional] **Color Images**: (@ref nvblox::Image) These can be used to color the reconstruction for visualization.
+* [Optional] **LIDAR**: (@ref nvblox::Lidar) Optionally or instead of Depth images, we can also consume LIDAR pointclouds which are reprojected to a cylindrical depth image based on the LIDAR intrinsics specified in this datatype.
+
+_Outputs_:
+* **Distance map slice**: (@ref nvblox::EsdfLayer) A 2D map that expresses the distance at each point from objects reconstructed in the environment. This is typically used by a planner to compute a collision cost map.
+* **Mesh**: (@ref nvblox::MeshLayer) We output a mesh for visualization in RVIZ.
+
+The figure below shows a simple system utilizing nvblox for path planning.
+
+
+
+## (Brief) Technical Details
+
+Nvblox builds the reconstructed map in the form of a Truncated Signed Distance Function (TSDF) stored in a 3D voxel grid. This approach is similar to 3D occupancy grid mapping approaches in which occupancy probabilities are stored at each voxel. In contrast however, TSDF-based approaches (like nvblox) store the (signed) distance to the closest surface at each voxel. The surface of the environment can then be extracted as the zero-level set of this voxelized function. Typically TSDF-based reconstructions provide higher quality surface reconstructions.
+
+In addition to their use in reconstruction, distance fields are also useful for path planning because they provide an immediate means of checking whether potential future robot positions are in collision. This fact, the utility of distance functions for both reconstruction and planning, motivates their use in nvblox (a reconstruction library for path planning).
diff --git a/docs/pages/technical.rst b/docs/pages/technical.rst
deleted file mode 100644
index ab2be9b6..00000000
--- a/docs/pages/technical.rst
+++ /dev/null
@@ -1,25 +0,0 @@
-# Input/Outputs
-
-Here we discuss the inputs you have to provide to nvblox, and the outputs it produces for downstream tasks.
-
-Inputs:
-* **Depth Images**: We require input from a sensor supplying depth per pixel. Examples of such sensors are the Intel Realsense series and Kinect cameras.
-* **Sensor Pose**: We require localization of the depth sensor as input to nvblox
-* [Optional] **Color Images**: These can be used to color the reconstruction for visualization.
-
-Outputs:
-* **Distance map slice**: A 2D map that expresses the distance at each point from objects reconstructed in the environment. This is typically used by a planner to compute a collision cost map.
-* **Mesh**: We output a mesh for visualization in RVIZ.
-
-The figure below shows a simple system utilizing nvblox for path planning.
-
-
-
-
-
-
-# (Brief) Technical Details
-
-Nvblox builds the reconstructed map in the form of a Truncated Signed Distance Function (TSDF) stored in a 3D voxel grid. This approach is similar to 3D occupancy grid mapping approaches in which occupancy probabilities are stored at each voxel. In contrast however, TSDF-based approaches (like nvblox) store the (signed) distance to the closest surface at each voxel. The surface of the environment can then be extracted as the zero-level set of this voxelized function. Typically TSDF-based reconstructions provide higher quality surface reconstructions.
-
-In addition to their use in reconstruction, distance fields are also useful for path planning because they provide an immediate means of checking whether potential future robot positions are in collision. This fact, the utility of distance functions for both reconstruction and planning, motivates their use in nvblox (a reconstruction library for path planning).
diff --git a/docs/redistributable.md b/docs/redistributable.md
new file mode 100644
index 00000000..cb345766
--- /dev/null
+++ b/docs/redistributable.md
@@ -0,0 +1,51 @@
+# Redistributable Building
+This will walk you through how to build a single libnvblox shared library with only static dependencies. The dependencies must be built manually with -fPIC (position independent code) and linked as static libraries.
+You generally shouldn't need to do this except for building nvblox to be included in Bazel for example; the exported targets should work fine for inclusion in CMake projects.
+
+## Step 1: Install dependencies
+**Sqlite3**:
+Download here: https://www.sqlite.org/2022/sqlite-autoconf-3390400.tar.gz
+
+Unzip & build using this command:
+```
+CFLAGS=-fPIC ./configure --prefix=/home/USERNAME/code/external/sqlite-autoconf-3390400/install/
+```
+
+**glog**:
+Download here: https://github.com/google/glog/archive/refs/tags/v0.4.0.zip
+
+Unzip & build using this command:
+```
+mkdir build install
+cd build
+cmake .. -DCMAKE_POSITION_INDEPENDENT_CODE=ON \
+-DCMAKE_INSTALL_PREFIX=/home/USERNAME/code/external/glog-0.4.0/install/ \
+-DWITH_GFLAGS=OFF -DBUILD_SHARED_LIBS=OFF \
+&& make -j8 && make install
+```
+
+**gflags**:
+Download here: https://github.com/gflags/gflags/archive/refs/tags/v2.2.2.zip
+
+Unzip & build using this command:
+```
+mkdir build install
+cd build
+cmake .. -DCMAKE_POSITION_INDEPENDENT_CODE=ON \
+-DCMAKE_INSTALL_PREFIX=/home/USERNAME/code/external/gflags-2.2.2/install/ \
+-DGFLAGS_BUILD_STATIC_LIBS=ON -DGFLAGS=google \
+&& make -j8 && make install
+```
+
+## Step 2: Build nvblox
+Build nvblox with the paths set to where you've unzipped and built the dependencies:
+```
+cd nvblox/
+mkdir build install
+cmake .. -DCMAKE_INSTALL_PREFIX=/home/helen/code/nvblox/nvblox/install/ \
+-DBUILD_FOR_ALL_ARCHS=TRUE -DBUILD_REDISTRIBUTABLE=TRUE \
+-DSQLITE3_BASE_PATH="/home/USERNAME/code/external/sqlite-autoconf-3390400/install/" \
+-DGLOG_BASE_PATH="/home/USERNAME/code/external/glog-0.4.0/install/" \
+-DGFLAGS_BASE_PATH="/home/USERNAME/code/external/gflags-2.2.2/install/" \
+&& make -j8 && make install
+```
\ No newline at end of file
diff --git a/docs/root.rst b/docs/root.rst
new file mode 100644
index 00000000..50cba87b
--- /dev/null
+++ b/docs/root.rst
@@ -0,0 +1,8 @@
+=======
+Table of Contents
+=======
+
+.. toctree::
+ :maxdepth: 1
+ :glob:
+ root
diff --git a/docs/pages/examples/core_example.rst b/docs/rst/examples/core_example.rst
similarity index 100%
rename from docs/pages/examples/core_example.rst
rename to docs/rst/examples/core_example.rst
diff --git a/docs/pages/examples/index.rst b/docs/rst/examples/index.rst
similarity index 100%
rename from docs/pages/examples/index.rst
rename to docs/rst/examples/index.rst
diff --git a/docs/pages/examples/ros_example.rst b/docs/rst/examples/ros_example.rst
similarity index 100%
rename from docs/pages/examples/ros_example.rst
rename to docs/rst/examples/ros_example.rst
diff --git a/docs/index.rst b/docs/rst/index.rst
similarity index 74%
rename from docs/index.rst
rename to docs/rst/index.rst
index 7e1f72d6..384ed59d 100644
--- a/docs/index.rst
+++ b/docs/rst/index.rst
@@ -1,5 +1,5 @@
=======
-nvblox
+Introduction to nvblox
=======
Nvblox is a package for building a 3D reconstruction of the environment around your robot from sensor observations in real-time. The reconstruction is intended to be used by path planners to generate collision-free paths. Under the hood, nvblox uses NVIDIA CUDA to accelerate this task to allow operation at real-time rates. This repository contains ROS2 integration for the nvblox core library.
@@ -12,27 +12,14 @@ Nvblox is a package for building a 3D reconstruction of the environment around y
.. |pic2| image:: /images/nvblox_navigation_trim.gif
:width: 45%
-**Left**: nvblox used for reconstruction on a scan from the `Sun3D Dataset
`_. **Right**: the nvblox ROS2 wrapper used to construct a costmap for `ROS2 Nav2
`_ for navigating of a robot inside `Isaac Sim
`_.
+**Left**: nvblox used for reconstruction on a scan from the `Sun3D Dataset `_.
+**Right**: the nvblox ROS2 wrapper used to construct a costmap for `ROS2 Nav2 `_ for navigating of a robot inside `Isaac Sim `_.
Nvblox is composed of two packages
* `nvblox Core Library `_ Contains the core C++/CUDA reconstruction library.
* `nvblox ROS2 Interface `_ Contains a ROS2 wrapper and integrations for simulation and path planning. Internally builds the core library.
-.. toctree::
- :maxdepth: 1
- :caption: Getting Started
-
- pages/installation/index
- pages/examples/index
-
- .. :glob:
-
- .. pages/*
-
- .. api/library_root
-
-
diff --git a/docs/pages/installation/core.rst b/docs/rst/installation/core.rst
similarity index 100%
rename from docs/pages/installation/core.rst
rename to docs/rst/installation/core.rst
diff --git a/docs/pages/installation/index.rst b/docs/rst/installation/index.rst
similarity index 100%
rename from docs/pages/installation/index.rst
rename to docs/rst/installation/index.rst
diff --git a/docs/pages/installation/ros.rst b/docs/rst/installation/ros.rst
similarity index 100%
rename from docs/pages/installation/ros.rst
rename to docs/rst/installation/ros.rst
diff --git a/docs/pages/integrators.rst b/docs/rst/integrators.rst
similarity index 100%
rename from docs/pages/integrators.rst
rename to docs/rst/integrators.rst
diff --git a/docs/pages/map.rst b/docs/rst/map.rst
similarity index 98%
rename from docs/pages/map.rst
rename to docs/rst/map.rst
index 5220da2b..bcc12d90 100644
--- a/docs/pages/map.rst
+++ b/docs/rst/map.rst
@@ -37,7 +37,7 @@ Storage of the mesh in this way allows us to perform incremental updates (we onl
Layers
======
-Template classes for layers containing blocks, and the child class for layers containing ``VoxelBlock``s.
+Template classes for layers containing blocks, and the child class for layers containing ``VoxelBlock`` s.
* :ref:`BlockLayer `
* :ref:`VoxelBlockLayer `
diff --git a/docs/pages/mapper.rst b/docs/rst/mapper.rst
similarity index 100%
rename from docs/pages/mapper.rst
rename to docs/rst/mapper.rst
diff --git a/docs/pages/rays.rst b/docs/rst/rays.rst
similarity index 100%
rename from docs/pages/rays.rst
rename to docs/rst/rays.rst
diff --git a/nvblox/CMakeLists.txt b/nvblox/CMakeLists.txt
index fab37aa3..9dc56a97 100644
--- a/nvblox/CMakeLists.txt
+++ b/nvblox/CMakeLists.txt
@@ -1,7 +1,33 @@
cmake_minimum_required(VERSION 3.16)
# set the project name and version
-project(nvblox VERSION 0.0.2 LANGUAGES CXX CUDA)
+project(nvblox VERSION 0.0.4 LANGUAGES CXX CUDA)
+
+########################
+# OPTIONS AND SETTINGS #
+########################
+# Build options
+option(BUILD_EXPERIMENTS "Build performance experimentation binaries" OFF)
+option(BUILD_TESTING "Build tests" ON)
+
+# Whether to build redistributable artifacts using static libraries. In this
+# case, the pre-build static libraries for stdgpu, glog, and gflags MUST be
+# specified, and built with -fPIC. (See README instructions).
+option(BUILD_REDISTRIBUTABLE "Whether to build redistributable artifacts,
+ linking in static versions of the dependent libraries" OFF)
+
+# These are only used if BUILD_REDISTRIBUTABLE is ON:
+# Set these from the command line by setting -DSQLITE3_BASE_PATH="YOUR/PATH/HERE" when calling cmake.
+set(SQLITE3_BASE_PATH "" CACHE STRING "Base path to static output for Sqlite3, build with -fPIC.")
+set(GLOG_BASE_PATH "" CACHE STRING "Base path to static output for glog, build with -fPIC.")
+set(GFLAGS_BASE_PATH "" CACHE STRING "Base path to static output for gflags, build with -fPIC.")
+
+# GPU architectures
+# By default this flag is NOT set. Cmake then detects the architecture of the build computer
+# and compiles for that architecture only. This can be an issue if you're building on one
+# machine, and running on machines with a different GPU achitecture. In this case, set the flag.
+# The penalty for doing this is increased build times.
+option(BUILD_FOR_ALL_ARCHS "Build for all GPU architectures" OFF)
# specify the C++ standard
set(CMAKE_CXX_STANDARD 14)
@@ -11,19 +37,14 @@ set(CMAKE_BUILD_TYPE RelWithDebInfo)
# This link flag replaces "runpath" with "rpath" in executables and shared objects.
# This is important because it means the search paths are passed down the shared object tree.
# https://stackoverflow.com/questions/58997230/cmake-project-fails-to-find-shared-library
-SET(nvblox_link_options "-Wl,--disable-new-dtags")
+set(nvblox_link_options "-Wl,--disable-new-dtags")
-# Build options
-option(BUILD_EXPERIMENTS "Build performance experimentation binaries" OFF)
-
-# GPU architectures
-# By default this flag is NOT set. Cmake then detects the architecture of the build computer
-# and compiles for that architecture only. This can be an issue if you're building on one
-# machine, and running on machines with a different GPU achitecture. In this case, set the flag.
-# The penalty for doing this is increased build times.
-option(BUILD_FOR_ALL_ARCHS "Build for all GPU architectures" OFF)
+# Cmake -fPIC flag.
+# NOTE(alexmillane): I needed to add this when I changed to linking against a static version
+# of stdgpu. Without is we get an error from thrust/cub linking.
+set(CMAKE_POSITION_INDEPENDENT_CODE TRUE)
-# Only used if the flag above is true.
+# Only used if the BUILD_FOR_ALL_ARCHS flag above is true.
if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.18)
set(CUDA_ARCHITECTURE_FLAGS "86;80;75;72;70;61;60")
else()
@@ -36,7 +57,6 @@ else()
-gencode=arch=compute_60,code=sm_60")
endif()
-
# Suppress spammy Eigen CUDA warnings.
# "expt-relaxed-constexpr" allows sharing constexpr between host and device code.
# "display_error_number" shows a warning number with all warnings, and the
@@ -45,6 +65,9 @@ endif()
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --disable-warnings --generate-line-info -lineinfo -Xcudafe --display_error_number")
set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} --compiler-options -fPIC")
+################
+# DEPENDENCIES #
+################
# Download thirdparty deps
message(STATUS "Downloading 3rdparty dependencies")
message(STATUS "Downloading Eigen")
@@ -58,41 +81,83 @@ include(thirdparty/stdgpu/stdgpu.cmake)
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
find_package(CUDA REQUIRED)
-# In the case of ROS builds, glog will likely be found at a higher level.
-# We want to link against that version in that case.
-if(NOT Glog_FOUND)
- find_package(Glog REQUIRED)
+# Treat redistributable and non-redistributable builds differently for
+# the other dependencies.
+if (BUILD_REDISTRIBUTABLE)
+ if (SQLITE3_BASE_PATH STREQUAL "")
+ message(FATAL_ERROR "No SQLite3 static libraries specified! Set the SQLITE3_BASE_PATH variable!")
+ endif()
+ if (GLOG_BASE_PATH STREQUAL "")
+ message(FATAL_ERROR "No Glog static libraries specified! Set the GLOG_BASE_PATH variable!")
+ endif()
+ if (GFLAGS_BASE_PATH STREQUAL "")
+ message(FATAL_ERROR "No gflags static libraries specified! Set the GFLAGS_BASE_PATH variable!")
+ endif()
+
+ # SQLite static lib:
+ set(SQLite3_LIBRARIES "${SQLITE3_BASE_PATH}/lib/libsqlite3.a")
+ set(SQLite3_INCLUDE_DIRS "${SQLITE3_BASE_PATH}/include/")
+
+ # GLOG STATIC LIB
+ set(GLOG_LIBRARIES "${GLOG_BASE_PATH}/lib/libglog.a")
+ set(GLOG_INCLUDE_DIRS "${GLOG_BASE_PATH}/include")
+
+ # GFLAGS STATIC LIB
+ set(gflags_LIBRARIES "${GFLAGS_BASE_PATH}/lib/libgflags.a")
+ set(gflags_INCLUDE_DIRS "${GFLAGS_BASE_PATH}/include")
+
+ # Set dependent settings, like turning testing off.
+ set(BUILD_TESTING OFF)
+else()
+ # By default, use system dependencies for these.
+ # In the case of ROS builds, glog will likely be found at a higher level.
+ # We want to link against that version in that case.
+ set(GLOG_PREFER_EXPORTED_GLOG_CMAKE_CONFIGURATION FALSE)
+ if(NOT Glog_FOUND)
+ find_package(Glog REQUIRED)
+ endif()
+ find_package(gflags REQUIRED)
+ find_package(SQLite3 REQUIRED)
+
+ set(gflags_INCLUDE_DIRS ${gflag_INCLUDE_DIR})
+ set(gflags_LIBRARIES "gflags")
endif()
-find_package(gflags REQUIRED)
-
-find_package(SQLite3 REQUIRED)
+############
+# INCLUDES #
+############
+message(INFO "SQLite3_LIBRARIES: ${SQLite3_LIBRARIES}\n"
+ "GLOG_LIBRARIES: ${GLOG_LIBRARIES}\n"
+ "gflags_LIBRARIES: ${gflags_LIBRARIES}\n"
+ "SQLite3_INCLUDE_DIRS: ${SQLite3_INCLUDE_DIRS}\n"
+ "GLOG_INCLUDE_DIRS: ${GLOG_INCLUDE_DIRS}\n"
+ "gflags_INCLUDE_DIRS: ${gflags_INCLUDE_DIRS}")
# Include dirs
include_directories(${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
include_directories(include)
+include_directories(${GLOG_INCLUDE_DIRS} ${gflags_INCLUDE_DIRS})
#############
# LIBRARIES #
#############
-add_library(nvblox_cuda_check SHARED
- src/core/cuda/error_check.cu
-)
-
-add_library(nvblox_gpu_hash SHARED
+# NOTE(alexmillane): nvblox (unfortunately) is split into two libraries, "nvblox_gpu_hash" which
+# wraps our interactions with stdgpu and "nvblox_lib" which only interacts with our gpu hash wrapper.
+# Ideally, I would have liked to have just a single object, however this caused run-time errors.
+# We compile the first *without* the separable code flag, and the second with. This is the only way
+# I've managed to get things working so far.
+add_library(nvblox_gpu_hash STATIC
src/gpu_hash/cuda/gpu_layer_view.cu
src/gpu_hash/cuda/gpu_set.cu
src/utils/timing.cpp
src/utils/nvtx_ranges.cpp
+ src/core/cuda/error_check.cu
)
add_dependencies(nvblox_gpu_hash nvblox_eigen stdgpu)
target_link_libraries(nvblox_gpu_hash PUBLIC
stdgpu
- nvblox_cuda_check
nvblox_eigen
${CUDA_nvToolsExt_LIBRARY}
- gflags
- ${GLOG_LIBRARIES}
)
target_link_options(nvblox_gpu_hash PUBLIC ${nvblox_link_options})
@@ -103,10 +168,12 @@ add_library(nvblox_lib SHARED
src/core/color.cpp
src/core/cuda/blox.cu
src/core/cuda/image_cuda.cu
+ src/core/cuda/layer.cu
src/core/cuda/warmup.cu
src/core/image.cpp
src/core/interpolation_3d.cpp
src/core/mapper.cpp
+ src/core/cuda/pointcloud.cu
src/integrators/cuda/view_calculator.cu
src/integrators/cuda/projective_tsdf_integrator.cu
src/integrators/cuda/projective_color_integrator.cu
@@ -119,6 +186,7 @@ add_library(nvblox_lib SHARED
src/io/mesh_io.cpp
src/io/ply_writer.cpp
src/io/layer_cake_io.cpp
+ src/io/pointcloud_io.cpp
src/mesh/marching_cubes.cu
src/mesh/mesh_block.cu
src/mesh/mesh_integrator_color.cu
@@ -132,15 +200,19 @@ add_library(nvblox_lib SHARED
src/serialization/serializer.cpp
src/serialization/sqlite_database.cpp
src/serialization/layer_type_register.cpp
+ src/semantics/cuda/image_masker.cu
+ src/semantics/cuda/image_projector.cu
+ src/core/cuda/error_check.cu
)
-target_link_libraries(nvblox_lib PUBLIC
+target_link_libraries(nvblox_lib
+ PUBLIC
${GLOG_LIBRARIES}
- gflags
+ ${gflags_LIBRARIES}
+ nvblox_eigen
${CUDA_LIBRARIES}
- ${CUDA_nvToolsExt_LIBRARY}
nvblox_gpu_hash
- nvblox_cuda_check
- nvblox_eigen
+ PRIVATE
+ ${CUDA_nvToolsExt_LIBRARY}
${SQLite3_LIBRARIES}
)
target_link_options(nvblox_lib PUBLIC ${nvblox_link_options})
@@ -192,23 +264,6 @@ if(BUILD_EXPERIMENTS)
add_subdirectory(experiments)
endif()
-#############################
-# INTERFACE LIBRARY FOR ROS #
-#############################
-# TODO: delete
-add_library(nvblox_interface INTERFACE)
-target_link_libraries(nvblox_interface INTERFACE
- nvblox_lib
- nvblox_gpu_hash
- nvblox_cuda_check
- ${GLOG_LIBRARIES}
- ${CUDA_LIBRARIES}
- Eigen3::Eigen
- gflags
- SQLite::SQLite3
-)
-target_include_directories(nvblox_interface INTERFACE include ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
-
##########
# EXPORT #
##########
@@ -217,7 +272,7 @@ include(GNUInstallDirs)
set_target_properties(stdgpu PROPERTIES INTERFACE_LINK_LIBRARIES "")
install(
- TARGETS nvblox_lib nvblox_gpu_hash nvblox_cuda_check nvblox_datasets stdgpu nvblox_eigen fuse_3dmatch fuse_replica
+ TARGETS nvblox_lib nvblox_gpu_hash nvblox_datasets stdgpu nvblox_eigen fuse_3dmatch fuse_replica
EXPORT nvbloxTargets
ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME}
LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}
@@ -229,12 +284,28 @@ install(
DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
)
+install(
+ DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/executables/include/
+ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
+)
install(
DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/eigen/include/eigen3
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
PATTERN "*/unsupported**" EXCLUDE
)
+if(BUILD_REDISTRIBUTABLE)
+ # Only re-distribute these headers if we're building the redistributable artifacts.
+ install(
+ DIRECTORY ${gflags_INCLUDE_DIRS}/gflags
+ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
+ )
+ install(
+ DIRECTORY ${GLOG_INCLUDE_DIRS}/glog
+ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
+ )
+endif()
+
include(CMakePackageConfigHelpers)
# generate the config file that is includes the exports
diff --git a/nvblox/evaluation/replica/evaluation_utils/voxel_grid.py b/nvblox/evaluation/replica/evaluation_utils/voxel_grid.py
index 14a2db26..3908a167 100644
--- a/nvblox/evaluation/replica/evaluation_utils/voxel_grid.py
+++ b/nvblox/evaluation/replica/evaluation_utils/voxel_grid.py
@@ -20,7 +20,7 @@
import numpy as np
import open3d as o3d
-from plyfile import PlyData
+from plyfile import PlyData, PlyElement
from matplotlib import pyplot as plt
from moviepy.editor import ImageSequenceClip
@@ -80,7 +80,7 @@ def get_valid_voxel_values(self) -> np.ndarray:
"""Get the centers of all valid voxels as an Nx3 numpy array
Returns:
- np.ndarray: Nx3 numpy array containing the center locations of valid voxels.
+ np.ndarray: Nx1 numpy array containing the center locations of valid voxels.
"""
return self.voxels[self.voxels != VoxelGrid.unobserved_sentinal]
@@ -102,6 +102,20 @@ def createFromPly(ply_path: Path) -> 'VoxelGrid':
str(ply_path)).elements[0]['intensity'])
return VoxelGrid.createFromSparseVoxels(sdf_pointcloud_xyz, sdf_pointcloud_values)
+ def writeToPly(self, ply_path: Path) -> None:
+ """Writes the ESDF as a pointcloud ply to file.
+
+ Args:
+ ply_path (Path): Path to the ply file to write
+ """
+ xyz = self.get_valid_voxel_centers()
+ distances = self.get_valid_voxel_values()
+ xyzi = np.hstack([xyz, distances.reshape((-1, 1))])
+ xyzi_structured = np.array([tuple(row) for row in xyzi], dtype=[(
+ 'x', 'f4'), ('y', 'f4'), ('z', 'f4'), ('intensity', 'f4')])
+ point_elements = PlyElement.describe(xyzi_structured, 'vertex')
+ PlyData([point_elements], text=True).write(str(ply_path))
+
@staticmethod
def createFromSparseVoxels(voxels_xyz: np.ndarray, voxel_values: np.ndarray) -> 'VoxelGrid':
"""Creates an VoxelGrid object from a list of valid voxels values and their locations.
@@ -135,6 +149,8 @@ def get_slice_mesh_at_ratio(self, slice_level_ratio: float, axis: str = 'x', cub
Args:
slice_level_ratio (float): Where to slice.
axis (str, optional): The axis to slice along. Defaults to 'x'.
+ cube_size (float, optional): Size of the mesh cube that will be used to
+ represent voxels. Given as a fraction of voxel size (i.e between 0.0 and 1.0).
Returns:
o3d.geometry.TriangleMesh: Mesh representing the slice.
@@ -156,6 +172,8 @@ def get_slice_mesh_at_index(self, slice_level_idx: int, axis: str = 'x', cube_si
Args:
slice_level_idx (int): The index to slice at.
axis (str, optional): The axis to slice along. Defaults to 'x'.
+ cube_size (float, optional): Size of the mesh cube that will be used to
+ represent voxels. Given as a fraction of voxel size (i.e between 0.0 and 1.0).
Returns:
o3d.geometry.TriangleMesh: Mesh representing the slice.
@@ -228,11 +246,12 @@ def __repr__(self) -> str:
return "VoxelGrid of voxels with shape: " + str(self.voxels.shape) \
+ " and " + str(self.num_valid_voxels()) + " valid voxels."
- def get_z_slice_animation_clip(self, mesh: o3d.geometry.TriangleMesh = None) -> ImageSequenceClip:
+ def get_z_slice_animation_clip(self, mesh: o3d.geometry.TriangleMesh = None, viewpoint: o3d.camera.PinholeCameraParameters = None) -> ImageSequenceClip:
"""Creates a image sequence containing horizontal slices moving through the z dimension of the VoxelGrid
Args:
mesh (o3d.geometry.TriangleMesh, optional): Additional mesh to add to the animation. Defaults to None.
+ viewpoint (o3d.camera.PinholeCameraParameters, optional): Viewpoint to record the slice from. Defaults to None.
Returns:
ImageSequenceClip: sequence of images of the slicing results
@@ -240,6 +259,9 @@ def get_z_slice_animation_clip(self, mesh: o3d.geometry.TriangleMesh = None) ->
images = []
vis = o3d.visualization.Visualizer()
vis.create_window()
+ if viewpoint is not None:
+ ctr = vis.get_view_control()
+ ctr.convert_from_pinhole_camera_parameters(viewpoint)
if mesh is not None:
vis.add_geometry(mesh)
slice_mesh = o3d.geometry.TriangleMesh()
@@ -255,6 +277,8 @@ def get_z_slice_animation_clip(self, mesh: o3d.geometry.TriangleMesh = None) ->
first = False
else:
vis.add_geometry(slice_mesh, reset_bounding_box=False)
+ if viewpoint is not None:
+ ctr.convert_from_pinhole_camera_parameters(viewpoint)
vis.poll_events()
vis.update_renderer()
image_float = np.asarray(vis.capture_screen_float_buffer())
diff --git a/nvblox/evaluation/replica/replica_esdf_evaluation.py b/nvblox/evaluation/replica/replica_esdf_evaluation.py
index 57c22b53..18b6942b 100755
--- a/nvblox/evaluation/replica/replica_esdf_evaluation.py
+++ b/nvblox/evaluation/replica/replica_esdf_evaluation.py
@@ -92,6 +92,7 @@ def evaluate_esdf(reconstructed_esdf_path: Path,
print('Calculating ESDF errors')
sdf_abs_diff = esdf_evaluation.get_sdf_abs_error_grid(
reconstructed_sdf, gt_sdf)
+ sdf_abs_diff.writeToPly(output_dir / 'error_esdf.ply')
abs_errors = sdf_abs_diff.get_valid_voxel_values()
# Statistics
@@ -117,7 +118,9 @@ def evaluate_esdf(reconstructed_esdf_path: Path,
json.dump(statistics_dict, statistics_file, indent=4)
# Error histogram.
- fig = px.histogram(sdf_abs_diff.get_valid_voxel_values())
+ sdf_diff_abs_np = sdf_abs_diff.get_valid_voxel_values()
+ np.savetxt(output_dir / 'esdf_errors.txt', sdf_diff_abs_np)
+ fig = px.histogram(sdf_diff_abs_np)
fig.write_image(output_dir / 'esdf_error_histogram.png')
fig.show()
diff --git a/nvblox/evaluation/replica/replica_reconstruction.py b/nvblox/evaluation/replica/replica_reconstruction.py
index 3579695f..795c21f8 100755
--- a/nvblox/evaluation/replica/replica_reconstruction.py
+++ b/nvblox/evaluation/replica/replica_reconstruction.py
@@ -22,6 +22,7 @@
from pathlib import Path
from typing import Tuple
import json
+import pandas as pd
import replica
import timing
@@ -29,7 +30,9 @@
def replica_reconstruction(dataset_path: Path,
output_root_path: Path = None,
- fuse_replica_binary_path: Path = None) -> Tuple[Path, Path]:
+ fuse_replica_binary_path: Path = None,
+ esdf_frame_subsampling : int = 1,
+ mesh_frame_subsampling : int = -1) -> Tuple[Path, Path]:
"""Builds a reconstruction for the replica dataset
Args:
@@ -43,6 +46,12 @@ def replica_reconstruction(dataset_path: Path,
fuse_replica_binary_path (Path, optional): The path to the binary which does the
fusion. Defaults to the build folder. Defaults to None.
+ esdf_frame_subsampling (int, optional): How often to compute the ESDF. We compute
+ Every N frames. Defaults to 1 (every frame).
+
+ mesh_frame_subsampling (int, optional): How often to compute the Mesh. We compute
+ Every N frames. Defaults to -1, which means compute once at the end.
+
Raises:
Exception: If the binary is not found.
@@ -71,21 +80,26 @@ def replica_reconstruction(dataset_path: Path,
esdf_output_path_flag = "--esdf_output_path"
timing_output_path_flag = "--timing_output_path"
esdf_frame_subsampling_flag = "--esdf_frame_subsampling"
+ mesh_frame_subsampling_flag = "--mesh_frame_subsampling"
subprocess.run([f"{fuse_replica_binary_path}", f"{dataset_path}",
mesh_output_path_flag, f"{reconstructed_mesh_path}",
esdf_output_path_flag, f"{reconstructed_esdf_path}",
timing_output_path_flag, f"{timing_path}",
- esdf_frame_subsampling_flag, f"{1}"])
+ esdf_frame_subsampling_flag, f"{esdf_frame_subsampling}",
+ mesh_frame_subsampling_flag, f"{mesh_frame_subsampling}"])
# Extract the means of the timers
timings_df = timing.get_timings_as_dataframe(timing_path)
means_series = timings_df['mean']
+ means_series.index = ['mean/' + row_name for row_name in means_series.index]
+ total_series = timings_df['total_time']
+ total_series.index = ['total/' + row_name for row_name in total_series.index]
# Write the results to a JSON
output_timings_path = output_dir / 'timing.json'
print(f"Writing the timings to: {output_timings_path}")
with open(output_timings_path, "w") as timings_file:
- json.dump(means_series.to_dict(), timings_file, indent=4)
+ json.dump(pd.concat([means_series, total_series]).to_dict(), timings_file, indent=4)
return reconstructed_mesh_path, reconstructed_esdf_path
@@ -93,8 +107,7 @@ def replica_reconstruction(dataset_path: Path,
if __name__ == '__main__':
parser = argparse.ArgumentParser(
- description="""Reconstruct a mesh from the replica dataset and test it
- against ground-truth geometry.""")
+ description="""Reconstruct a mesh from the replica dataset.""")
parser.add_argument("dataset_path", type=Path,
help="Path to the dataset root folder.")
diff --git a/nvblox/evaluation/replica/replica_surface_evaluation.py b/nvblox/evaluation/replica/replica_surface_evaluation.py
index ec545b2e..25065cbc 100755
--- a/nvblox/evaluation/replica/replica_surface_evaluation.py
+++ b/nvblox/evaluation/replica/replica_surface_evaluation.py
@@ -115,6 +115,11 @@ def evaluate_mesh(reconstructed_mesh_path: Path,
with open(output_statistics_path, "w") as statistics_file:
json.dump(statistics_dict, statistics_file, indent=4)
+ # Write raw errors to a file
+ errors_path = output_dir / 'surface_errors.txt'
+ print(f'Writing errors to: {errors_path}')
+ np.savetxt(errors_path, per_vertex_errors)
+
# Visualization
if do_error_visualization:
o3d.visualization.draw_geometries([error_mesh])
diff --git a/nvblox/examples/CMakeLists.txt b/nvblox/examples/CMakeLists.txt
index 4395744f..90808864 100644
--- a/nvblox/examples/CMakeLists.txt
+++ b/nvblox/examples/CMakeLists.txt
@@ -6,3 +6,13 @@ target_link_libraries(esdf_query
nvblox_lib
)
set_target_properties(esdf_query PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+# An example for map loading
+add_executable(load_map_and_mesh
+ src/load_map_and_mesh.cpp
+)
+target_link_libraries(load_map_and_mesh
+ nvblox_lib
+)
+
+set_target_properties(load_map_and_mesh PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
diff --git a/nvblox/examples/src/esdf_query.cu b/nvblox/examples/src/esdf_query.cu
index 9045dbc4..1e959279 100644
--- a/nvblox/examples/src/esdf_query.cu
+++ b/nvblox/examples/src/esdf_query.cu
@@ -14,6 +14,7 @@ See the License for the specific language governing permissions and
limitations under the License.
*/
#include
+#include
#include "nvblox/core/indexing.h"
#include "nvblox/core/layer.h"
diff --git a/nvblox/examples/src/load_map_and_mesh.cpp b/nvblox/examples/src/load_map_and_mesh.cpp
new file mode 100644
index 00000000..5186b31d
--- /dev/null
+++ b/nvblox/examples/src/load_map_and_mesh.cpp
@@ -0,0 +1,79 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+
+#include "gflags/gflags.h"
+#include "glog/logging.h"
+
+#include "nvblox/nvblox.h"
+
+#include "nvblox/serialization/layer_type_register.h"
+
+using namespace nvblox;
+
+int main(int argc, char* argv[]) {
+ gflags::ParseCommandLineFlags(&argc, &argv, true);
+ google::InitGoogleLogging(argv[0]);
+ FLAGS_alsologtostderr = true;
+ google::InstallFailureSignalHandler();
+
+ CHECK_GT(argc, 0) << "Please pass the path to the map to load.";
+ const std::string base_path = argv[1];
+ LOG(INFO) << "Loading map from " << base_path;
+
+ // Load
+ LayerCake cake = io::loadLayerCakeFromFile(base_path, MemoryType::kDevice);
+ LOG(INFO) << "Loaded cake with voxel size: " << cake.voxel_size();
+
+ // Print the names of the loaded layers.
+ LOG(INFO) << "Loaded layers:";
+ for (const auto& id_layer_pair : cake.get_layers()) {
+ LOG(INFO) << "Loaded: "
+ << LayerTypeRegister::getLayerName(id_layer_pair.first);
+ }
+
+ // Mesh
+ LOG(INFO) << "Meshing";
+ MeshIntegrator mesh_integrator;
+ MeshLayer mesh_layer(cake.block_size(), MemoryType::kDevice);
+ mesh_integrator.integrateMeshFromDistanceField(cake.get(),
+ &mesh_layer);
+ mesh_integrator.colorMesh(cake.get(), &mesh_layer);
+ LOG(INFO) << "Done";
+
+ std::string output_path;
+ if (argc > 2) {
+ output_path = argv[2];
+ } else {
+ output_path = "./mesh.ply";
+ }
+ LOG(INFO) << "Writing mesh to: " << output_path;
+ CHECK(io::outputMeshLayerToPly(mesh_layer, output_path));
+ LOG(INFO) << "Done";
+
+ // Esdf
+ if (cake.exists()) {
+ if (argc > 3) {
+ output_path = argv[3];
+ } else {
+ output_path = "./esdf.ply";
+ }
+ LOG(INFO) << "Writing esdf to: " << output_path;
+ CHECK(io::outputVoxelLayerToPly(cake.get(), output_path));
+ LOG(INFO) << "Done";
+ }
+
+ LOG(INFO) << "Finished running example.";
+}
diff --git a/nvblox/executables/CMakeLists.txt b/nvblox/executables/CMakeLists.txt
index 23cf07a9..16d1753f 100644
--- a/nvblox/executables/CMakeLists.txt
+++ b/nvblox/executables/CMakeLists.txt
@@ -3,6 +3,7 @@ add_library(nvblox_datasets SHARED
src/datasets/3dmatch.cpp
src/datasets/image_loader.cpp
src/datasets/replica.cpp
+ src/datasets/redwood.cpp
src/fuser.cpp
)
target_include_directories(nvblox_datasets PUBLIC
@@ -31,3 +32,12 @@ target_link_libraries(fuse_replica
nvblox_lib nvblox_datasets
)
set_target_properties(fuse_replica PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+# Redwood executable
+add_executable(fuse_redwood
+ src/fuse_redwood.cpp
+)
+target_link_libraries(fuse_redwood
+ nvblox_lib nvblox_datasets
+)
+set_target_properties(fuse_redwood PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
diff --git a/nvblox/executables/include/nvblox/datasets/3dmatch.h b/nvblox/executables/include/nvblox/datasets/3dmatch.h
index dd68918d..2086d29e 100644
--- a/nvblox/executables/include/nvblox/datasets/3dmatch.h
+++ b/nvblox/executables/include/nvblox/datasets/3dmatch.h
@@ -36,6 +36,7 @@ class DataLoader : public RgbdDataLoaderInterface {
public:
DataLoader(const std::string& base_path, const int seq_id,
bool multithreaded = true);
+ virtual ~DataLoader() = default;
/// Interface for a function that loads the next frames in a dataset
///@param[out] depth_frame_ptr The loaded depth frame.
diff --git a/nvblox/executables/include/nvblox/datasets/redwood.h b/nvblox/executables/include/nvblox/datasets/redwood.h
new file mode 100644
index 00000000..2ed1b342
--- /dev/null
+++ b/nvblox/executables/include/nvblox/datasets/redwood.h
@@ -0,0 +1,70 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+#include
+#include
+#include
+
+#include "nvblox/core/types.h"
+#include "nvblox/datasets/data_loader.h"
+#include "nvblox/datasets/image_loader.h"
+#include "nvblox/executables/fuser.h"
+
+namespace nvblox {
+namespace datasets {
+namespace redwood {
+
+// Build a Fuser for the Redwood dataset
+std::unique_ptr createFuser(const std::string base_path);
+
+///@brief A class for loading Redwood data
+class DataLoader : public RgbdDataLoaderInterface {
+ public:
+ DataLoader(const std::string& base_path, bool multithreaded = true);
+ virtual ~DataLoader() = default;
+
+ /// Interface for a function that loads the next frames in a dataset
+ ///@param[out] depth_frame_ptr The loaded depth frame.
+ ///@param[out] T_L_C_ptr Transform from Camera to the Layer frame.
+ ///@param[out] camera_ptr The intrinsic camera model.
+ ///@param[out] color_frame_ptr Optional, load color frame.
+ ///@return Whether loading succeeded.
+ DataLoadResult loadNext(DepthImage* depth_frame_ptr, // NOLINT
+ Transform* T_L_C_ptr, // NOLINT
+ Camera* camera_ptr, // NOLINT
+ ColorImage* color_frame_ptr = nullptr) override;
+
+ protected:
+ const std::string base_path_;
+ const std::string dataset_name_;
+
+ // The pose file.
+ // Note(alexmillane): Note that all the poses are in a single file so we keep
+ // the file open here.
+ std::ifstream trajectory_file_;
+
+ // The camera intrinsics for this dataset are just described on the website,
+ // not by a file, so we hardcode them into this class.
+ Camera camera_;
+
+ // The next frame to be loaded
+ int frame_number_ = 0;
+};
+
+} // namespace redwood
+} // namespace datasets
+} // namespace nvblox
diff --git a/nvblox/executables/include/nvblox/datasets/replica.h b/nvblox/executables/include/nvblox/datasets/replica.h
index 7cc3f4be..24015c4d 100644
--- a/nvblox/executables/include/nvblox/datasets/replica.h
+++ b/nvblox/executables/include/nvblox/datasets/replica.h
@@ -17,6 +17,7 @@ limitations under the License.
#include
#include
+#include
#include "nvblox/core/types.h"
#include "nvblox/datasets/data_loader.h"
@@ -34,6 +35,7 @@ std::unique_ptr createFuser(const std::string base_path);
class DataLoader : public RgbdDataLoaderInterface {
public:
DataLoader(const std::string& base_path, bool multithreaded = true);
+ virtual ~DataLoader() = default;
/// Interface for a function that loads the next frames in a dataset
///@param[out] depth_frame_ptr The loaded depth frame.
diff --git a/nvblox/executables/src/datasets/3dmatch.cpp b/nvblox/executables/src/datasets/3dmatch.cpp
index e48db361..70aa508a 100644
--- a/nvblox/executables/src/datasets/3dmatch.cpp
+++ b/nvblox/executables/src/datasets/3dmatch.cpp
@@ -116,9 +116,7 @@ std::unique_ptr> createColorImageLoader(
std::unique_ptr createFuser(const std::string base_path,
const int seq_id) {
- // Object to load 3DMatch data
auto data_loader = std::make_unique(base_path, seq_id);
- // Fuser
return std::make_unique(std::move(data_loader));
}
diff --git a/nvblox/executables/src/datasets/image_loader.cpp b/nvblox/executables/src/datasets/image_loader.cpp
index b77e5559..b0eccef3 100644
--- a/nvblox/executables/src/datasets/image_loader.cpp
+++ b/nvblox/executables/src/datasets/image_loader.cpp
@@ -64,16 +64,17 @@ bool load8BitColorImage(const std::string& filename,
timing::Timer stbi_timer("file_loading/color_image/stbi");
int width, height, num_channels;
uint8_t* image_data =
- stbi_load(filename.c_str(), &width, &height, &num_channels, 0);
+ stbi_load(filename.c_str(), &width, &height, &num_channels, 4);
stbi_timer.Stop();
if (image_data == nullptr) {
return false;
}
- CHECK_EQ(num_channels, 3);
+ // Currently we only support loading 3 channel (rgb) or 4 channel (rgba) images.
+ CHECK(num_channels == 3 || num_channels == 4);
- CHECK_EQ(sizeof(Color), 3 * sizeof(uint8_t))
- << "Color is padded so image loading wont work.";
+ CHECK_EQ(sizeof(Color), 4 * sizeof(uint8_t))
+ << "Color struct was padded by the compiler so image loading wont work.";
*color_image_ptr = ColorImage::fromBuffer(
height, width, reinterpret_cast(image_data), memory_type);
diff --git a/nvblox/executables/src/datasets/redwood.cpp b/nvblox/executables/src/datasets/redwood.cpp
new file mode 100644
index 00000000..69347010
--- /dev/null
+++ b/nvblox/executables/src/datasets/redwood.cpp
@@ -0,0 +1,205 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#include "nvblox/datasets/redwood.h"
+
+#include
+
+#include
+#include
+#include
+#include
+#include
+
+#include "nvblox/utils/timing.h"
+
+namespace nvblox {
+namespace datasets {
+namespace redwood {
+namespace internal {
+
+std::string getDatasetNameFromBasePath(const std::string& base_path) {
+ // Remove trailing slash
+ std::string base_path_copy = base_path;
+ if (base_path_copy.back() == '/') {
+ base_path_copy.pop_back();
+ }
+ // Last part of the path is the dataset name.
+ size_t pos = base_path_copy.find_last_of('/');
+ if (pos == std::string::npos) {
+ pos = 0;
+ }
+ return base_path_copy.substr(pos + 1);
+}
+
+void parsePoseFromStream(std::ifstream* trajectory_file_stream_ptr,
+ Transform* transform_ptr) {
+ CHECK_NOTNULL(transform_ptr);
+ constexpr int kDimension = 4;
+ for (int row = 0; row < kDimension; row++) {
+ for (int col = 0; col < kDimension; col++) {
+ float item = 0.0;
+ (*trajectory_file_stream_ptr) >> item;
+ (*transform_ptr)(row, col) = item;
+ }
+ }
+ // NOTE(alexmillane): There's some weird trailing characters... just reading
+ // one extra char solves things
+ char extra_char;
+ (*trajectory_file_stream_ptr) >> extra_char;
+}
+
+std::string getPathForTrajectory(const std::string& base_path,
+ const std::string& dataset_name) {
+ return base_path + "/pose_" + dataset_name + "/" + dataset_name + ".log";
+}
+
+std::string getPathForDepthImage(const std::string& base_path,
+ const std::string& dataset_name,
+ const int frame_id) {
+ std::stringstream ss;
+ ss << base_path << "/rgbd_" << dataset_name << "/" << dataset_name
+ << "/depth/" << std::setfill('0') << std::setw(6) << frame_id << ".png";
+ return ss.str();
+}
+
+std::string getPathForColorImage(const std::string& base_path,
+ const std::string& dataset_name,
+ const int frame_id) {
+ std::stringstream ss;
+ ss << base_path << "/rgbd_" << dataset_name << "/" << dataset_name
+ << "/image/" << std::setfill('0') << std::setw(6) << frame_id << ".jpg";
+ return ss.str();
+}
+
+std::unique_ptr> createDepthImageLoader(
+ const std::string& base_path, const std::string& dataset_name,
+ const float depth_scaling_factor, const bool multithreaded) {
+ return createImageLoader(
+ std::bind(getPathForDepthImage, base_path, dataset_name,
+ std::placeholders::_1),
+ multithreaded, depth_scaling_factor);
+}
+
+std::unique_ptr> createColorImageLoader(
+ const std::string& base_path, const std::string& dataset_name,
+ const bool multithreaded) {
+ return createImageLoader(
+ std::bind(getPathForColorImage, base_path, dataset_name,
+ std::placeholders::_1),
+ multithreaded);
+}
+
+} // namespace internal
+
+std::unique_ptr createFuser(const std::string base_path) {
+ auto data_loader = std::make_unique(base_path);
+ return std::make_unique(std::move(data_loader));
+}
+
+DataLoader::DataLoader(const std::string& base_path, bool multithreaded)
+ : RgbdDataLoaderInterface(
+ internal::createDepthImageLoader(
+ base_path, internal::getDatasetNameFromBasePath(base_path),
+ datasets::kDefaultUintDepthScaleFactor, multithreaded),
+ internal::createColorImageLoader(
+ base_path, internal::getDatasetNameFromBasePath(base_path),
+ multithreaded)),
+ base_path_(base_path),
+ dataset_name_(internal::getDatasetNameFromBasePath(base_path)),
+ trajectory_file_(std::ifstream(internal::getPathForTrajectory(
+ base_path, internal::getDatasetNameFromBasePath(base_path)))) {
+ constexpr float fu = 525;
+ constexpr float fv = 525;
+ constexpr int width = 640;
+ constexpr int height = 480;
+ camera_ = Camera(fu, fv, width, height);
+}
+
+/// Interface for a function that loads the next frames in a dataset
+///@param[out] depth_frame_ptr The loaded depth frame.
+///@param[out] T_L_C_ptr Transform from Camera to the Layer frame.
+///@param[out] camera_ptr The intrinsic camera model.
+///@param[out] color_frame_ptr Optional, load color frame.
+///@return Whether loading succeeded.
+DataLoadResult DataLoader::loadNext(DepthImage* depth_frame_ptr,
+ Transform* T_L_C_ptr, Camera* camera_ptr,
+ ColorImage* color_frame_ptr) {
+ CHECK_NOTNULL(depth_frame_ptr);
+ CHECK_NOTNULL(T_L_C_ptr);
+ CHECK_NOTNULL(camera_ptr);
+ // CHECK_NOTNULL(color_frame_ptr); // Can be null
+
+ // Because we might fail along the way, increment the frame number before we
+ // start.
+ const int frame_number = frame_number_;
+ ++frame_number_;
+
+ // Load the image into a Depth Frame.
+ CHECK(depth_image_loader_);
+ timing::Timer timer_file_depth("file_loading/depth_image");
+ if (!depth_image_loader_->getNextImage(depth_frame_ptr)) {
+ LOG(INFO) << "Couldn't find depth image";
+ return DataLoadResult::kNoMoreData;
+ }
+ timer_file_depth.Stop();
+
+ // Load the color image into a ColorImage
+ if (color_frame_ptr) {
+ CHECK(color_image_loader_);
+ timing::Timer timer_file_color("file_loading/color_image");
+ if (!color_image_loader_->getNextImage(color_frame_ptr)) {
+ LOG(INFO) << "Couldn't find color image";
+ return DataLoadResult::kNoMoreData;
+ }
+ timer_file_color.Stop();
+ }
+
+ float scale;
+
+ // Get the camera for this frame.
+ timing::Timer timer_file_intrinsics("file_loading/camera");
+ *camera_ptr = camera_;
+ timer_file_intrinsics.Stop();
+
+ // Get the next pose
+ timing::Timer timer_file_pose("file_loading/pose");
+ CHECK(trajectory_file_.is_open());
+ std::string line;
+ if (std::getline(trajectory_file_, line)) {
+ redwood::internal::parsePoseFromStream(&trajectory_file_, T_L_C_ptr);
+ } else {
+ LOG(INFO) << "Couldn't find pose";
+ return DataLoadResult::kNoMoreData;
+ }
+
+ // Check that the loaded data doesn't contain NaNs or a faulty rotation
+ // matrix. This does occur. If we find one, skip that frame and move to the
+ // next.
+ constexpr float kRotationMatrixDetEpsilon = 1e-4;
+ if (!T_L_C_ptr->matrix().allFinite() ||
+ std::abs(T_L_C_ptr->matrix().block<3, 3>(0, 0).determinant() - 1.0f) >
+ kRotationMatrixDetEpsilon) {
+ LOG(WARNING) << "Bad CSV data.";
+ return DataLoadResult::kBadFrame; // Bad data, but keep going.
+ }
+
+ timer_file_pose.Stop();
+ return DataLoadResult::kSuccess;
+}
+
+} // namespace redwood
+} // namespace datasets
+} // namespace nvblox
diff --git a/nvblox/executables/src/datasets/replica.cpp b/nvblox/executables/src/datasets/replica.cpp
index 07f2c3e0..371f0ecb 100644
--- a/nvblox/executables/src/datasets/replica.cpp
+++ b/nvblox/executables/src/datasets/replica.cpp
@@ -142,9 +142,7 @@ std::unique_ptr> createColorImageLoader(
} // namespace internal
std::unique_ptr createFuser(const std::string base_path) {
- // Object to load 3DMatch data
auto data_loader = std::make_unique(base_path);
- // Fuser
return std::make_unique(std::move(data_loader));
}
diff --git a/nvblox/executables/src/fuse_redwood.cpp b/nvblox/executables/src/fuse_redwood.cpp
new file mode 100644
index 00000000..cb5f8f07
--- /dev/null
+++ b/nvblox/executables/src/fuse_redwood.cpp
@@ -0,0 +1,61 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+
+
+#include
+#include
+#include
+
+#include
+#include
+
+#include
+#include
+
+#include "nvblox/core/image.h"
+#include "nvblox/core/types.h"
+#include "nvblox/datasets/image_loader.h"
+#include "nvblox/datasets/redwood.h"
+#include "nvblox/executables/fuser.h"
+
+using namespace nvblox;
+
+int main(int argc, char* argv[]) {
+ gflags::ParseCommandLineFlags(&argc, &argv, true);
+ google::InitGoogleLogging(argv[0]);
+ FLAGS_alsologtostderr = true;
+ google::InstallFailureSignalHandler();
+
+ // Path to the dataset
+ if (argc < 2) {
+ LOG(ERROR) << "No path to data directory given, failing.";
+ return 1;
+ }
+ std::string base_path = argv[1];
+ LOG(INFO) << "Loading Redwood files from " << base_path;
+
+ // Fuser
+ std::unique_ptr fuser = datasets::redwood::createFuser(base_path);
+
+ // Mesh location (optional)
+ if (argc >= 3) {
+ fuser->mesh_output_path_ = argv[2];
+ LOG(INFO) << "Mesh location:" << fuser->mesh_output_path_;
+ }
+
+ // Make sure the layers are the correct resolution.
+ return fuser->run();
+}
diff --git a/nvblox/include/nvblox/core/accessors.h b/nvblox/include/nvblox/core/accessors.h
index 0c3046a0..9590dae3 100644
--- a/nvblox/include/nvblox/core/accessors.h
+++ b/nvblox/include/nvblox/core/accessors.h
@@ -24,47 +24,51 @@ limitations under the License.
namespace nvblox {
-// Get Block
+/// Get Block
template
bool getBlockAtPosition(const BlockLayer& layer,
const Eigen::Vector3f& position,
const BlockType** data);
-// Get Voxel
+/// Get Voxel
template
bool getVoxelAtPosition(const BlockLayer>& layer,
const Eigen::Vector3f& position,
const VoxelType** data);
-// Accessors for voxelized types
+/// Accessors for voxelized types
template
const VoxelType* getVoxelAtBlockAndVoxelIndex(
const BlockLayer>& layer, const Index3D& block_index,
const Index3D& voxel_index);
-// Accessors for calling a function on all voxels in a layer
+/// Accessors for calling a function on all voxels in a layer (const).
template
using ConstVoxelCallbackFunction =
std::function;
+/// Accessors for calling a function on all voxels in a layer (non-const).
template
using VoxelCallbackFunction = std::function;
+/// Call function on all voxels in a layer (const).
template
void callFunctionOnAllVoxels(const BlockLayer>& layer,
ConstVoxelCallbackFunction callback);
+/// Call function on all voxels in a layer (non-const).
template
void callFunctionOnAllVoxels(BlockLayer>* layer,
VoxelCallbackFunction callback);
-// Accessors for calling a function on all voxels in a block
+/// Accessors for calling a function on all voxels in a block (const).
template
using ConstBlockCallbackFunction =
std::function;
+/// Accessors for calling a function on all voxels in a block (non-const).
template
using BlockCallbackFunction =
std::function;
@@ -77,7 +81,7 @@ template
void callFunctionOnAllVoxels(VoxelBlock* block,
BlockCallbackFunction callback);
-// Accessors for calling a function on all blocks in a layer
+/// Accessors for calling a function on all blocks in a layer
template
using ConstBlockCallbackFunction =
std::function;
diff --git a/nvblox/include/nvblox/core/blox.h b/nvblox/include/nvblox/core/blox.h
index d7da7e37..977bac8d 100644
--- a/nvblox/include/nvblox/core/blox.h
+++ b/nvblox/include/nvblox/core/blox.h
@@ -22,6 +22,7 @@ limitations under the License.
namespace nvblox {
+/// A block that contains 8x8x8 voxels of a given type.
template
struct VoxelBlock {
typedef unified_ptr Ptr;
@@ -30,13 +31,18 @@ struct VoxelBlock {
static constexpr size_t kVoxelsPerSide = 8;
VoxelType voxels[kVoxelsPerSide][kVoxelsPerSide][kVoxelsPerSide];
+ /// Allocate a voxel block of a given memory type.
static Ptr allocate(MemoryType memory_type);
+ /// Initializes all the memory of the voxels to 0 by default, can be
+ /// specialized by voxel type.
static void initOnGPU(VoxelBlock* block_ptr);
};
// Initialization Utility Functions
+/// Set all the memory of the block to 0 on the GPU.
template
void setBlockBytesZeroOnGPU(BlockType* block_device_ptr);
+/// Set all of the default colors to gray on a GPU.
void setColorBlockGrayOnGPU(VoxelBlock* block_device_ptr);
} // namespace nvblox
diff --git a/nvblox/include/nvblox/core/bounding_boxes.h b/nvblox/include/nvblox/core/bounding_boxes.h
index 6d153fc2..f68b4a14 100644
--- a/nvblox/include/nvblox/core/bounding_boxes.h
+++ b/nvblox/include/nvblox/core/bounding_boxes.h
@@ -20,27 +20,43 @@ limitations under the License.
#include "nvblox/core/types.h"
namespace nvblox {
-
+/// Get all of the blocks that are touched by an AABB, no matter how little.
+/// @param block_size Metric size of the block.
+/// @param aabb Axis-Aligned Bounding Box that does the touching.
+/// @return Vector of all the touched block indices.
inline std::vector getBlockIndicesTouchedByBoundingBox(
const float block_size, const AxisAlignedBoundingBox& aabb);
+/// Gets the Axis-Aligned Bounding Box of a block.
+/// @param block_size Metric size of the block.
+/// @param block_index The index of the block.
+/// @return The AABB.
inline AxisAlignedBoundingBox getAABBOfBlock(const float block_size,
const Index3D& block_index);
+/// Get AABB that covers ALL blocks in the block index list.
AxisAlignedBoundingBox getAABBOfBlocks(const float block_size,
const std::vector& blocks);
+/// Get the outer AABB of all of the allocated blocks in the layer.
template
AxisAlignedBoundingBox getAABBOfAllocatedBlocks(
const BlockLayer& layer);
+/// Get all of the allocated blocks that are within a given AABB.
template
std::vector getAllocatedBlocksWithinAABB(
const BlockLayer& layer, const AxisAlignedBoundingBox& aabb);
+/// Get the outer AABB of all of the blocks that contain observed voxels in an
+/// ESDF layer.
AxisAlignedBoundingBox getAABBOfObservedVoxels(const EsdfLayer& layer);
+/// Get the outer AABB of all of the blocks that contain observed voxels in an
+/// TSDF layer.
AxisAlignedBoundingBox getAABBOfObservedVoxels(const TsdfLayer& layer,
const float min_weight = 1e-4);
+/// Get the outer AABB of all of the blocks that contain observed voxels in an
+/// Color layer.
AxisAlignedBoundingBox getAABBOfObservedVoxels(const ColorLayer& layer,
const float min_weight = 1e-4);
diff --git a/nvblox/include/nvblox/core/camera.h b/nvblox/include/nvblox/core/camera.h
index 01dc7214..990bde9b 100644
--- a/nvblox/include/nvblox/core/camera.h
+++ b/nvblox/include/nvblox/core/camera.h
@@ -30,6 +30,10 @@ class Camera {
int width, int height);
__host__ __device__ inline Camera(float fu, float fv, int width, int height);
+ /// Project a 3D point in camera image space to a 2D pixel coordinate.
+ /// @param p_C Input 3D point coordinate in image space.
+ /// @param u_C Output 2D pixel coordinate.
+ /// @return Whether the pixel is on the image.
__host__ __device__ inline bool project(const Vector3f& p_C,
Vector2f* u_C) const;
diff --git a/nvblox/include/nvblox/core/color.h b/nvblox/include/nvblox/core/color.h
index b5d84d7d..f59cf731 100644
--- a/nvblox/include/nvblox/core/color.h
+++ b/nvblox/include/nvblox/core/color.h
@@ -21,20 +21,25 @@ limitations under the License.
namespace nvblox {
+/// Color, stored as 8-bit RGBA, with helper functions for commonly-used colors.
struct Color {
- __host__ __device__ Color() : r(0), g(0), b(0) {}
+ __host__ __device__ Color() : r(0), g(0), b(0), a(0) {}
__host__ __device__ Color(uint8_t _r, uint8_t _g, uint8_t _b)
- : r(_r), g(_g), b(_b) {}
+ : r(_r), g(_g), b(_b), a(255) {}
+ __host__ __device__ Color(uint8_t _r, uint8_t _g, uint8_t _b, uint8_t _a)
+ : r(_r), g(_g), b(_b), a(_a) {}
uint8_t r;
uint8_t g;
uint8_t b;
+ uint8_t a;
+ /// Check if colors are exactly identical. Also checks alpha.
bool operator==(const Color& other) const {
- return (r == other.r) && (g == other.g) && (b == other.b);
+ return (r == other.r) && (g == other.g) && (b == other.b) && (a == other.a);
}
- // Static functions for working with colors
+ /// Static functions for working with colors
static Color blendTwoColors(const Color& first_color, float first_weight,
const Color& second_color, float second_weight);
diff --git a/nvblox/include/nvblox/core/cuda/atomic_float.cuh b/nvblox/include/nvblox/core/cuda/atomic_float.cuh
new file mode 100644
index 00000000..9fa52ba6
--- /dev/null
+++ b/nvblox/include/nvblox/core/cuda/atomic_float.cuh
@@ -0,0 +1,24 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+namespace nvblox {
+
+__device__ __forceinline__ float atomicMinFloat(float* addr, float value);
+
+} // namespace nvblox
+
+#include "nvblox/core/cuda/impl/atomic_float_impl.cuh"
diff --git a/nvblox/include/nvblox/core/cuda/impl/atomic_float_impl.cuh b/nvblox/include/nvblox/core/cuda/impl/atomic_float_impl.cuh
new file mode 100644
index 00000000..40afc83c
--- /dev/null
+++ b/nvblox/include/nvblox/core/cuda/impl/atomic_float_impl.cuh
@@ -0,0 +1,32 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+namespace nvblox {
+
+// From:
+// https://stackoverflow.com/questions/17399119/how-do-i-use-atomicmax-on-floating-point-values-in-cuda
+__device__ __forceinline__ float atomicMinFloat(float* addr, float value) {
+ float old;
+ old = (value >= 0)
+ ? __int_as_float(atomicMin((int*)addr, __float_as_int(value)))
+ : __uint_as_float(
+ atomicMax((unsigned int*)addr, __float_as_uint(value)));
+
+ return old;
+}
+
+} // namespace nvblox
diff --git a/nvblox/include/nvblox/core/cuda/impl/layer_impl.cuh b/nvblox/include/nvblox/core/cuda/impl/layer_impl.cuh
new file mode 100644
index 00000000..83afbfdc
--- /dev/null
+++ b/nvblox/include/nvblox/core/cuda/impl/layer_impl.cuh
@@ -0,0 +1,70 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+#include "nvblox/core/layer.h"
+#include "nvblox/gpu_hash/cuda/gpu_indexing.cuh"
+
+namespace nvblox {
+
+template
+__global__ void queryVoxelsKernel(
+ int num_queries, Index3DDeviceHashMapType> block_hash,
+ float block_size, const Vector3f* query_locations_ptr,
+ VoxelType* voxels_ptr, bool* success_flags_ptr) {
+ const int idx = threadIdx.x + blockIdx.x * blockDim.x;
+ if (idx >= num_queries) {
+ return;
+ }
+ const Vector3f query_location = query_locations_ptr[idx];
+
+ VoxelType* voxel;
+ if (!getVoxelAtPosition(block_hash, query_location, block_size,
+ &voxel)) {
+ success_flags_ptr[idx] = false;
+ } else {
+ success_flags_ptr[idx] = true;
+ voxels_ptr[idx] = *voxel;
+ }
+}
+
+template
+void VoxelBlockLayer::getVoxelsGPU(
+ const device_vector& positions_L,
+ device_vector* voxels_ptr,
+ device_vector* success_flags_ptr) const {
+ CHECK_NOTNULL(voxels_ptr);
+ CHECK_NOTNULL(success_flags_ptr);
+
+ const int num_queries = positions_L.size();
+ voxels_ptr->resize(num_queries);
+ success_flags_ptr->resize(num_queries);
+
+ cudaStream_t cuda_stream;
+ checkCudaErrors(cudaStreamCreate(&cuda_stream));
+ constexpr int kNumThreads = 512;
+ const int num_blocks = num_queries / kNumThreads + 1;
+
+ queryVoxelsKernel<<>>(
+ num_queries, this->getGpuLayerView().getHash().impl_, this->block_size_,
+ positions_L.data(), voxels_ptr->data(), success_flags_ptr->data());
+
+ checkCudaErrors(cudaStreamSynchronize(cuda_stream));
+ checkCudaErrors(cudaPeekAtLastError());
+ checkCudaErrors(cudaStreamDestroy(cuda_stream));
+}
+
+} // namespace nvblox
diff --git a/nvblox/include/nvblox/core/image.h b/nvblox/include/nvblox/core/image.h
index 453c1f47..c6fc0487 100644
--- a/nvblox/include/nvblox/core/image.h
+++ b/nvblox/include/nvblox/core/image.h
@@ -50,16 +50,17 @@ __host__ __device__ inline ElementType& access(int linear_idx,
} // namespace image
-// Row-major image.
-// - Note that in a row-major image, rows follow one another in linear memory,
-// which means the column index varied between subsequent elements.
-// - Images use corner based indexing such that the pixel with index (0,0) is
-// centered at (0.5px,0.5px) and spans (0.0px,0.0px) to (1.0px,1.0px).
-// - Points on the image plane are defined as: u_px = (u_px.x(), u_px.y()) =
-// (col_idx, row_idx) in pixels.
+
constexpr MemoryType kDefaultImageMemoryType = MemoryType::kDevice;
+/// Row-major image.
+/// - Note that in a row-major image, rows follow one another in linear memory,
+/// which means the column index varied between subsequent elements.
+/// - Images use corner based indexing such that the pixel with index (0,0) is
+/// centered at (0.5px,0.5px) and spans (0.0px,0.0px) to (1.0px,1.0px).
+/// - Points on the image plane are defined as: u_px = (u_px.x(), u_px.y()) =
+/// (col_idx, row_idx) in pixels.
template
class Image {
public:
@@ -77,12 +78,12 @@ class Image {
Image(Image&& other);
Image& operator=(Image&& other);
- // Deep copy constructor (second can be used to transition memory type)
+ /// Deep copy constructor (second can be used to transition memory type)
explicit Image(const Image& other);
Image(const Image& other, MemoryType memory_type);
Image& operator=(const Image& other);
- // Prefetch the data to the gpu (only makes sense for kUnified images)
+ /// Prefetch the data to the gpu (only makes sense for kUnified images)
void toGPU() const;
// Attributes
@@ -93,9 +94,9 @@ class Image {
inline int height() const { return rows_; }
inline MemoryType memory_type() const { return memory_type_; }
- // Access
- // NOTE(alexmillane): The guard-rails are off here. If you declare a kDevice
- // Image and try to access its data, you will get undefined behaviour.
+ /// Access
+ /// NOTE(alexmillane): The guard-rails are off here. If you declare a kDevice
+ /// Image and try to access its data, you will get undefined behaviour.
inline ElementType operator()(const int row_idx, const int col_idx) const {
return image::access(row_idx, col_idx, cols_, data_.get());
}
@@ -109,18 +110,18 @@ class Image {
return image::access(linear_idx, data_.get());
}
- // Raw pointer access
+ /// Raw pointer access
inline ElementType* dataPtr() { return data_.get(); }
inline const ElementType* dataConstPtr() const { return data_.get(); }
- // Reset the contents of the image. Reallocate if the image got larger.
+ /// Reset the contents of the image. Reallocate if the image got larger.
void populateFromBuffer(int rows, int cols, const ElementType* buffer,
MemoryType memory_type = kDefaultImageMemoryType);
- // Set the image to 0.
+ /// Set the image to 0.
void setZero();
- // Factories
+ /// Factories
static inline Image fromBuffer(
int rows, int cols, const ElementType* buffer,
MemoryType memory_type = kDefaultImageMemoryType);
@@ -134,8 +135,9 @@ class Image {
using DepthImage = Image;
using ColorImage = Image;
+using MonoImage = Image;
-// Image Reductions
+// Image Operations
namespace image {
float max(const DepthImage& image);
diff --git a/nvblox/include/nvblox/core/impl/image_impl.h b/nvblox/include/nvblox/core/impl/image_impl.h
index 6bceee2a..afbb0959 100644
--- a/nvblox/include/nvblox/core/impl/image_impl.h
+++ b/nvblox/include/nvblox/core/impl/image_impl.h
@@ -18,8 +18,6 @@ limitations under the License.
#include
#include "nvblox/core/cuda/image_cuda.h"
-#include "nvblox/core/interpolation_2d.h"
-#include "nvblox/io/csv.h"
namespace nvblox {
diff --git a/nvblox/include/nvblox/core/impl/interpolation_2d_impl.h b/nvblox/include/nvblox/core/impl/interpolation_2d_impl.h
index 705438eb..93d6d632 100644
--- a/nvblox/include/nvblox/core/impl/interpolation_2d_impl.h
+++ b/nvblox/include/nvblox/core/impl/interpolation_2d_impl.h
@@ -75,6 +75,12 @@ struct FloatPixelGreaterThanZero {
}
};
+struct ColorPixelAlphaGreaterThanZero {
+ __host__ __device__ static inline bool check(const Color& pixel_value) {
+ return pixel_value.a > 0;
+ }
+};
+
} // namespace checkers
// CPU interfaces
diff --git a/nvblox/include/nvblox/core/impl/unified_vector_impl.h b/nvblox/include/nvblox/core/impl/unified_vector_impl.h
index d0e1f20d..f08d68e9 100644
--- a/nvblox/include/nvblox/core/impl/unified_vector_impl.h
+++ b/nvblox/include/nvblox/core/impl/unified_vector_impl.h
@@ -166,7 +166,7 @@ inline std::vector unified_vector::toVector() const {
sizeof(bool) * buffer_size_, cudaMemcpyDefault));
// Now populate the vector
std::vector vect(buffer_size_);
- for (int i = 0; i < buffer_size_; i++) {
+ for (size_t i = 0; i < buffer_size_; i++) {
vect[i] = bool_buffer[i];
}
return vect;
diff --git a/nvblox/include/nvblox/core/indexing.h b/nvblox/include/nvblox/core/indexing.h
index 2b6da5ee..aff48736 100644
--- a/nvblox/include/nvblox/core/indexing.h
+++ b/nvblox/include/nvblox/core/indexing.h
@@ -22,8 +22,8 @@ namespace nvblox {
__host__ __device__ inline float voxelSizeToBlockSize(const float voxel_size);
__host__ __device__ inline float blockSizeToVoxelSize(const float block_size);
-// Assuming a fixed-size voxel block, get the voxel index of a voxel at that
-// position within a block.
+/// Assuming a fixed-size voxel block, get the voxel index of a voxel at that
+/// position within a block.
__host__ __device__ inline Index3D getVoxelIndexFromPositionInLayer(
const float block_size, const Vector3f& position);
diff --git a/nvblox/include/nvblox/core/interpolation_2d.h b/nvblox/include/nvblox/core/interpolation_2d.h
index b7c9d6a2..ebee63c5 100644
--- a/nvblox/include/nvblox/core/interpolation_2d.h
+++ b/nvblox/include/nvblox/core/interpolation_2d.h
@@ -23,13 +23,17 @@ namespace interpolation {
namespace checkers {
-// A checker that always returns that a pixel is valid (the default below)
+/// A checker that always returns that a pixel is valid (the default below)
template
struct PixelAlwaysValid;
-// A checker that returns true if a float pixel is greater than 0.0f.
+/// A checker that returns true if a float pixel is greater than 0.0f.
struct FloatPixelGreaterThanZero;
+/// A checker that returns true if the alpha channel of a color pixel is greater
+/// than 0.
+struct ColorPixelAlphaGreaterThanZero;
+
} // namespace checkers
template
diff --git a/nvblox/include/nvblox/core/interpolation_3d.h b/nvblox/include/nvblox/core/interpolation_3d.h
index 05e491e7..03651810 100644
--- a/nvblox/include/nvblox/core/interpolation_3d.h
+++ b/nvblox/include/nvblox/core/interpolation_3d.h
@@ -24,13 +24,13 @@ limitations under the License.
namespace nvblox {
namespace interpolation {
-// Single points
+/// Single points
bool interpolateOnCPU(const Vector3f& p_L, const TsdfLayer& layer,
float* distance);
bool interpolateOnCPU(const Vector3f& p_L, const EsdfLayer& layer,
float* distance);
-// Vectors of points
+/// Vectors of points
template
void interpolateOnCPU(const std::vector& points_L,
const VoxelBlockLayer& layer,
diff --git a/nvblox/include/nvblox/core/iterator.h b/nvblox/include/nvblox/core/iterator.h
index 83af41c2..f910ab52 100644
--- a/nvblox/include/nvblox/core/iterator.h
+++ b/nvblox/include/nvblox/core/iterator.h
@@ -21,6 +21,8 @@ limitations under the License.
namespace nvblox {
+/// Iterator class for unified_vectors that enables us to use thrust and STL
+/// operators on our own vectors.
template
struct RawIterator {
using iterator_category = std::forward_iterator_tag;
diff --git a/nvblox/include/nvblox/core/layer.h b/nvblox/include/nvblox/core/layer.h
index df42e402..d713734f 100644
--- a/nvblox/include/nvblox/core/layer.h
+++ b/nvblox/include/nvblox/core/layer.h
@@ -23,10 +23,12 @@ limitations under the License.
#include "nvblox/core/traits.h"
#include "nvblox/core/types.h"
#include "nvblox/core/unified_ptr.h"
+#include "nvblox/core/unified_vector.h"
#include "nvblox/gpu_hash/gpu_layer_view.h"
namespace nvblox {
+/// Base class for all layer objects.
class BaseLayer {
public:
typedef std::shared_ptr Ptr;
@@ -37,17 +39,18 @@ class BaseLayer {
// Just an interface class
};
+/// A layer that contains blocks, which are stored in a hash map.
template
class BlockLayer : public BaseLayer {
public:
typedef std::shared_ptr Ptr;
typedef std::shared_ptr ConstPtr;
- // Check that custom block types implement allocate
+ /// Check that custom block types implement allocate
static_assert(traits::has_allocate<_BlockType>::value,
"BlockType must implement an allocate() function.");
- // Allows inspection of the contained BlockType through LayerType::BlockType
+ /// Allows inspection of the contained BlockType through LayerType::BlockType
typedef _BlockType BlockType;
typedef BlockLayer LayerType;
typedef GPULayerView GPULayerViewType;
@@ -61,55 +64,55 @@ class BlockLayer : public BaseLayer {
gpu_layer_view_up_to_date_(false) {}
virtual ~BlockLayer() {}
- // Deep copies, with optionally changing the memory type.
+ /// Deep copies, with optionally changing the memory type.
BlockLayer(const BlockLayer& other);
BlockLayer(const BlockLayer& other, MemoryType memory_type);
BlockLayer& operator=(const BlockLayer& other);
- // Move operations
+ /// Move operations
BlockLayer(BlockLayer&& other) = default;
BlockLayer& operator=(BlockLayer&& other) = default;
- // Block accessors by index.
+ /// Block accessors by index.
typename BlockType::Ptr getBlockAtIndex(const Index3D& index);
typename BlockType::ConstPtr getBlockAtIndex(const Index3D& index) const;
typename BlockType::Ptr allocateBlockAtIndex(const Index3D& index);
- // Block accessors by position.
+ /// Block accessors by position.
typename BlockType::Ptr getBlockAtPosition(const Vector3f& position);
typename BlockType::ConstPtr getBlockAtPosition(
const Vector3f& position) const;
typename BlockType::Ptr allocateBlockAtPosition(const Vector3f& position);
- // Get all blocks indices.
+ /// Get all blocks indices.
std::vector getAllBlockIndices() const;
- // Check if allocated
+ /// Check if allocated
bool isBlockAllocated(const Index3D& index) const;
__host__ __device__ float block_size() const { return block_size_; }
int numAllocatedBlocks() const { return blocks_.size(); }
- // Clear the layer of all data
+ /// Clear the layer of all data
void clear() { blocks_.clear(); }
- // Clear (deallocate) blocks passed in
- // Note if a block does not exist, this function just (silently)
- // continues trying the rest of the list.
+ /// Clear (deallocate) blocks passed in
+ /// Note if a block does not exist, this function just (silently)
+ /// continues trying the rest of the list.
void clearBlocks(const std::vector& indices);
MemoryType memory_type() const { return memory_type_; }
- // GPU Hash
- // Note(alexmillane): The hash returned here is invalidated by calls to
- // allocateBlock
+ /// GPU Hash
+ /// Note(alexmillane): The hash returned here is invalidated by calls to
+ /// allocateBlock
GPULayerViewType getGpuLayerView() const;
protected:
float block_size_;
MemoryType memory_type_;
- // CPU Hash (Index3D -> BlockType::Ptr)
+ /// CPU Hash (Index3D -> BlockType::Ptr)
BlockHash blocks_;
/// GPU Hash
@@ -124,6 +127,8 @@ class BlockLayer : public BaseLayer {
mutable std::unique_ptr gpu_layer_view_;
};
+/// Specialization for BlockLayer that exclusively contains VoxelBlocks to make
+/// access easier.
template
class VoxelBlockLayer : public BlockLayer> {
public:
@@ -145,13 +150,13 @@ class VoxelBlockLayer : public BlockLayer> {
VoxelBlockLayer() = delete;
virtual ~VoxelBlockLayer() {}
- // Deep copies
+ /// Deep copies
VoxelBlockLayer(const VoxelBlockLayer& other);
VoxelBlockLayer(const VoxelBlockLayer& other, MemoryType memory_type);
// Assignment retains the current layer's memory type.
VoxelBlockLayer& operator=(const VoxelBlockLayer& other);
- // Move operations
+ /// Move operations
VoxelBlockLayer(VoxelBlockLayer&& other) = default;
VoxelBlockLayer& operator=(VoxelBlockLayer&& other) = default;
@@ -172,6 +177,10 @@ class VoxelBlockLayer : public BlockLayer> {
std::vector* voxels_ptr,
std::vector* success_flags_ptr) const;
+ void getVoxelsGPU(const device_vector& positions_L,
+ device_vector* voxels_ptr,
+ device_vector* success_flags_ptr) const;
+
/// Get a voxel by copy by (closest) position
/// The position is given with respect to the layer frame (L). The function
/// returns the closest voxels to the passed points.
diff --git a/nvblox/include/nvblox/core/layer_cake.h b/nvblox/include/nvblox/core/layer_cake.h
index ba5c51d0..b8ae20a9 100644
--- a/nvblox/include/nvblox/core/layer_cake.h
+++ b/nvblox/include/nvblox/core/layer_cake.h
@@ -24,33 +24,36 @@ limitations under the License.
namespace nvblox {
+/// Holds a collection of layers. Currently a restriction that only 1 layer
+/// of each type can be stored, and all layers should have the same voxel
+/// (or block) size.
class LayerCake {
public:
LayerCake() = default;
LayerCake(float voxel_size) : voxel_size_(voxel_size) {}
- // Deep Copy (disallowed for now)
+ /// Deep Copy (disallowed for now)
LayerCake(const LayerCake& other) = delete;
LayerCake& operator=(const LayerCake& other) = delete;
- // Move
+ /// Move
LayerCake(LayerCake&& other) = default;
LayerCake& operator=(LayerCake&& other) = default;
template
LayerType* add(MemoryType memory_type);
- // Moves the ownership of the layer to the LayerCake.
+ /// Moves the ownership of the layer to the LayerCake.
inline void insert(const std::type_index& type_index,
std::unique_ptr&& layer);
- // Retrieve layers (as pointers)
+ /// Retrieve layers (as pointers)
template
LayerType* getPtr();
template
const LayerType* getConstPtr() const;
- // Retrieve layers (as reference) (will fail if the layer doesn't exist)
+ /// Retrieve layers (as reference) (will fail if the layer doesn't exist)
template
const LayerType& get() const;
@@ -61,7 +64,7 @@ class LayerCake {
void clear() { layers_.clear(); }
- // Factory accepting a list of LayerTypes (and MemoryTypes)
+ /// Factory accepting a list of LayerTypes (and MemoryTypes)
template
static LayerCake create(float voxel_size, MemoryType memory_type);
template
@@ -73,14 +76,15 @@ class LayerCake {
}
float voxel_size() const { return voxel_size_; }
+ float block_size() const { return voxel_size_ * VoxelBlock::kVoxelsPerSide; }
private:
// Params
float voxel_size_ = 0.0f;
- // Stored layers
- // Note(alexmillane): Currently we restrict the cake to storing a single layer
- // of each type.
+ /// Stored layers
+ /// Note(alexmillane): Currently we restrict the cake to storing a single layer
+ /// of each type.
std::unordered_map> layers_;
};
diff --git a/nvblox/include/nvblox/core/lidar.h b/nvblox/include/nvblox/core/lidar.h
index 7c07e0c3..7654b77a 100644
--- a/nvblox/include/nvblox/core/lidar.h
+++ b/nvblox/include/nvblox/core/lidar.h
@@ -19,6 +19,8 @@ limitations under the License.
namespace nvblox {
+/// Helper class for handling input LIDAR pointclouds and storing the LIDAR
+/// intrinsics. This helps convert a LIDAR pointcloud into a depth image.
class Lidar {
public:
Lidar() = delete;
@@ -27,37 +29,37 @@ class Lidar {
float vertical_fov_rad);
__host__ __device__ inline ~Lidar() = default;
- // Projects a 3D point to the (floating-point) image plane
+ /// Projects a 3D point to the (floating-point) image plane
__host__ __device__ inline bool project(const Vector3f& p_C,
Vector2f* u_C) const;
- // Projects a 3D point to the (index-based) image plane
+ /// Projects a 3D point to the (index-based) image plane
__host__ __device__ inline bool project(const Vector3f& p_C,
Index2D* u_C) const;
- // Gets the depth of a point
+ /// Gets the depth of a point
__host__ __device__ inline float getDepth(const Vector3f& p_C) const;
- // Back projection (image plane point to 3D point)
+ /// Back projection (image plane point to 3D point)
__host__ __device__ inline Vector3f unprojectFromImagePlaneCoordinates(
const Vector2f& u_C, const float depth) const;
__host__ __device__ inline Vector3f unprojectFromPixelIndices(
const Index2D& u_C, const float depth) const;
- // Back projection (image plane point to ray)
- // NOTE(alexmillane): These return normalized vectors
+ /// Back projection (image plane point to ray)
+ /// NOTE(alexmillane): These return normalized vectors
__host__ __device__ inline Vector3f vectorFromImagePlaneCoordinates(
const Vector2f& u_C) const;
__host__ __device__ inline Vector3f vectorFromPixelIndices(
const Index2D& u_C) const;
- // Conversions between pixel indices and image plane coordinates
+ /// Conversions between pixel indices and image plane coordinates
__host__ __device__ inline Vector2f pixelIndexToImagePlaneCoordsOfCenter(
const Index2D& u_C) const;
__host__ __device__ inline Index2D imagePlaneCoordsToPixelIndex(
const Vector2f& u_C) const;
- // View
+ /// View
__host__ inline AxisAlignedBoundingBox getViewAABB(
const Transform& T_L_C, const float min_depth,
const float max_depth) const;
@@ -69,10 +71,10 @@ class Lidar {
__host__ __device__ inline int rows() const;
__host__ __device__ inline int cols() const;
- // Equality
+ /// Equality
__host__ inline friend bool operator==(const Lidar& lhs, const Lidar& rhs);
- // Hash
+ /// Hash
struct Hash {
__host__ inline size_t operator()(const Lidar& lidar) const;
};
diff --git a/nvblox/include/nvblox/core/mapper.h b/nvblox/include/nvblox/core/mapper.h
index ab4e4cc2..9c431f38 100644
--- a/nvblox/include/nvblox/core/mapper.h
+++ b/nvblox/include/nvblox/core/mapper.h
@@ -29,6 +29,7 @@ limitations under the License.
#include "nvblox/integrators/projective_color_integrator.h"
#include "nvblox/integrators/projective_tsdf_integrator.h"
#include "nvblox/mesh/mesh_integrator.h"
+#include "nvblox/semantics/image_masker.h"
namespace nvblox {
@@ -38,6 +39,7 @@ namespace nvblox {
class MapperBase {
public:
MapperBase() = default;
+ virtual ~MapperBase() = default;
/// Move
MapperBase(MapperBase&& other) = default;
@@ -68,7 +70,7 @@ class RgbdMapper : public MapperBase {
/// @param voxel_size_m The voxel size in meters for the contained layers.
/// @param memory_type In which type of memory the layers should be stored.
RgbdMapper(float voxel_size_m, MemoryType memory_type = MemoryType::kDevice);
- virtual ~RgbdMapper() {}
+ virtual ~RgbdMapper() = default;
/// Constructor which initializes from a saved map.
/// @param map_filepath Path to the serialized map to be loaded.
@@ -266,4 +268,65 @@ class RgbdMapper : public MapperBase {
Index3DSet esdf_blocks_to_update_;
};
+/// Sub-class of the RgbdMapper which extends it to deal with humans in the
+/// scene.
+class HumanMapper : public RgbdMapper {
+ public:
+ HumanMapper(float voxel_size_m, MemoryType memory_type = MemoryType::kDevice);
+ ~HumanMapper() = default;
+
+ /// Integrates a depth frame into the tsdf reconstruction
+ /// assuming identity transformation between depth and mask frame and
+ /// identical intrinsics.
+ ///@param depth_frame Depth frame to integrate. Depth in the image is
+ /// specified as a float representing meters.
+ ///@param mask Human mask. Interpreted as 0=non-human, >0=human.
+ ///@param T_L_C Pose of the camera, specified as a transform from Camera-frame
+ /// to Layer-frame transform.
+ ///@param camera Intrinsics model of the camera.
+ void integrateDepth(const DepthImage& depth_frame, const MonoImage& mask,
+ const Transform& T_L_C, const Camera& camera);
+
+ /// Integrates a depth frame into the tsdf reconstruction
+ /// using the transformation between depth and mask frame and their
+ /// intrinsics.
+ ///@param depth_frame Depth frame to integrate. Depth in the image is
+ /// specified as a float representing meters.
+ ///@param mask Human mask. Interpreted as 0=non-human, >0=human.
+ ///@param T_L_CD Pose of the depth camera, specified as a transform from
+ /// Camera-frame to Layer-frame transform.
+ ///@param T_CM_CD Transform from depth camera to mask camera frame.
+ ///@param depth_camera Intrinsics model of the depth camera.
+ ///@param mask_camera Intrinsics model of the mask camera.
+ void integrateDepth(const DepthImage& depth_frame, const MonoImage& mask,
+ const Transform& T_L_CD, const Transform& T_CM_CD,
+ const Camera& depth_camera, const Camera& mask_camera);
+
+ /// Integrates a color frame into the reconstruction.
+ ///@param color_frame Color image to integrate.
+ ///@param mask Human mask. Interpreted as 0=non-human, >0=human.
+ ///@param T_L_C Pose of the camera, specified as a transform from Camera-frame
+ /// to Layer-frame transform.
+ ///@param camera Intrinsics model of the camera.
+ void integrateColor(const ColorImage& color_frame, const MonoImage& mask,
+ const Transform& T_L_C, const Camera& camera);
+
+ // These functions return a reference to the masked images generated during
+ // the preceeding calls to integrateColor() and integrateDepth().
+ const DepthImage& getLastDepthFrameWithoutHumans();
+ const DepthImage& getLastDepthFrameOnlyHumans();
+ const ColorImage& getLastColorFrameWithoutHumans();
+ const ColorImage& getLastColorFrameOnlyHumans();
+
+ protected:
+ // Split depth images based on a mask.
+ // Note that we internally pre-allocate space for the split images on the
+ // first call.
+ ImageMasker image_masker_;
+ DepthImage depth_frame_no_humans_;
+ DepthImage depth_frame_only_humans_;
+ ColorImage color_frame_no_humans_;
+ ColorImage color_frame_only_humans_;
+};
+
} // namespace nvblox
\ No newline at end of file
diff --git a/nvblox/include/nvblox/core/pointcloud.h b/nvblox/include/nvblox/core/pointcloud.h
new file mode 100644
index 00000000..63b054be
--- /dev/null
+++ b/nvblox/include/nvblox/core/pointcloud.h
@@ -0,0 +1,91 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+#include "nvblox/core/camera.h"
+#include "nvblox/core/image.h"
+#include "nvblox/core/types.h"
+#include "nvblox/core/unified_vector.h"
+
+namespace nvblox {
+
+constexpr MemoryType kDefaultPointcloudMemoryType = MemoryType::kDevice;
+
+/// Pointcloud that lives in either device, host or unified memory.
+/// We represent a pointcloud as a vector of 3D vectors.
+class Pointcloud {
+ public:
+ /// Construct an empty pointcloud
+ Pointcloud(int num_points,
+ MemoryType memory_type = kDefaultPointcloudMemoryType);
+ Pointcloud(MemoryType memory_type = kDefaultPointcloudMemoryType);
+
+ /// Constructor from points
+ Pointcloud(const std::vector& points,
+ MemoryType memory_type = kDefaultPointcloudMemoryType);
+
+ /// Move operations
+ Pointcloud(Pointcloud&& other) = default;
+ Pointcloud& operator=(Pointcloud&& other) = default;
+
+ /// Deep copy constructor (second can be used to transition memory type)
+ Pointcloud(const Pointcloud& other);
+ Pointcloud(const Pointcloud& other, MemoryType memory_type);
+ Pointcloud& operator=(const Pointcloud& other);
+
+ /// Expand memory available
+ void resize(int num_points) { points_.resize(num_points); }
+
+ /// Attributes
+ inline int num_points() const { return points_.size(); }
+ inline int size() const { return points_.size(); }
+ inline MemoryType memory_type() const { return points_.memory_type(); }
+ inline bool empty() const { return points_.empty(); }
+
+ /// Access
+ /// NOTE(alexmillane): The guard-rails are off here. If you declare a kDevice
+ /// Image and try to access its data, you will get undefined behaviour. If you
+ /// access out of bounds, you're gonna have a bad time.
+ inline const Vector3f& operator()(const int idx) const {
+ return points_[idx];
+ }
+ inline Vector3f& operator()(const int idx) { return points_[idx]; }
+ const unified_vector& points() const { return points_; }
+ unified_vector& points() { return points_; }
+
+ /// Add points
+ void push_back(Vector3f&& point) { points_.push_back(point); }
+
+ /// Raw pointer access
+ inline Vector3f* dataPtr() { return points_.data(); }
+ inline const Vector3f* dataConstPtr() const { return points_.data(); }
+
+ /// Set the image to 0.
+ void setZero() { points_.setZero(); }
+
+ protected:
+ unified_vector points_;
+};
+
+/// Transforms the points in a pointcloud into another frame
+/// @param T_out_in Transform that takes a point in frame "in" to frame "out".
+/// @param pointcloud_in Pointcloud in frame "in".
+/// @param[out] pointcloud_out Pointer to pointcloud in frame "out".
+void transformPointcloudOnGPU(const Transform& T_out_in, // NOLINT
+ const Pointcloud& pointcloud_in, // NOLINT
+ Pointcloud* pointcloud_out_ptr);
+
+} // namespace nvblox
diff --git a/nvblox/include/nvblox/core/types.h b/nvblox/include/nvblox/core/types.h
index 2ac3d103..ac25e15e 100644
--- a/nvblox/include/nvblox/core/types.h
+++ b/nvblox/include/nvblox/core/types.h
@@ -25,12 +25,12 @@ limitations under the License.
namespace nvblox {
-// Whether the storage or processing is happening on CPU, GPU, or any future
-// amazing hardware- accelerated platform.
+/// Whether the storage or processing is happening on CPU, GPU, or any future
+/// amazing hardware- accelerated platform.
enum class DeviceType { kCPU, kGPU };
-// How GPU data is stored, either in Device-only or unified (both) memory.
-// NOTE(alexmillane): tag: c++17, switch to constexpr when we move to c++17.
+/// How GPU data is stored, either in Device-only or unified (both) memory.
+/// NOTE(alexmillane): tag: c++17, switch to constexpr when we move to c++17.
enum class MemoryType { kDevice, kUnified, kHost };
inline std::string toString(MemoryType memory_type) {
switch (memory_type) {
@@ -50,10 +50,10 @@ typedef Eigen::AlignedBox3f AxisAlignedBoundingBox;
typedef Eigen::Isometry3f Transform;
-// This can be replaced with std::byte once we go to C++17.
+/// This can be replaced with std::byte once we go to C++17.
typedef uint8_t Byte;
-// Aligned Eigen containers
+/// Aligned Eigen containers
template
using AlignedVector = std::vector>;
diff --git a/nvblox/include/nvblox/core/unified_ptr.h b/nvblox/include/nvblox/core/unified_ptr.h
index 97fd8ca9..d9673083 100644
--- a/nvblox/include/nvblox/core/unified_ptr.h
+++ b/nvblox/include/nvblox/core/unified_ptr.h
@@ -24,19 +24,19 @@ limitations under the License.
namespace nvblox {
-// shared_ptr for device, unified memory, and pinned host memory
-// Things to be aware of
-// - Single objects
-// - Constructor and Destructor are not called when memory_type==kDevice
-// (these are CPU functions). Therefore unified_ptr generates an error if
-// used in device mode with non-trivially destructible types.
-// - Both Constructor and Destructor are called when memory_type==kUnified ||
-// kHost.
-// - Arrays
-// - Default Constructor called when storing arrays in kUnified || kHost.
-// - No Constructor called in kDevice mode.
-// - No destructor called for ANY memory setting so we make a static check
-// that objects are trivially destructable.
+/// shared_ptr for device, unified memory, and pinned host memory
+/// Things to be aware of
+/// - Single objects
+/// - Constructor and Destructor are not called when memory_type==kDevice
+/// (these are CPU functions). Therefore unified_ptr generates an error if
+/// used in device mode with non-trivially destructible types.
+/// - Both Constructor and Destructor are called when memory_type==kUnified ||
+/// kHost.
+/// - Arrays
+/// - Default Constructor called when storing arrays in kUnified || kHost.
+/// - No Constructor called in kDevice mode.
+/// - No destructor called for ANY memory setting so we make a static check
+/// that objects are trivially destructable.
template
class unified_ptr {
public:
@@ -52,24 +52,24 @@ class unified_ptr {
unified_ptr(unified_ptr&& other);
~unified_ptr();
- // Operator =
+ /// Operator =
unified_ptr& operator=(const unified_ptr& other);
unified_ptr& operator=(unified_ptr&& other);
- // Operator bool
+ /// Operator bool
operator bool() const;
- // Operator comparison
+ /// Operator comparison
bool operator==(const unified_ptr& other) const;
bool operator!=(const unified_ptr& other) const;
- // Operator dereference
+ /// Operator dereference
T_noextent* operator->();
const T_noextent* operator->() const;
T& operator*();
const T& operator*() const;
- // Operator array access
- // Only enabled if the underlying type is an array
+ /// Operator array access
+ /// Only enabled if the underlying type is an array
template ::value, bool> = true>
T_noextent& operator[](size_t i) {
@@ -83,31 +83,31 @@ class unified_ptr {
return ptr_[i];
}
- // Operator convert
- // Only enabled if the base type is NOT const, otherwise adds a second
- // trivial converter.
+ /// Operator convert
+ /// Only enabled if the base type is NOT const, otherwise adds a second
+ /// trivial converter.
template ::value, bool> = true>
operator unified_ptr() const;
- // Get the raw pointer.
+ /// Get the raw pointer.
T_noextent* get();
const T_noextent* get() const;
- // Reset the pointer to point to nothing.
+ /// Reset the pointer to point to nothing.
void reset();
- // Copy the underlying object (potentially to another memory location)
- // NOTE: This is implemented as a memcpy at the pointed to location.
+ /// Copy the underlying object (potentially to another memory location)
+ /// NOTE: This is implemented as a memcpy at the pointed to location.
unified_ptr clone() const;
unified_ptr clone(MemoryType memory_type) const;
- // Copy memory between two unified ptrs, potentially of different memory locations.
+ /// Copy memory between two unified ptrs, potentially of different memory locations.
void copyTo(unified_ptr& ptr) const;
MemoryType memory_type() const { return memory_type_; }
- // Helper function to memset all the memory to 0.
+ /// Helper function to memset all the memory to 0.
void setZero();
friend class unified_ptr;
diff --git a/nvblox/include/nvblox/core/unified_vector.h b/nvblox/include/nvblox/core/unified_vector.h
index 4cab8e24..b0a23b1e 100644
--- a/nvblox/include/nvblox/core/unified_vector.h
+++ b/nvblox/include/nvblox/core/unified_vector.h
@@ -34,7 +34,7 @@ class unified_vector {
static constexpr MemoryType kDefaultMemoryType = MemoryType::kUnified;
- // Static asserts on the type.
+ /// Static asserts on the type.
static_assert(
std::is_default_constructible::value,
"Objects stored in unified vector should be default constructible.");
@@ -43,50 +43,50 @@ class unified_vector {
unified_vector(size_t size, MemoryType memory_type = kDefaultMemoryType);
unified_vector(size_t size, const T& initial,
MemoryType memory_type = kDefaultMemoryType);
- // Copy constructor.
+ /// Copy constructor.
unified_vector(const unified_vector& other,
MemoryType memory_type = kDefaultMemoryType);
unified_vector(const std::vector& other,
MemoryType memory_type = kDefaultMemoryType);
- // Move constructor.
+ /// Move constructor.
unified_vector(unified_vector&& other);
- // Destructor
+ /// Destructor
virtual ~unified_vector();
- // Operators.
+ /// Operators.
T& operator[](size_t index);
const T& operator[](size_t index) const;
unified_vector& operator=(const unified_vector& other);
unified_vector& operator=(unified_vector&& other);
unified_vector& operator=(const std::vector& other);
- // Convert to an std::vector. Creates a copy.
+ /// Convert to an std::vector. Creates a copy.
std::vector toVector() const;
- // Get raw pointers. This is also for GPU pointers.
+ /// Get raw pointers. This is also for GPU pointers.
T* data();
const T* data() const;
- // Hint to move the memory to the GPU or CPU.
+ /// Hint to move the memory to the GPU or CPU.
void toGPU();
void toCPU();
- // Access information.
+ /// Access information.
size_t capacity() const;
size_t size() const;
bool empty() const;
- // Changing the size.
+ /// Changing the size.
void reserve(size_t capacity);
void resize(size_t size);
void clear();
- // Adding elements.
+ /// Adding elements.
void push_back(const T& value);
- // Iterator access.
+ /// Iterator access.
iterator begin();
iterator end();
const_iterator begin() const { return cbegin(); }
@@ -94,10 +94,10 @@ class unified_vector {
const_iterator cbegin() const;
const_iterator cend() const;
- // Get the memory type.
+ /// Get the memory type.
MemoryType memory_type() const { return memory_type_; }
- // Set the entire *memory* of the vector to zero.
+ /// Set the entire *memory* of the vector to zero.
void setZero();
private:
@@ -108,6 +108,7 @@ class unified_vector {
size_t buffer_capacity_;
};
+/// Specialization for unified_vector on device memory only.
template
class device_vector : public unified_vector {
public:
@@ -140,6 +141,7 @@ class device_vector : public unified_vector {
}
};
+/// Specialization for unified_vector on host memory only.
template
class host_vector : public unified_vector {
public:
diff --git a/nvblox/include/nvblox/core/voxels.h b/nvblox/include/nvblox/core/voxels.h
index 02fc8b3a..ab4e3199 100644
--- a/nvblox/include/nvblox/core/voxels.h
+++ b/nvblox/include/nvblox/core/voxels.h
@@ -21,31 +21,35 @@ limitations under the License.
namespace nvblox {
+/// A voxel storing TSDF (truncated signed distance field) values.
struct TsdfVoxel {
- // Signed projective distance of the voxel from a surface.
+ /// Signed projective distance of the voxel from a surface.
float distance = 0.0f;
- // How many observations/how confident we are in this observation.
+ /// How many observations/how confident we are in this observation.
float weight = 0.0f;
};
+/// Voxels that stores the distance and full direction to the nearest surface.
struct EsdfVoxel {
// TODO(helen): optimize the memory layout here.
- // Cached squared distance towards the parent.
+ /// Cached squared distance towards the parent.
float squared_distance_vox = 0.0f;
- // Direction towards the parent, *in units of voxels*.
+ /// Direction towards the parent, *in units of voxels*.
Eigen::Vector3i parent_direction = Eigen::Vector3i::Zero();
- // Whether this voxel is inside the surface or not.
+ /// Whether this voxel is inside the surface or not.
bool is_inside = false;
- // Whether this voxel has been observed.
+ /// Whether this voxel has been observed.
bool observed = false;
- // Whether this voxel is a "site": i.e., near the zero-crossing and is
- // eligible to be considered a parent.
+ /// Whether this voxel is a "site": i.e., near the zero-crossing and is
+ /// eligible to be considered a parent.
bool is_site = false;
};
+/// Voxel that stores the color near the surface.
struct ColorVoxel {
+ /// The color!
Color color = Color::Gray();
- // How many observations/how confident we are in this observation.
+ /// How many observations/how confident we are in this observation.
float weight = 0.0f;
};
diff --git a/nvblox/include/nvblox/gpu_hash/cuda/impl/gpu_layer_view_impl.cuh b/nvblox/include/nvblox/gpu_hash/cuda/impl/gpu_layer_view_impl.cuh
index 0d38574f..fb12a72c 100644
--- a/nvblox/include/nvblox/gpu_hash/cuda/impl/gpu_layer_view_impl.cuh
+++ b/nvblox/include/nvblox/gpu_hash/cuda/impl/gpu_layer_view_impl.cuh
@@ -66,6 +66,7 @@ GPULayerView& GPULayerView::operator=(
template
void GPULayerView::reset(LayerType* layer_ptr) {
CHECK_NOTNULL(layer_ptr);
+ timing::Timer timer("gpu_hash/transfer");
// Allocate gpu hash if not allocated already
if (!gpu_hash_ptr_) {
@@ -88,8 +89,6 @@ void GPULayerView::reset(LayerType* layer_ptr) {
CHECK_LT(layer_ptr->numAllocatedBlocks(), gpu_hash_ptr_->max_num_blocks_);
}
- timing::Timer timer("gpu_hash/transfer");
-
gpu_hash_ptr_->impl_.clear();
// This is necessary for bug-free operation, as clear does not sync
@@ -131,7 +130,7 @@ void GPULayerView::reset(LayerType* layer_ptr) {
template
void GPULayerView::reset(size_t new_max_num_blocks) {
- timing::Timer timer("gpu_hash/reallocation");
+ timing::Timer timer("gpu_hash/transfer/reallocation");
gpu_hash_ptr_ = std::make_shared>(new_max_num_blocks);
}
diff --git a/nvblox/include/nvblox/io/csv.h b/nvblox/include/nvblox/io/csv.h
index 151108e7..dcab4ac1 100644
--- a/nvblox/include/nvblox/io/csv.h
+++ b/nvblox/include/nvblox/io/csv.h
@@ -24,6 +24,7 @@ namespace io {
void writeToCsv(const std::string& filepath, const DepthImage& frame);
void writeToCsv(const std::string& filepath, const ColorImage& frame);
+void writeToCsv(const std::string& filepath, const MonoImage& frame);
template
void writeToCsv(const std::string& filepath,
diff --git a/nvblox/include/nvblox/io/impl/pointcloud_io_impl.h b/nvblox/include/nvblox/io/impl/pointcloud_io_impl.h
index ff8a4018..0caa5216 100644
--- a/nvblox/include/nvblox/io/impl/pointcloud_io_impl.h
+++ b/nvblox/include/nvblox/io/impl/pointcloud_io_impl.h
@@ -60,32 +60,5 @@ bool outputVoxelLayerToPly(
return writer.write();
}
-/// Specializations for the TSDF type.
-template <>
-bool outputVoxelLayerToPly(const TsdfLayer& layer,
- const std::string& filename) {
- constexpr float kMinWeight = 0.1f;
- auto lambda = [&kMinWeight](const TsdfVoxel* voxel, float* distance) -> bool {
- *distance = voxel->distance;
- return voxel->weight > kMinWeight;
- };
- return outputVoxelLayerToPly(layer, filename, lambda);
-}
-
-/// Specialization for the ESDF type.
-template <>
-bool outputVoxelLayerToPly(const EsdfLayer& layer,
- const std::string& filename) {
- const float voxel_size = layer.voxel_size();
- auto lambda = [&voxel_size](const EsdfVoxel* voxel, float* distance) -> bool {
- *distance = voxel_size * std::sqrt(voxel->squared_distance_vox);
- if (voxel->is_inside) {
- *distance = -*distance;
- }
- return voxel->observed;
- };
- return outputVoxelLayerToPly(layer, filename, lambda);
-}
-
} // namespace io
} // namespace nvblox
\ No newline at end of file
diff --git a/nvblox/include/nvblox/io/pointcloud_io.h b/nvblox/include/nvblox/io/pointcloud_io.h
index 01318728..f4de4747 100644
--- a/nvblox/include/nvblox/io/pointcloud_io.h
+++ b/nvblox/include/nvblox/io/pointcloud_io.h
@@ -38,14 +38,6 @@ template
bool outputVoxelLayerToPly(const VoxelBlockLayer& layer,
const std::string& filename);
-/// Specializations for the TSDF type.
-template <>
-bool outputVoxelLayerToPly(const TsdfLayer& layer, const std::string& filename);
-
-/// Specialization for the ESDF type.
-template <>
-bool outputVoxelLayerToPly(const EsdfLayer& layer, const std::string& filename);
-
} // namespace io
} // namespace nvblox
diff --git a/nvblox/include/nvblox/mesh/marching_cubes.h b/nvblox/include/nvblox/mesh/marching_cubes.h
index a6714b3b..f0185fd3 100644
--- a/nvblox/include/nvblox/mesh/marching_cubes.h
+++ b/nvblox/include/nvblox/mesh/marching_cubes.h
@@ -30,25 +30,25 @@
namespace nvblox {
namespace marching_cubes {
-// This (internal) struct contains intermediate results of marching cubes.
+/// This (internal) struct contains intermediate results of marching cubes.
struct PerVoxelMarchingCubesResults {
- // The 3D positions of the corners of a 2x2x2 cube of voxels formed by the
- // surrounding voxel neighbours in the positive direction of each coordinate.
+ /// The 3D positions of the corners of a 2x2x2 cube of voxels formed by the
+ /// surrounding voxel neighbours in the positive direction of each coordinate.
Vector3f vertex_coords[8];
- // The value of the TSDF at each of the neighbouring voxels described above.
+ /// The value of the TSDF at each of the neighbouring voxels described above.
float vertex_sdf[8];
- // Does this voxel contain a mesh? (Does it straddle a zero level set?)
+ /// Does this voxel contain a mesh? (Does it straddle a zero level set?)
bool contains_mesh = false;
- // The index into the marching cubes triangle table (found in
- // nvblox/mesh/impl/marching_cubes_table.h). This index is determined based on
- // the tsdf configuration (+/-) of the surrounding voxels and is the main
- // contribution of marching cubes algorithm.
+ /// The index into the marching cubes triangle table (found in
+ /// nvblox/mesh/impl/marching_cubes_table.h). This index is determined based on
+ /// the tsdf configuration (+/-) of the surrounding voxels and is the main
+ /// contribution of marching cubes algorithm.
uint8_t marching_cubes_table_index = 0;
- // At the end of marching cubes, vertices calculated for this and other voxels
- // in this MeshBlock are stored in a single vector. This member indicates where
- // in this vertex vector the vertices associated with this voxel begin. It is
- // calculated through an exclusive prefix sum of the numbers of vertices
- // in each voxel of this MeshBlock.
+ /// At the end of marching cubes, vertices calculated for this and other voxels
+ /// in this MeshBlock are stored in a single vector. This member indicates where
+ /// in this vertex vector the vertices associated with this voxel begin. It is
+ /// calculated through an exclusive prefix sum of the numbers of vertices
+ /// in each voxel of this MeshBlock.
int vertex_vector_start_index;
};
diff --git a/nvblox/include/nvblox/mesh/mesh.h b/nvblox/include/nvblox/mesh/mesh.h
index 4e00d9c1..b67565fe 100644
--- a/nvblox/include/nvblox/mesh/mesh.h
+++ b/nvblox/include/nvblox/mesh/mesh.h
@@ -21,10 +21,10 @@ limitations under the License.
namespace nvblox {
-// A structure which holds a Mesh.
-// Generally produced as the result of fusing MeshBlocks in a Layer
-// into a single mesh.
-// NOTE(alexmillane): Currently only on the CPU.
+/// A structure which holds a combined Mesh for CPU access.
+/// Generally produced as the result of fusing MeshBlocks in a Layer
+/// into a single mesh.
+/// NOTE(alexmillane): Currently only on the CPU.
struct Mesh {
// Data
std::vector vertices;
@@ -32,7 +32,8 @@ struct Mesh {
std::vector triangles;
std::vector colors;
- // Factory
+ /// Create a combined Mesh object from a MeshBlock layer. Useful for mesh
+ /// output.
static Mesh fromLayer(const BlockLayer& layer);
};
diff --git a/nvblox/include/nvblox/mesh/mesh_block.h b/nvblox/include/nvblox/mesh/mesh_block.h
index 468c8167..ffc791e6 100644
--- a/nvblox/include/nvblox/mesh/mesh_block.h
+++ b/nvblox/include/nvblox/mesh/mesh_block.h
@@ -26,17 +26,19 @@ limitations under the License.
namespace nvblox {
-// A mesh block containing all of the triangles from this block.
-// Each block contains only the UPPER part of its neighbors: i.e., the max
-// x, y, and z axes. Its neighbors are responsible for the rest.
+/// A mesh block containing all of the triangles from this block.
+/// Each block contains only the UPPER part of its neighbors: i.e., the max
+/// x, y, and z axes. Its neighbors are responsible for the rest.
struct MeshBlock {
typedef std::shared_ptr Ptr;
typedef std::shared_ptr ConstPtr;
+ /// Create a mesh block of the specified memory type.
MeshBlock(MemoryType memory_type = MemoryType::kDevice);
- // "Clone" copy constructor, with the possibility to change device type.
+ /// "Clone" copy constructor, with the possibility to change device type.
MeshBlock(const MeshBlock& mesh_block);
+ /// "Clone" to a different memory type.
MeshBlock(const MeshBlock& mesh_block, MemoryType memory_type);
// Mesh Data
@@ -49,17 +51,21 @@ struct MeshBlock {
unified_vector colors;
unified_vector triangles;
+ /// Clear all data within the mesh block.
void clear();
- // Resize and reserve.
+ /// Resize vertices and normals to the correct number of vertices.
void resizeToNumberOfVertices(size_t new_size);
+ /// Reserve space in the vertices and normals vectors.
void reserveNumberOfVertices(size_t new_capacity);
+ /// Size of the vertices vector.
size_t size() const;
+ /// Capacity (allocated size) of the vertices vector.
size_t capacity() const;
- // Resize colors/intensities such that:
- // `colors.size()/intensities.size() == vertices.size()`
+ /// Resize colors/intensities such that:
+ /// `colors.size()/intensities.size() == vertices.size()`
void expandColorsToMatchVertices();
// Copy mesh data to the CPU.
@@ -68,13 +74,13 @@ struct MeshBlock {
std::vector getTriangleVectorOnCPU() const;
std::vector getColorVectorOnCPU() const;
- // Note(alexmillane): Memory type ignored, MeshBlocks live in CPU memory.
+ /// Note(alexmillane): Memory type ignored, MeshBlocks live in CPU memory.
static Ptr allocate(MemoryType memory_type);
};
-// Helper struct for mesh blocks on CUDA.
-// NOTE: We need this because we cant pass MeshBlock to kernel functions because
-// of the presence of unified_vector members.
+/// Helper struct for mesh blocks on CUDA.
+/// NOTE: We need this because we can't pass MeshBlock to kernel functions because
+/// of the presence of unified_vector members.
struct CudaMeshBlock {
CudaMeshBlock() = default;
CudaMeshBlock(MeshBlock* block);
diff --git a/nvblox/include/nvblox/mesh/mesh_integrator.h b/nvblox/include/nvblox/mesh/mesh_integrator.h
index 36c2d38f..49395b9c 100644
--- a/nvblox/include/nvblox/mesh/mesh_integrator.h
+++ b/nvblox/include/nvblox/mesh/mesh_integrator.h
@@ -22,6 +22,7 @@ limitations under the License.
namespace nvblox {
+/// Class to integrate TSDF data into a mesh using marching cubes.
class MeshIntegrator {
public:
MeshIntegrator();
@@ -32,18 +33,25 @@ class MeshIntegrator {
const TsdfLayer& distance_layer, BlockLayer* mesh_layer,
const DeviceType device_type = DeviceType::kGPU);
- /// Integrates only the selected blocks from the distance layer.
+ /// Integrates only the selected blocks from the distance layer on the CPU.
+ /// Prefer to use the GPU version.
bool integrateBlocksCPU(const TsdfLayer& distance_layer,
const std::vector& block_indices,
BlockLayer* mesh_layer);
+ /// @brief Integrate new TSDF blocks into a mesh on the GPU.
+ /// @param distance_layer The TSDF layer to integrate.
+ /// @param block_indices Which block indices to integrate, can either be
+ /// updated ones or all blocks in the TSDF layer.
+ /// @param mesh_layer The mesh layer for output.
+ /// @return Whether the integration succeeded.
bool integrateBlocksGPU(const TsdfLayer& distance_layer,
const std::vector& block_indices,
BlockLayer* mesh_layer);
- // Color mesh layer.
- // TODO(alexmillane): Currently these functions color vertices by taking the
- // CLOSEST color. Would be good to have an option at least for interpolation.
+ /// Color mesh layer.
+ /// TODO(alexmillane): Currently these functions color vertices by taking the
+ /// CLOSEST color. Would be good to have an option at least for interpolation.
void colorMesh(const ColorLayer& color_layer, MeshLayer* mesh_layer);
void colorMesh(const ColorLayer& color_layer,
const std::vector& block_indices,
diff --git a/nvblox/include/nvblox/nvblox.h b/nvblox/include/nvblox/nvblox.h
index bb13f1ab..dc3cf4dc 100644
--- a/nvblox/include/nvblox/nvblox.h
+++ b/nvblox/include/nvblox/nvblox.h
@@ -20,9 +20,17 @@ limitations under the License.
#include "nvblox/core/mapper.h"
#include "nvblox/core/layer.h"
#include "nvblox/core/voxels.h"
+#include "nvblox/core/pointcloud.h"
#include "nvblox/integrators/esdf_integrator.h"
#include "nvblox/integrators/projective_color_integrator.h"
#include "nvblox/integrators/projective_integrator_base.h"
#include "nvblox/integrators/projective_tsdf_integrator.h"
+#include "nvblox/mesh/mesh_integrator.h"
+#include "nvblox/mesh/mesh_block.h"
+#include "nvblox/mesh/mesh.h"
#include "nvblox/io/csv.h"
+#include "nvblox/io/layer_cake_io.h"
+#include "nvblox/io/mesh_io.h"
+#include "nvblox/io/ply_writer.h"
+#include "nvblox/io/pointcloud_io.h"
#include "nvblox/mesh/mesh_integrator.h"
diff --git a/nvblox/include/nvblox/primitives/primitives.h b/nvblox/include/nvblox/primitives/primitives.h
index d532914c..b348918a 100644
--- a/nvblox/include/nvblox/primitives/primitives.h
+++ b/nvblox/include/nvblox/primitives/primitives.h
@@ -57,7 +57,7 @@ class Primitive {
public:
enum class Type { kPlane, kCube, kSphere, kCylinder };
- // Epsilon for ray intersection and computation.
+ /// Epsilon for ray intersection and computation.
static constexpr float kEpsilon = 1e-4;
Primitive(const Vector3f& center, Type type)
@@ -84,6 +84,7 @@ class Primitive {
Color color_;
};
+/// Primitive sphere, given a center and a radius.
class Sphere : public Primitive {
public:
Sphere(const Vector3f& center, float radius)
@@ -102,6 +103,8 @@ class Sphere : public Primitive {
float radius_;
};
+/// Primitive cube, given a center and an X,Y,Z size (can be a rectangular
+/// prism).
class Cube : public Primitive {
public:
Cube(const Vector3f& center, const Vector3f& size)
@@ -119,6 +122,7 @@ class Cube : public Primitive {
Vector3f size_;
};
+/// Primitive plane, given a center and a normal.
/// Requires normal being passed in to ALREADY BE NORMALIZED!!!!
class Plane : public Primitive {
public:
@@ -142,7 +146,7 @@ class Plane : public Primitive {
Vector3f normal_;
};
-/// Cylinder centered on the XY plane!
+/// Cylinder centered on the XY plane, with a given radius and height (in Z).
class Cylinder : public Primitive {
public:
Cylinder(const Vector3f& center, float radius, float height)
diff --git a/nvblox/include/nvblox/primitives/scene.h b/nvblox/include/nvblox/primitives/scene.h
index 57a9bead..734bd508 100644
--- a/nvblox/include/nvblox/primitives/scene.h
+++ b/nvblox/include/nvblox/primitives/scene.h
@@ -60,11 +60,15 @@ class Scene {
public:
Scene();
- /// === Creating an environment ===
+ /// Create an environment by adding primitives, which are then owned by the
+ /// scene.
void addPrimitive(std::unique_ptr primitive);
- /// Convenience functions for setting up bounded areas.
+ /// @brief Adds a ground level (normal up) at a given height on the Z axis.
+ /// @param height Height, in meters, of the ground.
void addGroundLevel(float height);
+ /// @brief Adds a ceiling at a given height. All heights on the Z axis.
+ /// @param height Height, in meters, of the ceiling.
void addCeiling(float height);
/// Add 4 walls (infinite planes) bounding the space. In case this is not the
@@ -75,14 +79,14 @@ class Scene {
/// Deletes all objects!
void clear();
- /// === Generating synthetic data from environment ===
/// Generates a synthetic view given camera parameters and a transformation
/// of the camera to the scene.
void generateDepthImageFromScene(const Camera& camera, const Transform& T_S_C,
float max_dist,
DepthImage* depth_frame) const;
- /// === Computing ground truth SDFs ===
+ /// Computes the ground truth SDFs (either TSDF or ESDF depending on template
+ /// parameter).
template
void generateSdfFromScene(float max_dist,
VoxelBlockLayer* layer) const;
diff --git a/nvblox/include/nvblox/rays/ray_caster.h b/nvblox/include/nvblox/rays/ray_caster.h
index fd92acf3..48001c67 100644
--- a/nvblox/include/nvblox/rays/ray_caster.h
+++ b/nvblox/include/nvblox/rays/ray_caster.h
@@ -19,18 +19,19 @@ limitations under the License.
namespace nvblox {
+/// Class for casting rays through a voxelized space from origin to destination.
class RayCaster {
public:
__host__ __device__ inline RayCaster(const Vector3f& origin,
const Vector3f& destination,
float scale = 1.0f);
- // Returns the index, so in "unscaled" coordinates.
+ /// Returns the index, so in "unscaled" coordinates.
__host__ __device__ inline bool nextRayIndex(Index3D* ray_index);
- // Returns scaled coordinates. Just the above multiplied by the scale factor.
+ /// Returns scaled coordinates. Just the above multiplied by the scale factor.
__host__ __device__ inline bool nextRayPositionScaled(Vector3f* ray_position);
- // Just raycasts over the whole thing and puts them in a vector for you.
+ /// Just raycasts over the whole thing and puts them in a vector for you.
__host__ void getAllIndices(std::vector* indices);
private:
diff --git a/nvblox/include/nvblox/semantics/image_masker.h b/nvblox/include/nvblox/semantics/image_masker.h
new file mode 100644
index 00000000..ad240110
--- /dev/null
+++ b/nvblox/include/nvblox/semantics/image_masker.h
@@ -0,0 +1,106 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+#include "nvblox/core/camera.h"
+#include "nvblox/core/image.h"
+#include "nvblox/core/types.h"
+#include "nvblox/utils/timing.h"
+
+namespace nvblox {
+
+/// A class to mask images based on a binary mask.
+class ImageMasker {
+ public:
+ ImageMasker();
+ ~ImageMasker();
+
+ /// A parameter getter
+ /// The occlusion threshold parameter associated with the image splitter.
+ /// A point is considered to be occluded on the mask image only if it lies
+ /// more than the occlusion threshold behind the point occluding it.
+ /// Occluded points are always assumed to be unmasked.
+ /// @returns the occlusion threshold in meters
+ float occlusion_threshold() const;
+
+ /// A parameter setter
+ /// See occlusion_threshold().
+ /// @param occlusion_threshold the occlusion threshold in meters.
+ void occlusion_threshold(float occlusion_threshold);
+
+ /// Splitting the input depth image according to a mask image
+ /// assuming depth and mask images are coming from the same camera.
+ ///@param input Depth image to be split according to mask.
+ ///@param mask Mask image.
+ ///@param unmasked_output Depth image containing the depth values
+ /// of all unmasked input pixels.
+ /// Masked pixels are set to -1.
+ ///@param masked_output Depth image containing the depth values
+ /// of all masked input pixels.
+ /// Unmasked pixels are set to -1.
+ void splitImageOnGPU(const DepthImage& input, const MonoImage& mask,
+ DepthImage* unmasked_output, DepthImage* masked_output);
+
+ /// Splitting the input color image according to a mask image
+ /// assuming color and mask images are coming from the same camera.
+ ///@param input Color image to be split according to mask.
+ ///@param mask Mask image.
+ ///@param unmasked_output Color image containing the color values of all
+ /// unmasked input pixels.
+ /// Masked pixels are set to black.
+ ///@param masked_output Color image containing the color values
+ /// of all masked input pixels.
+ /// Unmasked pixels are set to black.
+ void splitImageOnGPU(const ColorImage& input, const MonoImage& mask,
+ ColorImage* unmasked_output, ColorImage* masked_output);
+
+ /// Splitting the input depth image according to a mask image taking occlusion
+ /// into account.
+ ///@param depth_input Depth image to be split according to mask.
+ ///@param mask Mask image.
+ ///@param T_CM_CD Transform from depth camera to mask camera frame.
+ ///@param depth_camera Intrinsics model of the depth camera.
+ ///@param mask_camera Intrinsics model of the mask camera.
+ ///@param unmasked_depth_output Depth image containing the depth values
+ /// of all unmasked input pixels.
+ /// Masked pixels are set to -1.
+ ///@param masked_depth_output Depth image containing the depth values
+ /// of all masked input pixels.
+ /// Unmasked pixels are set to -1.
+ void splitImageOnGPU(const DepthImage& depth_input, const MonoImage& mask,
+ const Transform& T_CM_CD, const Camera& depth_camera,
+ const Camera& mask_camera,
+ DepthImage* unmasked_depth_output,
+ DepthImage* masked_depth_output);
+
+ private:
+ // Templatized internal version
+ template
+ void splitImageOnGPUTemplate(const ImageType& input, const MonoImage& mask,
+ ImageType* unmasked_output,
+ ImageType* masked_output);
+
+ template
+ void allocateOutput(const ImageType& input, ImageType* unmasked_output,
+ ImageType* masked_output);
+
+ // Params
+ float occlusion_threshold_m_ = 0.25;
+
+ cudaStream_t cuda_stream_;
+};
+
+} // namespace nvblox
diff --git a/nvblox/include/nvblox/semantics/image_projector.h b/nvblox/include/nvblox/semantics/image_projector.h
new file mode 100644
index 00000000..5b83a3ae
--- /dev/null
+++ b/nvblox/include/nvblox/semantics/image_projector.h
@@ -0,0 +1,58 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#pragma once
+
+#include "nvblox/core/camera.h"
+#include "nvblox/core/image.h"
+#include "nvblox/core/pointcloud.h"
+#include "nvblox/core/types.h"
+#include "nvblox/core/unified_ptr.h"
+#include "nvblox/core/unified_vector.h"
+
+namespace nvblox {
+
+/// A class which takes care of back projecting images on the GPU.
+class DepthImageBackProjector {
+ public:
+ DepthImageBackProjector();
+ ~DepthImageBackProjector();
+
+ /// Back projects a depth image to a pointcloud in the camera frame.
+ ///@param image DepthImage to be back projected
+ ///@param camera Pinhole camera intrinsics model
+ ///@param pointcloud_C Pointer to the output pointcloud. Must be in either
+ /// device or unified memory.
+ void backProjectOnGPU(const DepthImage& image, const Camera& camera,
+ Pointcloud* pointcloud_C_ptr);
+
+ /// Takes a collection of points, and returns the center of the voxels that
+ /// contain the points. Note that this function deletes duplicates, that is
+ /// that the output pointcloud may have less points than the input pointcloud.
+ ///@param pointcloud_L A collection of points in the Layer frame.
+ ///@param voxel_size The side length of the voxels in the layer.
+ ///@param voxel_center_pointcloud_L Voxel centers stored as a pointcloud.
+ void pointcloudToVoxelCentersOnGPU(const Pointcloud& pointcloud_L,
+ float voxel_size,
+ Pointcloud* voxel_center_pointcloud_L);
+
+ private:
+ cudaStream_t cuda_stream_ = nullptr;
+
+ unified_ptr pointcloud_size_device_;
+ unified_ptr pointcloud_size_host_;
+};
+
+} // namespace nvblox
\ No newline at end of file
diff --git a/nvblox/include/nvblox/serialization/common_types.h b/nvblox/include/nvblox/serialization/common_types.h
index 69e90c46..1297f4ca 100644
--- a/nvblox/include/nvblox/serialization/common_types.h
+++ b/nvblox/include/nvblox/serialization/common_types.h
@@ -22,12 +22,12 @@ limitations under the License.
namespace nvblox {
-// Bind default functions for the given types.
+/// Bind default functions for the given types.
template
LayerSerializationFunctions bindDefaultFunctions();
-// Registers all the common layer types for serialization. Must be called before
-// you serialize or de-serialize anything.
+/// Registers all the common layer types for serialization. Must be called before
+/// you serialize or de-serialize anything.
inline void registerCommonTypes();
} // namespace nvblox
diff --git a/nvblox/include/nvblox/serialization/layer_type_register.h b/nvblox/include/nvblox/serialization/layer_type_register.h
index f5e7be9e..20ac283c 100644
--- a/nvblox/include/nvblox/serialization/layer_type_register.h
+++ b/nvblox/include/nvblox/serialization/layer_type_register.h
@@ -26,12 +26,15 @@ limitations under the License.
namespace nvblox {
+/// Struct of various layer parameters of different types.
struct LayerParameterStruct {
std::map string_params;
std::map int_params;
std::map float_params;
};
+/// Struct holding callbacks for serialization functions for various layer
+/// types.
struct LayerSerializationFunctions {
// Serialization functions.
typedef std::function
@@ -69,23 +72,24 @@ class LayerTypeRegister {
// No constructor!
LayerTypeRegister() = delete;
- // Register a type and how to construct one.
+ /// Register a type and how to construct one.
static void registerType(
const std::string& type_name, const std::type_index& type_index,
const LayerSerializationFunctions& serialization_functions);
- // Create a layer of a given type.
+ /// Create a layer of a given type.
static std::unique_ptr createLayer(
const std::string& type_name, MemoryType memory_type,
const LayerParameterStruct& layer_params);
- // Get the struct of function callbacks.
+ /// Get the struct of function callbacks.
static LayerSerializationFunctions getSerializationFunctions(
const std::string& type_name);
- // Get the name of the layer based on its type index.
+ /// Get the name of the layer based on its type index.
static std::string getLayerName(const std::type_index& type_index);
+ /// Look up the type index based on a layer name.
static std::type_index getLayerTypeIndex(const std::string& type_name);
private:
diff --git a/nvblox/include/nvblox/serialization/serializer.h b/nvblox/include/nvblox/serialization/serializer.h
index 8893abc5..72678d80 100644
--- a/nvblox/include/nvblox/serialization/serializer.h
+++ b/nvblox/include/nvblox/serialization/serializer.h
@@ -25,6 +25,7 @@ limitations under the License.
namespace nvblox {
+/// Class to serialize and read a layer cake from an SQLite database.
class Serializer {
public:
/// Default constructor for invalid file. Must call open before using.
@@ -44,26 +45,29 @@ class Serializer {
bool open(const std::string& filename,
std::ios_base::openmode openmode = std::ios::in);
+ /// Load a layer cake from the opened file of a given memory type.
LayerCake loadLayerCake(MemoryType memory_type);
+
+ /// Write out a layer cake to the opened file, return success.
bool writeLayerCake(const LayerCake& cake);
/// Close the file.
bool close();
- // Below use only if you know WTF you're doing!
+ // ======= Below use only if you know WTF you're doing! =====================
- // Create a layer table & metadata table.
+ /// Create a layer table & metadata table.
bool createLayerTables(const std::string& layer_name);
- // Serialize layer parameters.
+ /// Serialize layer parameters.
bool setLayerParameters(const std::string& layer_name,
const LayerParameterStruct& layer_params);
- // Deserialize layer parameters.
+ /// Deserialize layer parameters.
bool getLayerParameters(const std::string& layer_name,
LayerParameterStruct* layer_params);
- // Write a new data parameter to the table.
+ /// Write a new data parameter to the table.
bool addLayerData(const std::string& layer_name, const Index3D& index,
const std::vector& data);
diff --git a/nvblox/include/nvblox/serialization/sqlite_database.h b/nvblox/include/nvblox/serialization/sqlite_database.h
index d5487e68..2df09841 100644
--- a/nvblox/include/nvblox/serialization/sqlite_database.h
+++ b/nvblox/include/nvblox/serialization/sqlite_database.h
@@ -22,6 +22,8 @@ limitations under the License.
struct sqlite3;
namespace nvblox {
+/// Class to wrap access to the C interface of SQLite in a slightly more
+/// usable format.
class SqliteDatabase {
public:
/// Default constructor for invalid file. Must call open before using.
@@ -39,22 +41,23 @@ class SqliteDatabase {
/// Close the file.
bool close();
- // Run a statement that does not have a return value.
+ /// Run a statement that does not have a return value.
bool runStatement(const std::string& statement);
+ /// Run a return-value-less statement on a byte blob.
bool runStatementWithBlob(const std::string& statement,
const std::vector& blob);
- // Run a query that has a SINGLE return value of the given type:
+ /// Run a query that has a SINGLE return value of the given type:
bool runSingleQueryString(const std::string& sql_query, std::string* result);
bool runSingleQueryInt(const std::string& sql_query, int* result);
bool runSingleQueryFloat(const std::string& sql_query, float* result);
bool runSingleQueryBlob(const std::string& sql_query,
std::vector* result);
- // Returns MULTIPLE values.
+ /// Returns MULTIPLE values.
bool runMultipleQueryString(const std::string& sql_query,
std::vector* result);
- // Return multiple values in an index.
+ /// Return multiple values in an index.
bool runMultipleQueryIndex3D(const std::string& sql_query,
std::vector* result);
diff --git a/nvblox/include/nvblox/utils/nvtx_ranges.h b/nvblox/include/nvblox/utils/nvtx_ranges.h
index 6be20dc3..e2c6be45 100644
--- a/nvblox/include/nvblox/utils/nvtx_ranges.h
+++ b/nvblox/include/nvblox/utils/nvtx_ranges.h
@@ -24,6 +24,8 @@ limitations under the License.
namespace nvblox {
namespace timing {
+/// Instrument our timers with NvtxRanges, which can be visualized in Nsight
+/// Systems to aid with debugging and profiling.
class NvtxRange {
public:
NvtxRange(const std::string& message, const Color& color,
diff --git a/nvblox/src/core/color.cpp b/nvblox/src/core/color.cpp
index 3132cc3c..d61a9161 100644
--- a/nvblox/src/core/color.cpp
+++ b/nvblox/src/core/color.cpp
@@ -31,6 +31,8 @@ Color Color::blendTwoColors(const Color& first_color, float first_weight,
first_color.g * first_weight + second_color.g * second_weight));
new_color.b = static_cast(std::round(
first_color.b * first_weight + second_color.b * second_weight));
+ new_color.a = static_cast(std::round(
+ first_color.a * first_weight + second_color.a * second_weight));
return new_color;
}
diff --git a/nvblox/src/core/cuda/layer.cu b/nvblox/src/core/cuda/layer.cu
new file mode 100644
index 00000000..c9b3515c
--- /dev/null
+++ b/nvblox/src/core/cuda/layer.cu
@@ -0,0 +1,36 @@
+/*
+Copyright 2022 NVIDIA CORPORATION
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+*/
+#include "nvblox/core/cuda/impl/layer_impl.cuh"
+#include "nvblox/core/layer.h"
+
+namespace nvblox {
+
+template void VoxelBlockLayer::getVoxelsGPU(
+ const device_vector& positions_L,
+ device_vector* voxels_ptr,
+ device_vector* success_flags_ptr) const;
+
+template void VoxelBlockLayer::getVoxelsGPU(
+ const device_vector& positions_L,
+ device_vector* voxels_ptr,
+ device_vector* success_flags_ptr) const;
+
+template void VoxelBlockLayer::getVoxelsGPU(
+ const device_vector& positions_L,
+ device_vector* voxels_ptr,
+ device_vector* success_flags_ptr) const;
+
+} // namespace nvblox
diff --git a/nvblox/src/core/cuda/pointcloud.cu b/nvblox/src/core/cuda/pointcloud.cu
new file mode 100644
index 00000000..b2cb45f1
--- /dev/null
+++ b/nvblox/src/core/cuda/pointcloud.cu
@@ -0,0 +1,71 @@
+#include "nvblox/core/pointcloud.h"
+
+#include
+
+#include "nvblox/core/cuda/error_check.cuh"
+
+namespace nvblox {
+
+Pointcloud::Pointcloud(int num_points, MemoryType memory_type)
+ : points_(unified_vector(num_points, memory_type)) {}
+
+Pointcloud::Pointcloud(MemoryType memory_type) {}
+
+Pointcloud::Pointcloud(const Pointcloud& other)
+ : points_(other.points_, other.memory_type()) {
+ LOG(WARNING) << "Deep copy of Pointcloud.";
+}
+
+Pointcloud::Pointcloud(const Pointcloud& other, MemoryType memory_type)
+ : points_(other.points_, memory_type) {}
+
+Pointcloud& Pointcloud::operator=(const Pointcloud& other) {
+ LOG(WARNING) << "Deep copy of Pointcloud.";
+ points_ = unified_vector(other.points_, other.memory_type());
+ return *this;
+}
+
+Pointcloud::Pointcloud(const std::vector& points,
+ MemoryType memory_type)
+ : points_(points, memory_type) {}
+
+// Pointcloud operations
+
+__global__ void transformPointcloudKernel(const Transform T_out_in,
+ int pointcloud_size,
+ const Vector3f* pointcloud_in,
+ Vector3f* pointcloud_out) {
+ const int index = threadIdx.x + blockIdx.x * blockDim.x;
+ if (index >= pointcloud_size) {
+ return;
+ }
+
+ pointcloud_out[index] = T_out_in * pointcloud_in[index];
+}
+
+void transformPointcloudOnGPU(const Transform& T_out_in,
+ const Pointcloud& pointcloud_in,
+ Pointcloud* pointcloud_out_ptr) {
+ CHECK_NOTNULL(pointcloud_out_ptr);
+ CHECK(pointcloud_out_ptr->memory_type() == MemoryType::kDevice ||
+ pointcloud_out_ptr->memory_type() == MemoryType::kUnified);
+
+ if (pointcloud_in.empty()) {
+ return;
+ }
+ pointcloud_out_ptr->resize(pointcloud_in.size());
+
+ cudaStream_t cuda_stream;
+ checkCudaErrors(cudaStreamCreate(&cuda_stream));
+ constexpr int kThreadsPerThreadBlock = 512;
+ const int num_blocks(pointcloud_in.size() / kThreadsPerThreadBlock + 1);
+ transformPointcloudKernel<<>>(T_out_in, pointcloud_in.size(),
+ pointcloud_in.dataConstPtr(),
+ pointcloud_out_ptr->dataPtr());
+ checkCudaErrors(cudaStreamSynchronize(cuda_stream));
+ checkCudaErrors(cudaPeekAtLastError());
+ checkCudaErrors(cudaStreamDestroy(cuda_stream));
+}
+
+} // namespace nvblox
diff --git a/nvblox/src/core/mapper.cpp b/nvblox/src/core/mapper.cpp
index d5351fe8..df04596f 100644
--- a/nvblox/src/core/mapper.cpp
+++ b/nvblox/src/core/mapper.cpp
@@ -202,4 +202,56 @@ bool RgbdMapper::loadMap(const std::string& filename) {
return true;
}
+// Human stuff below
+
+HumanMapper::HumanMapper(float voxel_size_m, MemoryType memory_type)
+ : RgbdMapper(voxel_size_m, memory_type) {}
+
+void HumanMapper::integrateDepth(const DepthImage& depth_frame,
+ const MonoImage& mask, const Transform& T_L_C,
+ const Camera& camera) {
+ // Remove humans from the depth image
+ image_masker_.splitImageOnGPU(depth_frame, mask, &depth_frame_no_humans_,
+ &depth_frame_only_humans_);
+ // Do normal integration with the humans removed
+ RgbdMapper::integrateDepth(depth_frame_no_humans_, T_L_C, camera);
+}
+
+void HumanMapper::integrateDepth(const DepthImage& depth_frame,
+ const MonoImage& mask, const Transform& T_L_CD,
+ const Transform& T_CM_CD,
+ const Camera& depth_camera,
+ const Camera& mask_camera) {
+ // Remove humans from the depth image
+ image_masker_.splitImageOnGPU(depth_frame, mask, T_CM_CD, depth_camera,
+ mask_camera, &depth_frame_no_humans_,
+ &depth_frame_only_humans_);
+
+ // Do normal integration with the humans removed
+ RgbdMapper::integrateDepth(depth_frame_no_humans_, T_L_CD, depth_camera);
+}
+
+void HumanMapper::integrateColor(const ColorImage& color_frame,
+ const MonoImage& mask, const Transform& T_L_C,
+ const Camera& camera) {
+ // Remove humans from the color image
+ image_masker_.splitImageOnGPU(color_frame, mask, &color_frame_no_humans_,
+ &color_frame_only_humans_);
+ // Do normal integration with the humans removed
+ RgbdMapper::integrateColor(color_frame_no_humans_, T_L_C, camera);
+}
+
+const DepthImage& HumanMapper::getLastDepthFrameWithoutHumans() {
+ return depth_frame_no_humans_;
+}
+const DepthImage& HumanMapper::getLastDepthFrameOnlyHumans() {
+ return depth_frame_only_humans_;
+}
+const ColorImage& HumanMapper::getLastColorFrameWithoutHumans() {
+ return color_frame_no_humans_;
+}
+const ColorImage& HumanMapper::getLastColorFrameOnlyHumans() {
+ return color_frame_only_humans_;
+}
+
} // namespace nvblox
diff --git a/nvblox/src/integrators/cuda/esdf_integrator.cu b/nvblox/src/integrators/cuda/esdf_integrator.cu
index 842e3f3e..14010653 100644
--- a/nvblox/src/integrators/cuda/esdf_integrator.cu
+++ b/nvblox/src/integrators/cuda/esdf_integrator.cu
@@ -15,6 +15,7 @@ limitations under the License.
*/
#include "nvblox/core/bounding_boxes.h"
#include "nvblox/core/bounding_spheres.h"
+#include "nvblox/core/cuda/atomic_float.cuh"
#include "nvblox/gpu_hash/cuda/gpu_hash_interface.cuh"
#include "nvblox/gpu_hash/cuda/gpu_indexing.cuh"
#include "nvblox/gpu_hash/cuda/gpu_set.cuh"
@@ -189,7 +190,7 @@ __global__ void markAllSitesCombinedKernel(
}
__syncthreads();
-
+
if (threadIdx.x == 1 && threadIdx.y == 1 && threadIdx.z == 1) {
if (updated) {
updated_vec[atomicAdd(updated_vec_size, 1)] = block_indices[block_idx];
@@ -200,18 +201,6 @@ __global__ void markAllSitesCombinedKernel(
}
}
-// From:
-// https://stackoverflow.com/questions/17399119/how-do-i-use-atomicmax-on-floating-point-values-in-cuda
-__device__ __forceinline__ float atomicMinFloat(float* addr, float value) {
- float old;
- old = (value >= 0)
- ? __int_as_float(atomicMin((int*)addr, __float_as_int(value)))
- : __uint_as_float(
- atomicMax((unsigned int*)addr, __float_as_uint(value)));
-
- return old;
-}
-
/// Thread size MUST be 8x8x8, block size can be anything.
__global__ void markSitesInSliceCombinedKernel(
int num_blocks, Index3D* block_indices,
@@ -665,6 +654,14 @@ void EsdfIntegrator::markSitesInSliceCombined(
cleared_counter_device_.get());
checkCudaErrors(cudaStreamSynchronize(cuda_stream_));
checkCudaErrors(cudaPeekAtLastError());
+
+ timing::Timer pack_out_timer("esdf/integrate/mark_sites/pack_out");
+ updated_counter_device_.copyTo(updated_counter_host_);
+ cleared_counter_device_.copyTo(cleared_counter_host_);
+
+ updated_blocks->resize(*updated_counter_host_);
+ cleared_blocks->resize(*cleared_counter_host_);
+ pack_out_timer.Stop();
}
__host__ __device__ void getDirectionAndVoxelIndicesFromThread(
diff --git a/nvblox/src/integrators/cuda/projective_color_integrator.cu b/nvblox/src/integrators/cuda/projective_color_integrator.cu
index 684531ef..2d70239e 100644
--- a/nvblox/src/integrators/cuda/projective_color_integrator.cu
+++ b/nvblox/src/integrators/cuda/projective_color_integrator.cu
@@ -15,6 +15,7 @@ limitations under the License.
*/
#include "nvblox/integrators/projective_color_integrator.h"
+#include "nvblox/core/interpolation_2d.h"
#include "nvblox/integrators/internal/cuda/projective_integrators_common.cuh"
#include "nvblox/integrators/internal/integrators_common.h"
#include "nvblox/utils/timing.h"
@@ -195,8 +196,9 @@ __global__ void integrateBlocks(
}
Color image_value;
- if (!interpolation::interpolate2DLinear(color_image, u_px, color_rows,
- color_cols, &image_value)) {
+ if (!interpolation::interpolate2DLinear<
+ Color, interpolation::checkers::ColorPixelAlphaGreaterThanZero>(
+ color_image, u_px, color_rows, color_cols, &image_value)) {
return;
}
diff --git a/nvblox/src/integrators/cuda/view_calculator.cu b/nvblox/src/integrators/cuda/view_calculator.cu
index 08bc589e..1af2124f 100644
--- a/nvblox/src/integrators/cuda/view_calculator.cu
+++ b/nvblox/src/integrators/cuda/view_calculator.cu
@@ -213,7 +213,8 @@ std::vector ViewCalculator::getBlocksInImageViewRaycastTemplate(
const DepthImage& depth_frame, const Transform& T_L_C,
const SensorType& camera, const float block_size,
const float truncation_distance_m, const float max_integration_distance_m) {
- timing::Timer setup_timer("in_view/setup");
+ timing::Timer total_timer("view_calculator/raycast");
+ timing::Timer setup_timer("view_calculator/raycast/setup");
// Aight so first we have to get the AABB of this guy.
const AxisAlignedBoundingBox aabb_L =
@@ -257,7 +258,7 @@ std::vector ViewCalculator::getBlocksInImageViewRaycastTemplate(
}
// Output vector.
- timing::Timer output_timer("in_view/output");
+ timing::Timer output_timer("view_calculator/raycast/output");
cudaMemcpyAsync(aabb_host_buffer_.data(), aabb_device_buffer_.data(),
sizeof(bool) * aabb_linear_size, cudaMemcpyDeviceToHost,
cuda_stream_);
@@ -270,9 +271,6 @@ std::vector ViewCalculator::getBlocksInImageViewRaycastTemplate(
&output_vector);
output_timer.Stop();
- // We have to manually destruct this. :(
- timing::Timer destory_timer("in_view/destroy");
- destory_timer.Stop();
return output_vector;
}
@@ -314,7 +312,7 @@ void ViewCalculator::getBlocksByRaycastingCorners(
dim3 block_dim(rounded_rows, rounded_cols);
dim3 thread_dim(kThreadDim, kThreadDim);
- timing::Timer image_blocks_timer("in_view/get_image_blocks");
+ timing::Timer image_blocks_timer("view_calculator/raycast/get_image_blocks");
getBlockIndicesInImageKernel<<>>(
T_L_C, camera, depth_frame.dataConstPtr(), depth_frame.rows(),
depth_frame.cols(), block_size, max_integration_distance_m,
@@ -324,7 +322,7 @@ void ViewCalculator::getBlocksByRaycastingCorners(
image_blocks_timer.Stop();
- timing::Timer image_blocks_copy_timer("in_view/image_blocks_copy");
+ timing::Timer image_blocks_copy_timer("view_calculator/raycast/image_blocks_copy");
unified_vector initial_vector;
const size_t aabb_linear_size = aabb_size.x() * aabb_size.y() * aabb_size.z();
@@ -335,7 +333,7 @@ void ViewCalculator::getBlocksByRaycastingCorners(
image_blocks_copy_timer.Stop();
// Call the kernel to do raycasting.
- timing::Timer raycast_blocks_timer("in_view/raycast_blocks");
+ timing::Timer raycast_blocks_timer("view_calculator/raycast/raycast_kernel");
int num_initial_blocks = initial_vector.size();
constexpr int kNumCorners = 9;
@@ -359,6 +357,7 @@ void ViewCalculator::getBlocksByRaycastingPixels(
const float truncation_distance_m, const float max_integration_distance_m,
const Index3D& min_index, const Index3D& aabb_size,
bool* aabb_updated_cuda) {
+ timing::Timer combined_kernel_timer("view_calculator/raycast/raycast_pixels_kernel");
// Number of rays per dimension. Depth frame size / subsampling rate.
const int num_subsampled_rows =
std::ceil(static_cast(depth_frame.rows() + 1) /
@@ -377,7 +376,6 @@ void ViewCalculator::getBlocksByRaycastingPixels(
dim3 block_dim(rounded_rows, rounded_cols);
dim3 thread_dim(kThreadDim, kThreadDim);
- timing::Timer combined_kernel_timer("in_view/combined_kernel");
combinedBlockIndicesInImageKernel<<>>(
T_L_C, camera, depth_frame.dataConstPtr(), depth_frame.rows(),
depth_frame.cols(), block_size, max_integration_distance_m,
diff --git a/nvblox/src/integrators/view_calculator.cpp b/nvblox/src/integrators/view_calculator.cpp
index ffb86007..33b0ec8c 100644
--- a/nvblox/src/integrators/view_calculator.cpp
+++ b/nvblox/src/integrators/view_calculator.cpp
@@ -16,6 +16,7 @@ limitations under the License.
#include "nvblox/integrators/view_calculator.h"
#include "nvblox/core/bounding_boxes.h"
+#include "nvblox/utils/timing.h"
namespace nvblox {
@@ -23,6 +24,7 @@ std::vector