Merge pull request #2670 from matkatz/align_filter_cuda_impl

Align processing block - CUDA implementation
IntelRealSense · Nov 12, 2018 · c48c88c · c48c88c
2 parents 00b89e0 + 15312ef
commit c48c88c
Show file tree

Hide file tree

Showing 15 changed files with 1,398 additions and 993 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -188,6 +188,7 @@ set(REALSENSE_CPP
     src/backend.cpp
     src/verify.c
     src/software-device.cpp
+    src/proc/processing-blocks-factory.cpp
     src/proc/align.cpp
     src/proc/colorizer.cpp
     src/proc/pointcloud.cpp
@@ -199,6 +200,9 @@ set(REALSENSE_CPP
     src/proc/temporal-filter.cpp
     src/proc/hole-filling-filter.cpp
     src/proc/disparity-transform.cpp
+
+    src/proc/sse/sse-align.cpp
+
     src/source.cpp
     src/ds5/ds5-options.cpp
     src/ds5/ds5-timestamp.cpp
@@ -262,11 +266,15 @@ if (BUILD_WITH_CUDA)
     set(REALSENSE_CU
         src/cuda/cuda-conversion.cu
         src/cuda/cuda-pointcloud.cu
+        src/proc/cuda/cuda-align.cu
         )
 
     set(REALSENSE_CUH
+        src/cuda/rscuda_utils.cuh
         src/cuda/cuda-conversion.cuh
         src/cuda/cuda-pointcloud.cuh
+        src/proc/cuda/cuda-align.cuh
+        src/proc/cuda/cuda-align.h
         )
 endif()
 
@@ -325,6 +333,7 @@ set(REALSENSE_HPP
     src/sync.h
     src/sensor.h
     src/stream.h
+    src/proc/processing-blocks-factory.h
     src/proc/align.h
     src/proc/colorizer.h
     src/proc/pointcloud.h
@@ -336,6 +345,9 @@ set(REALSENSE_HPP
     src/proc/hole-filling-filter.h
     src/proc/syncer-processing-block.h
     src/proc/disparity-transform.h
+
+    src/proc/sse/sse-align.h
+
     src/algo.h
     src/option.h
     src/metadata.h
@@ -610,6 +622,7 @@ if(WIN32)
         )
 
     source_group("Source Files\\Processing Blocks" FILES
+        src/proc/processing-blocks-factory.cpp
         src/proc/colorizer.cpp
         src/proc/synthetic-stream.cpp
         src/proc/align.cpp
@@ -624,6 +637,7 @@ if(WIN32)
         )
 
     source_group("Header Files\\Processing Blocks" FILES
+        src/proc/processing-blocks-factory.h
         src/proc/colorizer.h
         src/proc/align.h
         src/proc/pointcloud.h
@@ -637,6 +651,24 @@ if(WIN32)
         src/proc/hole-filling-filter.h
         )
 
+    source_group("Source Files\\Processing Blocks\\sse" FILES
+        src/proc/sse/sse-align.cpp
+    )
+
+    source_group("Header Files\\Processing Blocks\\sse" FILES
+        src/proc/sse/sse-align.h
+    )
+
+if (BUILD_WITH_CUDA)
+    source_group("Source Files\\Processing Blocks\\cuda" FILES
+        src/proc/cuda/cuda-align.cu
+    )
+
+    source_group("Header Files\\Processing Blocks\\cuda" FILES
+        src/proc/cuda/cuda-align.cuh
+        src/proc/cuda/cuda-align.h
+    )
+endif()
     if(BUILD_WITH_STATIC_CRT)
         foreach(flag_var
                 CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_DEBUG CMAKE_CXX_FLAGS_RELEASE
@@ -683,21 +715,21 @@ if(FORCE_WINUSB_UVC)
     set(BACKEND RS2_USE_WINUSB_UVC_BACKEND)
 #
     list(APPEND REALSENSE_CPP
-		src/win7/winusb_uvc/winusb_uvc.cpp
+        src/win7/winusb_uvc/winusb_uvc.cpp
     )
 #
     list(APPEND REALSENSE_HPP
-		src/libuvc/utlist.h
-		src/win7/winusb_uvc/winusb_uvc.h
+        src/libuvc/utlist.h
+        src/win7/winusb_uvc/winusb_uvc.h
     )
 #
-	source_group("Source Files\\WinUsbUvc" FILES
-		src/win7/winusb_uvc/winusb_uvc.cpp
+    source_group("Source Files\\WinUsbUvc" FILES
+        src/win7/winusb_uvc/winusb_uvc.cpp
     )
 #
     source_group("Header Files\\WinUsbUvc" FILES
-		src/libuvc/utlist.h
-		src/win7/winusb_uvc/winusb_uvc.h
+        src/libuvc/utlist.h
+        src/win7/winusb_uvc/winusb_uvc.h
     )
 #
 #    message( WARNING "Using winusb_uvc!" )

diff --git a/include/librealsense2/rsutil.h b/include/librealsense2/rsutil.h
@@ -14,7 +14,14 @@
 #include <stdint.h>
 #include <math.h>
 
+#ifdef RS2_USE_CUDA
+   // CUDA headers
+#include <cuda_runtime.h>
+#endif
 
+#ifdef RS2_USE_CUDA
+__device__
+#endif
 /* Given a point in 3D space, compute the corresponding pixel coordinates in an image with no distortion or forward distortion coefficients produced by the same camera */
 static void rs2_project_point_to_pixel(float pixel[2], const struct rs2_intrinsics * intrin, const float point[3])
 {
@@ -46,6 +53,9 @@ static void rs2_project_point_to_pixel(float pixel[2], const struct rs2_intrinsi
     pixel[1] = y * intrin->fy + intrin->ppy;
 }
 
+#ifdef RS2_USE_CUDA
+__device__
+#endif
 /* Given pixel coordinates and depth in an image with no distortion or inverse distortion coefficients, compute the corresponding point in 3D space relative to the same camera */
 static void rs2_deproject_pixel_to_point(float point[3], const struct rs2_intrinsics * intrin, const float pixel[2], float depth)
 {
@@ -69,6 +79,9 @@ static void rs2_deproject_pixel_to_point(float point[3], const struct rs2_intrin
     point[2] = depth;
 }
 
+#ifdef RS2_USE_CUDA
+__device__
+#endif
 /* Transform 3D coordinates relative to one sensor to 3D coordinates relative to another viewpoint */
 static void rs2_transform_point_to_point(float to_point[3], const struct rs2_extrinsics * extrin, const float from_point[3])
 {
@@ -113,6 +126,9 @@ static void adjust_2D_point_to_boundary(float p[2], int width, int height)
     if (p[1] > height) p[1] = height;
 }
 
+#ifdef RS2_USE_CUDA
+__device__
+#endif
 /* Find projected pixel with unknown depth search along line. */
 static void rs2_project_color_pixel_to_depth_pixel(float to_pixel[2],
     const uint16_t* data, float depth_scale,