From 9b8d5cd82841ef8edd346eab540743299ee4a670 Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Tue, 27 Aug 2019 21:48:21 -0700 Subject: [PATCH] have the api of the compu somewhat figured out. Need to weigh my options on texture, buffer, and image handles and how I'm going to get them from the vkprocessor --- src/compu_wip.rs | 184 +++++++++++++++++++++++++++++++++++++ src/main.rs | 120 +++++++++++------------- src/sprite.rs | 2 - src/util/shader_kernels.rs | 28 ------ 4 files changed, 236 insertions(+), 98 deletions(-) create mode 100644 src/compu_wip.rs diff --git a/src/compu_wip.rs b/src/compu_wip.rs new file mode 100644 index 00000000..e4908260 --- /dev/null +++ b/src/compu_wip.rs @@ -0,0 +1,184 @@ +use vulkano::buffer::{CpuAccessibleBuffer, BufferUsage}; +use std::sync::Arc; +use crate::canvas::Drawable; +use vulkano::framebuffer::RenderPassAbstract; +use vulkano::pipeline::{GraphicsPipelineAbstract, ComputePipeline}; +use vulkano::device::Device; +use crate::util::compute_kernel::ComputeKernel; +use image::{ImageBuffer}; +use image::GenericImageView; +use crate::util::compute_image::ComputeImage; +use vulkano::image::{ImageUsage, AttachmentImage}; +use vulkano::descriptor::descriptor_set::{PersistentDescriptorSetBuf, PersistentDescriptorSet}; +use vulkano::format::Format; +use vulkano::descriptor::pipeline_layout::PipelineLayout; + + +pub struct CompuSprite { + vertices: [(f32, f32); 6], + position: (f32, f32), + size: (f32, f32), + color: (f32, f32, f32, f32), + image_handle: Arc, +} + +impl CompuSprite { + + pub fn new(position: (f32, f32), + size: (f32, f32), + image_handle: Arc) -> CompuSprite { + + let fsize = (size.0 as f32, size.1 as f32); + + CompuSprite { + vertices: [ + (position.0, position.1 ), // top left + (position.0, position.1 + fsize.1), // bottom left + (position.0 + fsize.0, position.1 + fsize.1 ), // bottom right + (position.0, position.1 ), // top left + (position.0 + fsize.0, position.1 + fsize.1 ), // bottom right + (position.0 + fsize.0, position.1 ), // top right + ], + + position: position, + size: size, + color: (0.0,0.0,0.0,0.0), + image_handle: image_handle.clone() + } + } + + fn get_vertices(&self) -> Vec<(f32, f32)> { + self.vertices.to_vec() + } + + fn get_color(&self) -> (f32, f32, f32, f32) { + self.color.clone() + } + + fn get_image_handle(&self) -> Arc { + self.image_handle.clone() + } +} + + +pub struct CompuBuffers { + dimensions: (u32, u32), + device: Arc, + + io_buffers: Vec>>, + settings_buffer: Arc>, +} + +impl CompuBuffers { + + pub fn new(device: Arc, data: Vec, dimensions: (u32, u32), stride: u32) -> CompuBuffers { + + let data_length = dimensions.0 * dimensions.1 * stride; + + let input_buffer = { + let mut buff = data.iter(); + let data_iter = (0..data_length).map(|n| *(buff.next().unwrap())); + CpuAccessibleBuffer::from_iter(device.clone(), BufferUsage::all(), data_iter).unwrap() + }; + + let output_buffer = { + let mut buff = data.iter(); + let data_iter = (0..data_length).map(|n| *(buff.next().unwrap())); + CpuAccessibleBuffer::from_iter(device.clone(), BufferUsage::all(), data_iter).unwrap() + }; + + // Settings buffer which holds i32's + // Compile macros into the kernel eventually to index them + let settings_buffer = { + let vec = vec![dimensions.0, dimensions.1]; + let mut buff = vec.iter(); + let data_iter = + (0..2).map(|n| *(buff.next().unwrap())); + CpuAccessibleBuffer::from_iter(device.clone(), + BufferUsage::all(), + data_iter).unwrap() + }; + + CompuBuffers{ + dimensions: dimensions, + device: device.clone(), + + io_buffers: vec![input_buffer, output_buffer], + settings_buffer: settings_buffer + } + } + + + pub fn get_size(&self) -> (u32, u32) { + self.dimensions + } + + + pub fn get_descriptor_set(&self, compute_pipeline: std::sync::Arc>>) + -> Arc>>, ((((), + 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() + .add_buffer(self.settings_buffer.clone()).unwrap() + .build().unwrap()) + } +} + +// Canvas analog +pub struct CompuState { + compute_buffers: Vec, + compute_buffer_handles: Vec>, + + kernels: Vec, + kernel_handles: Vec>, +} + +impl CompuState { + fn new_compute_buffer() -> Arc { + + } + + fn new_kernel(&mut self, + filename: String, + device: &Arc) -> Arc { + + let kernel = ComputeKernel::new(filename, device.clone()); + + self.kernels.push(kernel); + + let id = Arc::new(self.kernels.len() as u32); + + self.kernel_handles.push(id.clone()); + + id + } +} + + +/* + ComputeFrame accumulates combos of : + Buffers and Kernels + Buffers, Images and Kernels + Buffers, Buffers and Kernels + + It will need to convert these into a logical list of command_buffer commands +*/ +pub struct ComputeFrame { + +} + +impl ComputeFrame { + fn add() { + + } + fn add_chained(output_buffer_id: Arc, input_buffer_id: Arc, kernel_id: Arc ) { + + } + fn add_with_image_swap(buffer_id: Arc, sprite: CompuSprite, ) { + + } +} diff --git a/src/main.rs b/src/main.rs index d366b817..ccec939a 100644 --- a/src/main.rs +++ b/src/main.rs @@ -20,7 +20,7 @@ use winit::dpi::LogicalSize; use vulkano_win::VkSurfaceBuild; use sprite::Sprite; use crate::canvas::CanvasFrame; - +use crate::compu_wip::{CompuSprite, ComputeFrame, CompuBuffers, CompuState}; mod util; mod slider; @@ -32,7 +32,7 @@ mod vertex_2d; mod vertex_3d; mod sprite; mod canvas; - +mod compu_wip; /* @@ -123,77 +123,54 @@ So to start with we need to know the compute data which we want to run and the k the descriptor set. We also need the XY for the compute run -I think I need to keep these kernels in their own Canvas type of structure - I want to modify their constants and recompile - - -let comp = Computable::new(device, data, kernel_id) - -I have canvas frame. Perhaps I should do something similar for compute - -ComputeFrame.add(device, data, kernel_id) - - - - -Not sure if I can do this... but inside it could be like -CompuSprite::new(&mut canvas, kernel); - -new(&mut canvas) -> CompuSprite { - let img_id = canvas.new_image(); - - Canvas { - img_id.clone() - } -} - - - - - -I want to be able to chain computes using the same data -So that would be a different pipeline using the same or similar descriptor set - - - - - -sprite = Sprite::with_texture(Canvas.get_texture_from_file()) - -(compute, sprite2) = Compute::with_swap_image(Canvas.get_new_image()) - - - -compute load shader -> shader object - -compute load buffers -> buffer object - -shader object + buffer object + maybe the swap buffer - -> command queue - - - - - - - -let mut canvas = CanvasFrame::new(); -canvas.draw(&sprite); -canvas.draw(&sprite2); - +CompuSprite + holds the image handle + impl's drawable +CompuBuffers + holds the buffer handles + creates the descriptor set? +CompuState + holds the compute buffers + holds the compute kernels +ComputeFrame + preserves order of entered elements + holds compute buffer ids and their tied kernel ids + holds compute buffer ids, their swap image ids and their kernel ids + holds compute buffer ids, their swap buffer ids, and their kernel ids -(frame_future) = processor.run(&surface, frame_future, canvas); +// Creates the compu sprite with an associated image stored in canvas +let compu_sprite = CompuSprite::new(&mut canvas) +// Compiles and stores the kernel in CompuState +// Returns the handle to the stored kernel (immutable at this state probably) +let KernelID = ComputeState.new_kernel(kernel_path) +// Pushes the data to the GPU and returns a handle +let ComputeBuffer_1 = ComputeState.new_compute_buffer(input_data, input_settings) +// One by one, assign the buffers to their kernels for run +ComputeFrame.add(ComputeBuffer_1, KernelID) +// If the buffer needs to be swapped in first +ComputeFrame.add_with_buffer_swap(ComputeBuffer_1, ComputeBuffer_2, KernelID) +// If the buffer should be swappd out to an image after it computes +ComputeFrame.add_with_image_swap(ComputeBuffer_1, compu_sprite, KernelID) +CompuState.compute() +// Full example if API: +CompuFrame.add(ComputeBuffer1ID, Kernel1ID) +CompuFrame.add(ComputeBuffer2ID, Kernel2ID) +CompuFrame.add_with_image_swap(ComputeBuffer1ID, CompuSprite, KernelID) +CanvasFrame.draw(Sprite1) +CanvasFrame.draw(Sprite2) +CanvasFrame.draw(CompuSprite) +vkprocessor.run(CanvasFrame, CompuFrame) @@ -201,7 +178,6 @@ canvas.draw(&sprite2); */ - fn main() { let instance = { @@ -232,9 +208,23 @@ fn main() { let mut mouse_xy = Vector2i::new(0,0); + + + let sprite = Sprite::new_with_color((0.,0.), (0.1,0.1), (1.,0.,0.,1.)); let sprite2 = Sprite::new_with_color((-1.,-0.5), (0.1,0.1), (0.,1.,0.,1.)); + /* + I need to get the image id from the CompuState + But the CompuState is owned by the vkprocessor I think... + + I will also need to get the texture id + */ + + let compu_sprite1 = CompuSprite::new((-1.,-0.5), (0.1,0.1)); + + + while let Some(p) = window.get_position() { elapsed_time = timer.elap_time(); @@ -283,12 +273,6 @@ fn main() { return; } - /* - - - - - */ let mut canvas = CanvasFrame::new(); canvas.draw(&sprite); diff --git a/src/sprite.rs b/src/sprite.rs index c9880bae..9bb7468c 100644 --- a/src/sprite.rs +++ b/src/sprite.rs @@ -11,8 +11,6 @@ pub struct Sprite { textured: bool, texture_id: Option, - - } impl Sprite { diff --git a/src/util/shader_kernels.rs b/src/util/shader_kernels.rs index e1a11c80..1c955691 100644 --- a/src/util/shader_kernels.rs +++ b/src/util/shader_kernels.rs @@ -1,41 +1,13 @@ - - - use vulkano::device::{Device, Queue}; use vulkano::instance::{PhysicalDevice, QueueFamily}; use vulkano::pipeline::{GraphicsPipeline, GraphicsPipelineAbstract, GraphicsPipelineBuilder}; - - - use std::sync::Arc; use std::ffi::CStr; use std::path::PathBuf; use shade_runner as sr; - - - - - - - use vulkano::framebuffer::{Subpass, RenderPassAbstract, Framebuffer, FramebufferAbstract}; use vulkano::pipeline::shader::{GraphicsShaderType, ShaderModule, SpecializationConstants, SpecializationMapEntry}; use vulkano::swapchain::{Capabilities}; - - - - - - - - - - - - - - - use crate::vertex_2d::ColoredVertex2D; /*