Completed Vulkan boilerplate code.

At the moment, most of it just tests copying memory and running kernels.
This should be enough to confirm that data transfers work as expected on
a real GPU.

src/correlation.rs

@@ -4,9 +4,9 @@ mod metal;
 mod vk;
 
 #[cfg(target_os = "macos")]
-use metal as gpu;
+type GpuContext = metal::GpuContext;
 #[cfg(not(target_os = "macos"))]
-use vk as gpu;
+type GpuContext = vk::GpuContext;
 
 use crate::data::{Grid, Point2D};
 use nalgebra::{Matrix3, Vector3};
@@ -67,10 +67,11 @@ pub struct PointCorrelations {
     correlation_threshold: f32,
     corridor_extend_range: f64,
     fundamental_matrix: Matrix3<f64>,
-    gpu_context: Option<gpu::GpuContext>,
+    gpu_context: Option<GpuContext>,
     selected_hardware: String,
 }
 
+#[derive(Copy, Clone)]
 enum CorrelationDirection {
     Forward,
     Reverse,
@@ -141,7 +142,7 @@ impl PointCorrelations {
         let gpu_context = if matches!(hardware_mode, HardwareMode::Gpu | HardwareMode::GpuLowPower)
         {
             let low_power = matches!(hardware_mode, HardwareMode::GpuLowPower);
-            match gpu::GpuContext::new(
+            match GpuContext::new(
                 (img1_dimensions.0 as usize, img1_dimensions.1 as usize),
                 (img2_dimensions.0 as usize, img2_dimensions.1 as usize),
                 projection_mode,
@@ -224,8 +225,8 @@ impl PointCorrelations {
             CorrelationDirection::Reverse,
         )?;
 
-        self.cross_check_filter(scale, CorrelationDirection::Forward);
-        self.cross_check_filter(scale, CorrelationDirection::Reverse);
+        self.cross_check_filter(scale, CorrelationDirection::Forward)?;
+        self.cross_check_filter(scale, CorrelationDirection::Reverse)?;
 
         self.first_pass = false;
 
@@ -543,10 +544,14 @@ impl PointCorrelations {
             .floor() as usize
     }
 
-    fn cross_check_filter(&mut self, scale: f32, dir: CorrelationDirection) {
+    fn cross_check_filter(
+        &mut self,
+        scale: f32,
+        dir: CorrelationDirection,
+    ) -> Result<(), Box<dyn error::Error>> {
         if let Some(gpu_context) = &mut self.gpu_context {
-            gpu_context.cross_check_filter(scale, dir);
-            return;
+            gpu_context.cross_check_filter(scale, dir)?;
+            return Ok(());
         };
 
         let (correlated_points, correlated_points_reverse) = match dir {
@@ -572,6 +577,7 @@ impl PointCorrelations {
                     }
                 }
             });
+        Ok(())
     }
 
     #[inline]

src/correlation/cross_check_filter.comp.glsl

@@ -0,0 +1,41 @@
+#version 450
+#pragma shader_stage(compute)
+
+layout(std430, push_constant) uniform readonly Parameters
+{
+    uint img1_width;
+    uint img1_height;
+    uint img2_width;
+    uint img2_height;
+    uint out_width;
+    uint out_height;
+    float scale;
+    uint iteration_pass;
+    mat3x3 fundamental_matrix;
+    int corridor_offset;
+    uint corridor_start;
+    uint corridor_end;
+    uint kernel_size;
+    float threshold;
+    float min_stdev;
+    uint neighbor_distance;
+    float extend_range;
+    float min_range;
+};
+layout(std430, set = 1, binding = 0) buffer Img1
+{
+    ivec2 img1[];
+};
+layout(std430, set = 1, binding = 1) buffer readonly Img2
+{
+    ivec2 img2[];
+};
+layout (local_size_x = 16, local_size_y = 16, local_size_z = 1) in;
+
+void main() {
+    const uint x = gl_GlobalInvocationID.x;
+    const uint y = gl_GlobalInvocationID.y;
+
+    // TODO: remove this debug code
+    img1[0] = ivec2(img1_width, img1_height);
+}

src/correlation/init_out_data.comp.glsl

@@ -0,0 +1,72 @@
+#version 450
+#pragma shader_stage(compute)
+
+layout(std430, push_constant) uniform readonly Parameters
+{
+    uint img1_width;
+    uint img1_height;
+    uint img2_width;
+    uint img2_height;
+    uint out_width;
+    uint out_height;
+    float scale;
+    uint iteration_pass;
+    mat3x3 fundamental_matrix;
+    int corridor_offset;
+    uint corridor_start;
+    uint corridor_end;
+    uint kernel_size;
+    float threshold;
+    float min_stdev;
+    uint neighbor_distance;
+    float extend_range;
+    float min_range;
+};
+layout(std430, set = 0, binding = 0) buffer readonly Images
+{
+    // Layout:
+    // img1; img2
+    float images[];
+};
+layout(std430, set = 0, binding = 1) buffer Internals_Img1
+{
+    // Layout:
+    // For searchdata: contains [min_corridor, stdev] for image1
+    // For cross_correlate: contains [avg, stdev] for image1
+    float internals_img1[];
+};
+layout(std430, set = 0, binding = 2) buffer Internals_Img2
+{
+    // Layout:
+    // Contains [avg, stdev] for image 2
+    float internals_img2[];
+};
+layout(std430, set = 0, binding = 3) buffer Internals_Int
+{
+    // Layout:
+    // Contains [min, max, neighbor_count] for the corridor range
+    int internals_int[];
+};
+layout(std430, set = 0, binding = 4) buffer Result_Matches
+{
+    ivec2 result_matches[];
+};
+layout(std430, set = 0, binding = 5) buffer  Result_Corr
+{
+    float result_corr[];
+};
+layout (local_size_x = 16, local_size_y = 16, local_size_z = 1) in;
+
+void main() {
+    const uint x = gl_GlobalInvocationID.x;
+    const uint y = gl_GlobalInvocationID.y;
+    /*
+    if (x < out_width && y < out_height) {
+        result_matches[out_width*y+x] = ivec2(-1, -1);
+        result_corr[out_width*y+x] = -1.0;
+    }
+    */
+
+    // TODO: remove this debug code
+    result_corr[0] = threshold;
+}

src/correlation/prepare_initialdata_correlation.comp.glsl

@@ -0,0 +1,5 @@
+#version 450
+#pragma shader_stage(compute)
+
+void main() {
+}

src/correlation/prepare_initialdata_searchdata.comp.glsl

@@ -0,0 +1,5 @@
+#version 450
+#pragma shader_stage(compute)
+
+void main() {
+}

src/correlation/prepare_searchdata.comp.glsl

@@ -0,0 +1,5 @@
+#version 450
+#pragma shader_stage(compute)
+
+void main() {
+}