just doing a naive layered compute for meow

src/canvas.rs

@@ -261,6 +261,10 @@ impl Canvas {
         id
     }
 
+    pub fn get_image(&self, image_handle: Arc<u32>) -> std::sync::Arc<vulkano::image::attachment::AttachmentImage> {
+        self.image_buffers.get((*image_handle).clone() as usize).unwrap().clone()
+    }
+
     // TODO Handle file not found gracefully
     fn get_texture_from_file(&self, image_filename: String) -> Arc<ImmutableImage<Format>> {
         let project_root =

src/compu_wip.rs

@@ -1,6 +1,6 @@
 use vulkano::buffer::{CpuAccessibleBuffer, BufferUsage};
 use std::sync::Arc;
-use crate::canvas::Drawable;
+use crate::canvas::{Drawable, Canvas};
 use vulkano::framebuffer::RenderPassAbstract;
 use vulkano::pipeline::{GraphicsPipelineAbstract, ComputePipeline};
 use vulkano::device::Device;
@@ -14,6 +14,7 @@ use vulkano::format::Format;
 use vulkano::descriptor::pipeline_layout::PipelineLayout;
 use std::borrow::Borrow;
 use image::Rgba;
+use vulkano::command_buffer::AutoCommandBufferBuilder;
 
 pub struct CompuSprite {
     vertices: [(f32, f32); 6],
@@ -87,7 +88,6 @@ pub struct CompuBuffers {
 }
 
 impl CompuBuffers {
-
     pub fn new(device: Arc<Device>, data: Vec<u8>, dimensions: (u32, u32), stride: u32) -> CompuBuffers {
         let data_length = dimensions.0 * dimensions.1 * stride;
 
@@ -130,9 +130,9 @@ impl CompuBuffers {
 
     pub fn get_descriptor_set(&self, compute_pipeline: std::sync::Arc<ComputePipeline<PipelineLayout<shade_runner::layouts::ComputeLayout>>>)
                               -> Arc<PersistentDescriptorSet<std::sync::Arc<ComputePipeline<PipelineLayout<shade_runner::layouts::ComputeLayout>>>, ((((),
-                           PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u8]>>>),
+                                                                                                                                                       PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u8]>>>),
-                          PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u8]>>>),
+                                                                                                                                                      PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u8]>>>),
-                         PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u32]>>>)>> {
+                                                                                                                                                     PersistentDescriptorSetBuf<std::sync::Arc<vulkano::buffer::cpu_access::CpuAccessibleBuffer<[u32]>>>)>> {
         Arc::new(PersistentDescriptorSet::start(compute_pipeline.clone(), 0)
             .add_buffer(self.io_buffers.get(0).unwrap().clone()).unwrap()
             .add_buffer(self.io_buffers.get(1).unwrap().clone()).unwrap()
@@ -140,8 +140,7 @@ impl CompuBuffers {
             .build().unwrap())
     }
 
-    pub fn read_output_buffer(&self) -> ImageBuffer<Rgba<u8>, Vec<u8>>{
+    pub fn read_output_buffer(&self) -> ImageBuffer<Rgba<u8>, Vec<u8>> {
-
         let xy = self.get_size();
 
         self.io_buffers.get(1).unwrap().write().unwrap().map(|x| x);
@@ -167,7 +166,6 @@ pub struct CompuState {
 }
 
 impl CompuState {
-
     pub fn new() -> CompuState {
         CompuState {
             compute_buffers: vec![],
@@ -182,7 +180,6 @@ impl CompuState {
                               dimensions: (u32, u32),
                               stride: u32,
                               device: Arc<Device>) -> Arc<u32> {
-
         self.compute_buffers.push(
             CompuBuffers::new(device.clone(), data, dimensions, stride));
 
@@ -197,9 +194,7 @@ impl CompuState {
         //compute_buffer.read_output_buffer().to_vec()
         Vec::new()
     }
-    pub fn write_compute_buffer(&self, handle: Arc<u32>, data: Vec<u8>) {
+    pub fn write_compute_buffer(&self, handle: Arc<u32>, data: Vec<u8>) {}
-        
-    }
 
     pub fn new_kernel(&mut self,
                       filename: String,
@@ -215,13 +210,70 @@ impl CompuState {
         id
     }
 
-    pub fn compute(&mut self, compute_frame: ComputeFrame) {
+    pub fn compute_commands(&mut self,
+                            compute_frame: ComputeFrame,
+                            mut command_buffer: AutoCommandBufferBuilder,
+                            canvas: &Canvas)
+                            -> AutoCommandBufferBuilder {
+/*
+        let mut command_buffer,,  = command_buffer.dispatch([100, 100, 1],
+                  self.compute_kernel.clone().unwrap().clone().get_pipeline(),
+                  self.compute_image.clone().unwrap().clone()
+                      .get_descriptor_set(self.compute_kernel.clone().unwrap().clone().get_pipeline()).clone(), ()).unwrap()
+
+            .copy_buffer_to_image(self.compute_image.clone().unwrap().clone().rw_buffers.get(0).unwrap().clone(),
+                                  self.compute_image.clone().unwrap().clone().get_swap_buffer().clone()).unwrap();
+*/
 
         // i = (Buffer, Kernel)
         for i in compute_frame.pure_compute {
+            let buffer_id = (*i.0).clone() as usize;
+            let kernel_id = (*i.1).clone() as usize;
+
+            let buffer = self.compute_buffers.get(buffer_id).unwrap();
+            let kernel = self.kernels.get(buffer_id).unwrap();
+
+            let p = kernel.clone().get_pipeline();
+            let d = buffer.get_descriptor_set(kernel.clone().get_pipeline());
+
+            command_buffer = command_buffer
+                .dispatch([100,100,1], p, d, ()).unwrap()
+        }
+
+        // i = (Buffer, Image, Kernel)
+        for i in compute_frame.swapped_to_image {
+            let buffer_id = (*i.0).clone() as usize;
+            let image_id  = i.1.clone();
+            let kernel_id = (*i.2).clone() as usize;
+
+            let buffer = self.compute_buffers.get(buffer_id).unwrap();
+            let image = canvas.get_image(image_id);
+            let kernel = self.kernels.get(buffer_id).unwrap();
+
+            let p = kernel.clone().get_pipeline();
+            let d = buffer.get_descriptor_set(kernel.clone().get_pipeline());
+
+            command_buffer = command_buffer
+                .dispatch([100,100,1], p, d, ()).unwrap()
+                .copy_buffer_to_image(buffer.io_buffers.get(0).unwrap().clone(), image).unwrap();
+        }
+
+
+        // i = (Input Buffer, Output Buffer, Kernel)
+        for i in compute_frame.swapped_to_buffer {
+            let buffer_id = (*i.0).clone() as usize;
+            let kernel_id = (*i.1).clone() as usize;
+
+            let buffer = self.compute_buffers.get(buffer_id).unwrap();
+            let kernel = self.kernels.get(buffer_id).unwrap();
+
+            let p = kernel.clone().get_pipeline();
+            let d = buffer.get_descriptor_set(kernel.clone().get_pipeline());
 
+            command_buffer = command_buffer.dispatch([100,100,1], p, d, ()).unwrap()
         }
 
+        command_buffer
     }
 }
 
@@ -242,7 +294,7 @@ impl ComputeFrame {
         ComputeFrame {
             pure_compute: vec![],
             swapped_to_image: vec![],
-            swapped_to_buffer: vec![]
+            swapped_to_buffer: vec![],
         }
     }
 

src/vkprocessor.rs

@@ -184,11 +184,9 @@ impl<'a> VkProcessor<'a> {
         -> Box<dyn GpuFuture> {
 
         // take the canvas frame and create the vertex buffers
+        // TODO: This performs gpu buffer creation. Shouldn't be in hotpath
         self.canvas.draw(canvas_frame);
 
-        // ditto with the compuframe
-        self.compute_state.compute(compute_frame);
-
         let mut framebuffers =
             self.canvas.window_size_dependent_setup(&self.swapchain_images.clone().unwrap().clone());
 
@@ -229,6 +227,9 @@ impl<'a> VkProcessor<'a> {
                 .copy_buffer_to_image(self.compute_image.clone().unwrap().clone().rw_buffers.get(0).unwrap().clone(),
                                       self.compute_image.clone().unwrap().clone().get_swap_buffer().clone()).unwrap();
 
+        // ditto with the compuframe
+        let mut command_buffer = self.compute_state.compute_commands(compute_frame, command_buffer, &self.canvas);
+
         let mut command_buffer = self.canvas.draw_commands(command_buffer, framebuffers, image_num);
 
         let command_buffer = command_buffer.build().unwrap();