old dir

9 files changed, 669 insertions, 0 deletions

diff --git a/old/vgcodec/src/app.rs b/old/vgcodec/src/app.rs
new file mode 100644
index 0000000..b51284c
--- /dev/null
+++ b/old/vgcodec/src/app.rs
@@ -0,0 +1,59 @@
+use std::sync::Arc;
+
+use wgpu::{Adapter, Device, Extent3d, ImageCopyTexture, Instance, Origin3d, Queue, Texture};
+
+pub struct App {
+    pub instance: Instance,
+    pub device: Device,
+    pub adapter: Adapter,
+    pub queue: Queue,
+}
+
+impl App {
+    pub async fn new() -> Arc<Self> {
+        let instance = wgpu::Instance::new(wgpu::Backends::all());
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions::default())
+            .await
+            .unwrap();
+        let (device, queue) = adapter
+            .request_device(
+                &wgpu::DeviceDescriptor {
+                    label: None,
+                    features: wgpu::Features::empty(),
+                    limits: wgpu::Limits::downlevel_defaults(),
+                },
+                None,
+            )
+            .await
+            .unwrap();
+        Arc::new(Self {
+            adapter,
+            device,
+            instance,
+            queue,
+        })
+    }
+
+    pub fn copy_texture(&self, source: &Texture, destination: &Texture, size: Extent3d) {
+        let mut encoder = self
+            .device
+            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        encoder.copy_texture_to_texture(
+            ImageCopyTexture {
+                aspect: wgpu::TextureAspect::All,
+                mip_level: 0,
+                origin: Origin3d::ZERO,
+                texture: source,
+            },
+            ImageCopyTexture {
+                aspect: wgpu::TextureAspect::All,
+                mip_level: 0,
+                origin: Origin3d::ZERO,
+                texture: destination,
+            },
+            size,
+        );
+        self.queue.submit(Some(encoder.finish()));
+    }
+}
diff --git a/old/vgcodec/src/approximate.rs b/old/vgcodec/src/approximate.rs
new file mode 100644
index 0000000..5bff325
--- /dev/null
+++ b/old/vgcodec/src/approximate.rs
@@ -0,0 +1,175 @@
+use crate::{
+    app::App,
+    diff::Differ,
+    export::Exporter,
+    paint::{PaintUniforms, Painter},
+};
+use log::{debug, info, warn};
+use rand::Rng;
+use std::sync::Arc;
+use wgpu::{Extent3d, Texture, TextureUsages};
+
+#[allow(unused)]
+pub struct Approximator {
+    app: Arc<App>,
+    size: Extent3d,
+    differ: Differ,
+    painter: Painter,
+    exporter: Exporter,
+
+    tex_approx: Texture,
+    tex_savestate: Texture,
+    tex_target: Texture,
+}
+
+impl Approximator {
+    pub fn new(app: &Arc<App>, tex_target: Texture, size: Extent3d) -> Self {
+        let App { device, .. } = app.as_ref();
+        let tex_approx = device.create_texture(&wgpu::TextureDescriptor {
+            size,
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Rgba8Unorm,
+            usage: TextureUsages::COPY_DST
+                | TextureUsages::TEXTURE_BINDING
+                | TextureUsages::STORAGE_BINDING
+                | TextureUsages::COPY_SRC,
+            label: None,
+        });
+        let tex_savestate = device.create_texture(&wgpu::TextureDescriptor {
+            size,
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Rgba8Unorm,
+            usage: TextureUsages::COPY_DST | TextureUsages::COPY_SRC,
+            label: None,
+        });
+
+        // let initial_init = (0..)
+        // write_texture(queue, &tex_approx, initial_init, size);
+
+        let differ = Differ::new(&app, size, &tex_approx, &tex_target);
+        let painter = Painter::new(&app, size, &tex_approx);
+        let exporter = Exporter::new(&app, size);
+
+        Approximator {
+            app: app.clone(),
+            size,
+            differ,
+            tex_savestate,
+            tex_approx,
+            tex_target,
+            painter,
+            exporter,
+        }
+    }
+
+    pub fn save(&self) {
+        self.app
+            .copy_texture(&self.tex_approx, &self.tex_savestate, self.size);
+    }
+    pub fn restore(&self) {
+        self.app
+            .copy_texture(&self.tex_savestate, &self.tex_approx, self.size);
+    }
+
+    pub async fn run(&mut self, iters: usize, out: &str) -> Vec<PaintUniforms> {
+        let rscale = self.size.width as f32 * 0.1;
+        let mut objects = vec![];
+        for i in 0..iters {
+            let mut p = PaintUniforms {
+                x: rand::thread_rng().gen_range(0.0..self.size.width as f32),
+                y: rand::thread_rng().gen_range(0.0..self.size.height as f32),
+                rx: rscale,
+                ry: rscale,
+                r: rand::thread_rng().gen_range(0.0..1.0),
+                g: rand::thread_rng().gen_range(0.0..1.0),
+                b: rand::thread_rng().gen_range(0.0..1.0),
+            };
+
+            // Find initial diff
+            self.save();
+            self.painter.run(p);
+            let initial_diff = pollster::block_on(self.differ.run());
+            debug!("initial diff={initial_diff}");
+            self.restore();
+
+            let mut cd = initial_diff;
+            self.save();
+            let ok = loop {
+                let mut q = 0;
+                let mut k = |k: usize| {
+                    q += k;
+                    q == 0
+                };
+                if k(self.optimize_param(&mut cd, &mut p, "more red", |p| p.r += 0.1)) {
+                    k(self.optimize_param(&mut cd, &mut p, "less red", |p| p.r -= 0.1));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "more green", |p| p.g += 0.1)) {
+                    k(self.optimize_param(&mut cd, &mut p, "less green", |p| p.g -= 0.1));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "more blue", |p| p.b += 0.1)) {
+                    k(self.optimize_param(&mut cd, &mut p, "less blue", |p| p.b -= 0.1));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "increase rx", |p| p.rx *= 1.5)) {
+                    k(self.optimize_param(&mut cd, &mut p, "decrease rx", |p| p.rx /= 1.5));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "increase ry", |p| p.ry *= 1.5)) {
+                    k(self.optimize_param(&mut cd, &mut p, "decrease ry", |p| p.ry /= 1.5));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "move right", |p| p.x += 8.0)) {
+                    k(self.optimize_param(&mut cd, &mut p, "move left", |p| p.x -= 8.0));
+                }
+                if k(self.optimize_param(&mut cd, &mut p, "move down", |p| p.y += 8.0)) {
+                    k(self.optimize_param(&mut cd, &mut p, "move up", |p| p.y -= 8.0));
+                }
+                if p.rx < 5.0 || p.ry < 5.0 {
+                    break false;
+                }
+                if q == 0 {
+                    break true;
+                }
+            };
+            if ok {
+                self.painter.run(p);
+                info!("{i} (improvement={})", initial_diff - cd);
+                objects.push(p);
+            } else {
+                warn!("object aborted");
+            }
+        }
+        self.exporter.run(&self.tex_approx, out).await;
+        return objects;
+    }
+
+    pub fn optimize_param<F>(
+        &self,
+        current_diff: &mut u32,
+        params: &mut PaintUniforms,
+        label: &'static str,
+        f: F,
+    ) -> usize
+    where
+        F: Fn(&mut PaintUniforms) -> (),
+    {
+        let mut p = params.clone();
+        let mut i = 0;
+        loop {
+            f(&mut p);
+            self.painter.run(p);
+            let diff = pollster::block_on(self.differ.run());
+            // pollster::block_on(self.exporter.run(&self.tex_approx, &format!("a/snap.png")));
+            debug!("try {label:?} ({})", diff as i64 - *current_diff as i64);
+            self.restore();
+            if diff >= *current_diff {
+                break i;
+            }
+            debug!("applied {label:?}");
+            i += 1;
+            *current_diff = diff;
+            *params = p
+        }
+    }
+}
diff --git a/old/vgcodec/src/diff.rs b/old/vgcodec/src/diff.rs
new file mode 100644
index 0000000..4a7453c
--- /dev/null
+++ b/old/vgcodec/src/diff.rs
@@ -0,0 +1,122 @@
+use std::{borrow::Cow, sync::Arc};
+use wgpu::{util::DeviceExt, BindGroup, Buffer, ComputePipeline, Extent3d, Texture};
+
+use crate::app::App;
+
+pub struct Differ {
+    app: Arc<App>,
+    size: Extent3d,
+
+    pipeline: ComputePipeline,
+    bind_group: BindGroup,
+
+    counter_buffer_size: u64,
+    counter_buffer: Buffer,
+    counter_staging_buffer: Buffer,
+}
+
+impl Differ {
+    pub fn new(app: &Arc<App>, extent: Extent3d, tex_a: &Texture, tex_b: &Texture) -> Self {
+        let App { device, .. } = app.as_ref();
+        let cs_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("diff.wgsl"))),
+        });
+
+        let counter_buffer_size = std::mem::size_of::<u32>() as wgpu::BufferAddress;
+        let counter_staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: counter_buffer_size,
+            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+        let counter_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: None,
+            contents: bytemuck::cast_slice(&[0u32]),
+            usage: wgpu::BufferUsages::STORAGE
+                | wgpu::BufferUsages::COPY_DST
+                | wgpu::BufferUsages::COPY_SRC,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: None,
+            module: &cs_module,
+            entry_point: "main",
+        });
+
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &pipeline.get_bind_group_layout(0),
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(
+                        &tex_a.create_view(&wgpu::TextureViewDescriptor::default()),
+                    ),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::TextureView(
+                        &tex_b.create_view(&wgpu::TextureViewDescriptor::default()),
+                    ),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: counter_buffer.as_entire_binding(),
+                },
+            ],
+        });
+        Self {
+            app: app.clone(),
+            size: extent,
+            pipeline,
+            counter_buffer_size,
+            counter_buffer,
+            bind_group,
+            counter_staging_buffer,
+        }
+    }
+
+    pub async fn run(&self) -> u32 {
+        let App { device, queue, .. } = self.app.as_ref();
+
+        queue.write_buffer(&self.counter_buffer, 0, bytemuck::cast_slice(&[0u32]));
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        {
+            let mut cpass =
+                encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: None });
+            cpass.set_pipeline(&self.pipeline);
+            cpass.set_bind_group(0, &self.bind_group, &[]);
+            cpass.dispatch_workgroups(self.size.width, self.size.height, 1);
+        }
+
+        encoder.copy_buffer_to_buffer(
+            &self.counter_buffer,
+            0,
+            &self.counter_staging_buffer,
+            0,
+            self.counter_buffer_size,
+        );
+
+        queue.submit(Some(encoder.finish()));
+
+        let buffer_slice = self.counter_staging_buffer.slice(..);
+        let (sender, receiver) = futures_intrusive::channel::shared::oneshot_channel();
+        buffer_slice.map_async(wgpu::MapMode::Read, move |v| {
+            sender.send(v.unwrap()).unwrap()
+        });
+        device.poll(wgpu::Maintain::Wait);
+        receiver.receive().await;
+
+        let data = buffer_slice.get_mapped_range();
+        let result: u32 = bytemuck::cast_slice(&data).to_vec()[0];
+
+        drop(data);
+        self.counter_staging_buffer.unmap();
+
+        result
+    }
+}
diff --git a/old/vgcodec/src/diff.wgsl b/old/vgcodec/src/diff.wgsl
new file mode 100644
index 0000000..82e169b
--- /dev/null
+++ b/old/vgcodec/src/diff.wgsl
@@ -0,0 +1,16 @@
+@group(0) @binding(0) 
+var tex_a: texture_2d<f32>;
+@group(0) @binding(1) 
+var tex_b: texture_2d<f32>;
+
+@group(0) @binding(2)
+var<storage, read_write> exp: atomic<u32>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    var col_a = textureLoad(tex_a, vec2(i32(global_id.x), i32(global_id.y)), 0);
+    var col_b = textureLoad(tex_b, vec2(i32(global_id.x), i32(global_id.y)), 0);
+    var diff = length(col_a - col_b);
+    atomicAdd(&exp, u32(diff * 1000.0));
+}
+
diff --git a/old/vgcodec/src/export.rs b/old/vgcodec/src/export.rs
new file mode 100644
index 0000000..e020c5d
--- /dev/null
+++ b/old/vgcodec/src/export.rs
@@ -0,0 +1,83 @@
+use crate::app::App;
+use image::RgbaImage;
+use std::{num::NonZeroU32, sync::Arc};
+use wgpu::{
+    Buffer, Extent3d, ImageCopyBuffer, ImageCopyTexture, ImageDataLayout, Origin3d, Texture,
+};
+
+pub struct Exporter {
+    app: Arc<App>,
+    size: Extent3d,
+
+    padded_bytes_per_row: u32,
+    export_buffer: Buffer,
+}
+
+impl Exporter {
+    pub fn new(app: &Arc<App>, size: Extent3d) -> Self {
+        let App { device, .. } = app.as_ref();
+
+        let bytes_per_pixel = std::mem::size_of::<u32>() as u32;
+        let unpadded_bytes_per_row = size.width * bytes_per_pixel;
+        let align = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
+        let padded_bytes_per_row_padding = (align - unpadded_bytes_per_row % align) % align;
+        let padded_bytes_per_row = unpadded_bytes_per_row + padded_bytes_per_row_padding;
+
+        let export_buffer_size = (padded_bytes_per_row * size.height) as u64;
+        let export_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: export_buffer_size,
+            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        Self {
+            padded_bytes_per_row,
+            app: app.clone(),
+            size,
+            export_buffer,
+        }
+    }
+    pub async fn run(&self, texture: &Texture, save_path: &str) {
+        let App { device, queue, .. } = self.app.as_ref();
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+        encoder.copy_texture_to_buffer(
+            ImageCopyTexture {
+                texture,
+                aspect: wgpu::TextureAspect::All,
+                mip_level: 0,
+                origin: Origin3d::ZERO,
+            },
+            ImageCopyBuffer {
+                buffer: &self.export_buffer,
+                layout: ImageDataLayout {
+                    offset: 0,
+                    bytes_per_row: Some(NonZeroU32::new(self.padded_bytes_per_row).unwrap()),
+                    rows_per_image: None,
+                },
+            },
+            self.size,
+        );
+
+        queue.submit(Some(encoder.finish()));
+
+        let buffer_slice = self.export_buffer.slice(..);
+        let (sender, receiver) = futures_intrusive::channel::shared::oneshot_channel();
+        buffer_slice.map_async(wgpu::MapMode::Read, move |v| {
+            sender.send(v.unwrap()).unwrap()
+        });
+        device.poll(wgpu::Maintain::Wait);
+        receiver.receive().await;
+
+        let data = buffer_slice.get_mapped_range();
+        let result: Vec<u8> = bytemuck::cast_slice(&data).to_vec();
+        let image = RgbaImage::from_raw(self.size.width, self.size.height, result).unwrap();
+        image.save(save_path).unwrap();
+
+        drop(data);
+        self.export_buffer.unmap();
+    }
+}
diff --git a/old/vgcodec/src/helper.rs b/old/vgcodec/src/helper.rs
new file mode 100644
index 0000000..13f9c1a
--- /dev/null
+++ b/old/vgcodec/src/helper.rs
@@ -0,0 +1,40 @@
+use log::info;
+use wgpu::{Extent3d, Queue, Texture};
+
+pub fn write_texture(queue: &Queue, target: &Texture, data: &[u8], size: Extent3d) {
+    info!("uploading texture {size:?} ({} bytes)", data.len());
+
+    let bytes_per_pixel = std::mem::size_of::<u32>() as u32;
+    let unpadded_bytes_per_row = size.width * bytes_per_pixel;
+    let align = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
+    let padded_bytes_per_row_padding = (align - unpadded_bytes_per_row % align) % align;
+    let padded_bytes_per_row = unpadded_bytes_per_row + padded_bytes_per_row_padding;
+
+    let mut padded = vec![];
+    for y in 0..(size.height as usize) {
+        for x in 0..(size.width as usize) {
+            for c in 0..4 {
+                padded.push(data[c + x * 4 + y * 4 * size.width as usize])
+            }
+        }
+        for _ in 0..padded_bytes_per_row_padding {
+            padded.push(0)
+        }
+    }
+
+    queue.write_texture(
+        wgpu::ImageCopyTexture {
+            texture: &target,
+            mip_level: 0,
+            origin: wgpu::Origin3d::ZERO,
+            aspect: wgpu::TextureAspect::All,
+        },
+        &padded,
+        wgpu::ImageDataLayout {
+            offset: 0,
+            bytes_per_row: Some(std::num::NonZeroU32::try_from(padded_bytes_per_row).unwrap()),
+            rows_per_image: None,
+        },
+        size,
+    );
+}
diff --git a/old/vgcodec/src/main.rs b/old/vgcodec/src/main.rs
new file mode 100644
index 0000000..9cae234
--- /dev/null
+++ b/old/vgcodec/src/main.rs
@@ -0,0 +1,60 @@
+pub mod app;
+pub mod approximate;
+pub mod diff;
+pub mod export;
+pub mod helper;
+pub mod paint;
+
+use app::App;
+use approximate::Approximator;
+use clap::Parser;
+use helper::write_texture;
+use log::info;
+use wgpu::TextureUsages;
+
+fn main() {
+    env_logger::init_from_env("LOG");
+    pollster::block_on(run());
+}
+
+#[derive(Parser)]
+#[clap(about)]
+struct Args {
+    #[clap(short = 'I', long)]
+    iterations: usize,
+    #[clap(short = 'o', long)]
+    outfile: String,
+    infile: String,
+}
+
+async fn run() {
+    let args = Args::parse();
+    let app = app::App::new().await;
+
+    let App { device, queue, .. } = app.as_ref();
+
+    let img_target = image::open(&args.infile).unwrap().into_rgba8();
+
+    let size = wgpu::Extent3d {
+        width: img_target.width(),
+        height: img_target.height(),
+        depth_or_array_layers: 1,
+    };
+
+    let tex_target = device.create_texture(&wgpu::TextureDescriptor {
+        size,
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        usage: TextureUsages::COPY_DST | TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_SRC,
+        label: None,
+    });
+
+    let img_raw = img_target.into_raw();
+    info!("{}", img_raw.len());
+    write_texture(queue, &tex_target, &img_raw, size);
+
+    let mut a = Approximator::new(&app, tex_target, size);
+    a.run(args.iterations, &args.outfile).await;
+}
diff --git a/old/vgcodec/src/paint.rs b/old/vgcodec/src/paint.rs
new file mode 100644
index 0000000..33ab572
--- /dev/null
+++ b/old/vgcodec/src/paint.rs
@@ -0,0 +1,92 @@
+use bytemuck::{Pod, Zeroable};
+use std::{borrow::Cow, sync::Arc};
+use wgpu::{util::DeviceExt, BindGroup, Buffer, ComputePipeline, Extent3d, Texture};
+
+use crate::app::App;
+
+pub struct Painter {
+    app: Arc<App>,
+    size: Extent3d,
+
+    pipeline: ComputePipeline,
+    bind_group: BindGroup,
+
+    uniform_buffer: Buffer,
+}
+
+#[repr(C)]
+#[derive(Pod, Zeroable, Clone, Copy)]
+pub struct PaintUniforms {
+    pub x: f32,
+    pub y: f32,
+    pub rx: f32,
+    pub ry: f32,
+    pub r: f32,
+    pub g: f32,
+    pub b: f32,
+}
+
+impl Painter {
+    pub fn new(app: &Arc<App>, extent: Extent3d, texture: &Texture) -> Self {
+        let App { device, .. } = app.as_ref();
+        let cs_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(include_str!("paint.wgsl"))),
+        });
+
+        // let uniform_buffer_size = std::mem::size_of::<PaintUniforms>() as wgpu::BufferAddress;
+        let uniform_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: None,
+            contents: bytemuck::cast_slice(&[PaintUniforms::zeroed()]),
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: None,
+            module: &cs_module,
+            entry_point: "main",
+        });
+
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &pipeline.get_bind_group_layout(0),
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(
+                        &texture.create_view(&wgpu::TextureViewDescriptor::default()),
+                    ),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: uniform_buffer.as_entire_binding(),
+                },
+            ],
+        });
+        Self {
+            app: app.clone(),
+            size: extent,
+            pipeline,
+            uniform_buffer,
+            bind_group,
+        }
+    }
+
+    pub fn run(&self, params: PaintUniforms) {
+        let App { device, queue, .. } = self.app.as_ref();
+
+        queue.write_buffer(&self.uniform_buffer, 0, bytemuck::cast_slice(&[params]));
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        {
+            let mut cpass =
+                encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: None });
+            cpass.set_pipeline(&self.pipeline);
+            cpass.set_bind_group(0, &self.bind_group, &[]);
+            cpass.dispatch_workgroups(self.size.width, self.size.height, 1);
+        }
+        queue.submit(Some(encoder.finish()));
+    }
+}
diff --git a/old/vgcodec/src/paint.wgsl b/old/vgcodec/src/paint.wgsl
new file mode 100644
index 0000000..7006310
--- /dev/null
+++ b/old/vgcodec/src/paint.wgsl
@@ -0,0 +1,22 @@
+struct Uniforms {
+    x: f32,
+    y: f32,
+    rx: f32,
+    ry: f32,
+    r: f32,
+    g: f32,
+    b: f32
+};
+
+@group(0) @binding(0) var tex: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(1) var<uniform> uniforms: Uniforms;
+
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let coords = global_id.xy;
+    let kc = vec2<f32>(coords) - vec2(uniforms.x, uniforms.y);
+    let d = length(kc / vec2(uniforms.rx, uniforms.ry));
+    if d < 1.0 {
+        textureStore(tex, vec2<i32>(coords), vec4<f32>(vec3<f32>(uniforms.r, uniforms.g, uniforms.b), 1.0));
+    }
+}