From 66d11e22c06c1b30e118c756b8be56c7557829ea Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Wed, 28 Aug 2019 22:55:44 -0700 Subject: [PATCH] just doing a naive layered compute for meow --- src/canvas.rs | 4 +++ src/compu_wip.rs | 82 +++++++++++++++++++++++++++++++++++++--------- src/vkprocessor.rs | 7 ++-- 3 files changed, 75 insertions(+), 18 deletions(-) diff --git a/src/canvas.rs b/src/canvas.rs index a6fdd4b6..c9374073 100644 --- a/src/canvas.rs +++ b/src/canvas.rs @@ -261,6 +261,10 @@ impl Canvas { id } + pub fn get_image(&self, image_handle: Arc) -> std::sync::Arc { + 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> { let project_root = diff --git a/src/compu_wip.rs b/src/compu_wip.rs index 05b09fb2..08cded7c 100644 --- a/src/compu_wip.rs +++ b/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, data: Vec, 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>>) -> Arc>>, ((((), - PersistentDescriptorSetBuf>>), - PersistentDescriptorSetBuf>>), - PersistentDescriptorSetBuf>>)>> { + PersistentDescriptorSetBuf>>), + PersistentDescriptorSetBuf>>), + PersistentDescriptorSetBuf>>)>> { 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, Vec>{ - + pub fn read_output_buffer(&self) -> ImageBuffer, Vec> { 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) -> Arc { - 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, data: Vec) { - - } + pub fn write_compute_buffer(&self, handle: Arc, data: Vec) {} 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![], } } diff --git a/src/vkprocessor.rs b/src/vkprocessor.rs index 6f8142bc..03477371 100644 --- a/src/vkprocessor.rs +++ b/src/vkprocessor.rs @@ -184,11 +184,9 @@ impl<'a> VkProcessor<'a> { -> Box { // 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();