KIND OF ARBITRARY ZOOM (hangs at len 5 for some reason)
This commit is contained in:
75
src/client/render/compute/data.rs
Normal file
75
src/client/render/compute/data.rs
Normal file
@@ -0,0 +1,75 @@
|
||||
use nalgebra::Vector2;
|
||||
|
||||
use crate::util::FixedDec;
|
||||
|
||||
use super::Camera;
|
||||
|
||||
const VIEW_ALIGN: usize = 4 * 2;
|
||||
|
||||
pub struct ComputeView {
|
||||
pub bytes: Vec<u8>,
|
||||
}
|
||||
|
||||
impl ComputeView {
|
||||
pub fn bytes(&self) -> &[u8] {
|
||||
&self.bytes
|
||||
}
|
||||
}
|
||||
|
||||
impl Default for ComputeView {
|
||||
fn default() -> Self {
|
||||
let val = FixedDec::from_parts(false, 0, vec![0, 0, 0]);
|
||||
Self::new(true, Vector2::zeros(), 0, &val, &val, &val)
|
||||
}
|
||||
}
|
||||
|
||||
impl ComputeView {
|
||||
fn new(
|
||||
reset: bool,
|
||||
dims: Vector2<u32>,
|
||||
level: i32,
|
||||
scale: &FixedDec,
|
||||
x: &FixedDec,
|
||||
y: &FixedDec,
|
||||
) -> Self {
|
||||
let mut bytes = Vec::new();
|
||||
bytes.extend((reset as u32).to_le_bytes());
|
||||
bytes.extend(level.to_le_bytes());
|
||||
bytes.extend(bytemuck::cast_slice(&[dims.x, dims.y]));
|
||||
scale.to_bytes(&mut bytes);
|
||||
x.to_bytes(&mut bytes);
|
||||
y.to_bytes(&mut bytes);
|
||||
let rem = bytes.len() % VIEW_ALIGN;
|
||||
if rem != 0 {
|
||||
bytes.extend((0..(VIEW_ALIGN - rem)).map(|_| 0));
|
||||
}
|
||||
Self { bytes }
|
||||
}
|
||||
|
||||
pub fn from_camera_size(camera: &Camera, size: &Vector2<u32>, reset: bool, len: usize) -> Self {
|
||||
let mut x = camera.pos.x.clone();
|
||||
x.set_whole_len(1);
|
||||
x.set_dec_len(len as i32 - 1);
|
||||
let mut y = camera.pos.y.clone();
|
||||
y.set_whole_len(1);
|
||||
y.set_dec_len(len as i32 - 1);
|
||||
|
||||
let fsize: Vector2<f32> = size.cast();
|
||||
let stretch = if size.x < size.y {
|
||||
Vector2::new(fsize.x / fsize.y, 1.0)
|
||||
} else {
|
||||
Vector2::new(1.0, fsize.y / fsize.x)
|
||||
};
|
||||
|
||||
let mut scale = camera.zoom.mult().clone();
|
||||
scale.set_precision(len);
|
||||
|
||||
Self::new(reset, *size, camera.zoom.level(), &scale, &x, &y)
|
||||
}
|
||||
}
|
||||
|
||||
impl PartialEq for ComputeView {
|
||||
fn eq(&self, other: &Self) -> bool {
|
||||
self.bytes[1..] == other.bytes[1..]
|
||||
}
|
||||
}
|
||||
283
src/client/render/compute/fixed.wgsl
Normal file
283
src/client/render/compute/fixed.wgsl
Normal file
@@ -0,0 +1,283 @@
|
||||
const POS: u32 = 0u;
|
||||
const NEG: u32 = 1u;
|
||||
|
||||
struct FixedDec {
|
||||
sign: u32,
|
||||
dec: i32,
|
||||
parts: array<u32, LEN>,
|
||||
}
|
||||
|
||||
fn shr(lhs: FixedDec, rhs: i32) -> FixedDec {
|
||||
var parts = array<u32, LEN>();
|
||||
let sr = u32(rem_euclid(rhs, 32));
|
||||
let sl = 32 - sr;
|
||||
let mask = (1u << sr) - 1;
|
||||
var rem = 0u;
|
||||
for (var i = 0; i < ILEN; i += 1) {
|
||||
let part = lhs.parts[i];
|
||||
parts[i] = (part >> sr) ^ rem;
|
||||
rem = (part & mask) << sl;
|
||||
}
|
||||
return FixedDec(lhs.sign, lhs.dec, parts);
|
||||
}
|
||||
|
||||
fn add(lhs: FixedDec, rhs: FixedDec) -> FixedDec {
|
||||
var dest = FixedDec();
|
||||
dest.sign = lhs.sign;
|
||||
dest.dec = max(lhs.dec, rhs.dec);
|
||||
var carry = false;
|
||||
let rhs_offset = rhs.dec - dest.dec;
|
||||
let lhs_offset = lhs.dec - dest.dec;
|
||||
var i = ILEN;
|
||||
if lhs.sign == rhs.sign {
|
||||
while i > 0 {
|
||||
i -= 1;
|
||||
let a = at(lhs, i + lhs_offset);
|
||||
let b = at(rhs, i + rhs_offset);
|
||||
let res = a + b + u32(carry);
|
||||
dest.parts[i] = res;
|
||||
carry = res < a;
|
||||
}
|
||||
if carry {
|
||||
var i = ILEN - 1;
|
||||
while i > 0 {
|
||||
i -= 1;
|
||||
dest.parts[i + 1] = dest.parts[i];
|
||||
}
|
||||
dest.parts[0] = 1u;
|
||||
dest.dec += 1;
|
||||
}
|
||||
} else {
|
||||
while i > 0 {
|
||||
i -= 1;
|
||||
let a = at(lhs, i + lhs_offset);
|
||||
let b = at(rhs, i + rhs_offset);
|
||||
let res = a - b - u32(carry);
|
||||
dest.parts[i] = res;
|
||||
carry = a < res;
|
||||
}
|
||||
if carry {
|
||||
dest.sign = u32(dest.sign == 0);
|
||||
var i = 0;
|
||||
while i < ILEN {
|
||||
dest.parts[i] = ~dest.parts[i];
|
||||
i += 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
return dest;
|
||||
}
|
||||
|
||||
fn sub(lhs: FixedDec, rhs: FixedDec) -> FixedDec {
|
||||
var r = rhs;
|
||||
r.sign = u32(r.sign == 0);
|
||||
return add(lhs, r);
|
||||
}
|
||||
|
||||
fn at(dec: FixedDec, i: i32) -> u32 {
|
||||
if i < 0 || i >= ILEN {
|
||||
return 0u;
|
||||
}
|
||||
var parts = dec.parts;
|
||||
return parts[i];
|
||||
}
|
||||
|
||||
fn mul(lhs: FixedDec, rhs: FixedDec) -> FixedDec {
|
||||
let sign = u32(lhs.sign != rhs.sign);
|
||||
var parts = array<u32, LEN2>();
|
||||
var dec = lhs.dec + rhs.dec;
|
||||
var lparts = lhs.parts;
|
||||
var rparts = rhs.parts;
|
||||
|
||||
var i = LEN;
|
||||
while i > 0 {
|
||||
i -= 1u;
|
||||
let x = lparts[i];
|
||||
var carry: u32 = 0;
|
||||
var j = LEN;
|
||||
while j > 0 {
|
||||
j -= 1u;
|
||||
let y = rparts[j];
|
||||
|
||||
// widening mul
|
||||
let lsb = x * y;
|
||||
let a = x & 0xffff;
|
||||
let b = x >> 16;
|
||||
let c = y & 0xffff;
|
||||
let d = y >> 16;
|
||||
let ad = a * d + ((a * c) >> 16);
|
||||
let bc = b * c;
|
||||
let car = ad > (0xffffffff - bc);
|
||||
let msb = ((ad + bc) >> 16) + (u32(car) << 16) + b * d;
|
||||
|
||||
let k = i + j + 1;
|
||||
let res = parts[k] + lsb;
|
||||
let carry1 = res < lsb;
|
||||
let res2 = res + carry;
|
||||
let carry2 = res2 < res;
|
||||
parts[k] = res2;
|
||||
carry = u32(carry1) + u32(carry2) + msb;
|
||||
}
|
||||
parts[i] = carry;
|
||||
}
|
||||
|
||||
var new_parts = array<u32, LEN>();
|
||||
i = 0u;
|
||||
while i < LEN2 && parts[i] == 0 {
|
||||
dec -= 1;
|
||||
i += 1u;
|
||||
}
|
||||
var j = 0u;
|
||||
while j < LEN && (i + j) < LEN2 {
|
||||
new_parts[j] = parts[i + j];
|
||||
j += 1u;
|
||||
}
|
||||
return FixedDec(sign, dec, new_parts);
|
||||
}
|
||||
|
||||
fn gt(x: FixedDec, y: FixedDec) -> bool {
|
||||
if x.dec > y.dec {
|
||||
return true;
|
||||
}
|
||||
if y.dec > x.dec {
|
||||
return false;
|
||||
}
|
||||
return x.parts[0] > y.parts[0];
|
||||
}
|
||||
|
||||
fn eq(x: FixedDec, y: FixedDec) -> bool {
|
||||
if x.sign != y.sign || x.dec != y.dec {
|
||||
return false;
|
||||
}
|
||||
var i = 0u;
|
||||
var xp = x.parts;
|
||||
var yp = y.parts;
|
||||
while i < LEN - 2 {
|
||||
if xp[i] != yp[i] {
|
||||
return false;
|
||||
}
|
||||
i += 1u;
|
||||
}
|
||||
if abs(f32(xp[i]) - f32(yp[i])) > 1024 * 16 {
|
||||
return false;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
fn to_f32(value: FixedDec) -> f32 {
|
||||
var parts = value.parts;
|
||||
|
||||
var sign = value.sign * (1u << 31);
|
||||
var skip_count = 0;
|
||||
|
||||
while skip_count < ILEN && parts[skip_count] == 0 {
|
||||
skip_count += 1;
|
||||
}
|
||||
|
||||
if skip_count == ILEN {
|
||||
if value.sign == POS {
|
||||
return 0.0;
|
||||
} else {
|
||||
return -0.0;
|
||||
}
|
||||
}
|
||||
let v = parts[skip_count];
|
||||
var start = countLeadingZeros(v) + 1;
|
||||
let exp_i = (value.dec - skip_count) * 32 - i32(start);
|
||||
var frac_sh = 0u;
|
||||
var exp = 0u;
|
||||
if exp_i >= -127 {
|
||||
if exp_i == -127 {
|
||||
start -= 1u;
|
||||
}
|
||||
exp = u32(exp_i + 127);
|
||||
} else {
|
||||
frac_sh = u32(-(exp_i + 32 * 4 - 1));
|
||||
if frac_sh < 23 {
|
||||
start -= 1u;
|
||||
} else {
|
||||
return 0.0;
|
||||
}
|
||||
};
|
||||
var frac: u32;
|
||||
if start > 9 {
|
||||
let sh = start - 9;
|
||||
let next_i = skip_count + 1;
|
||||
var v2 = 0u;
|
||||
if next_i < ILEN {
|
||||
v2 = parts[next_i] >> (32 - sh);
|
||||
}
|
||||
frac = (v << sh) + v2;
|
||||
} else {
|
||||
frac = v >> (9 - start);
|
||||
};
|
||||
frac &= ~(1u << 23);
|
||||
let res = (frac >> frac_sh) + (exp << 23) + sign;
|
||||
return bitcast<f32>(res);
|
||||
}
|
||||
|
||||
const INV_SIGN_MASK: u32 = (1u << 31) - 1;
|
||||
const FRAC_BIT: u32 = 1u << 23;
|
||||
const FRAC_MASK: u32 = FRAC_BIT - 1;
|
||||
|
||||
fn from_f32(value: f32) -> FixedDec {
|
||||
let raw = bitcast<u32>(value) & INV_SIGN_MASK;
|
||||
var exp = i32(raw >> 23) - 127;
|
||||
var frac = raw & FRAC_MASK;
|
||||
var start = -exp;
|
||||
if exp == -127 {
|
||||
exp = -126;
|
||||
start = -exp;
|
||||
} else {
|
||||
frac += FRAC_BIT;
|
||||
start -= 1;
|
||||
}
|
||||
let end = -exp + 23;
|
||||
let start_i = div_euclid(start, 32);
|
||||
let end_i = div_euclid(end - 1, 32);
|
||||
var parts = array<u32, LEN>();
|
||||
var dec = -start_i;
|
||||
if start_i == end_i {
|
||||
let val = frac << u32(8 - rem_euclid(start, 32));
|
||||
if val != 0 {
|
||||
parts[0] = val;
|
||||
}
|
||||
} else {
|
||||
let s = u32(rem_euclid(end, 32));
|
||||
let val_high = frac >> s;
|
||||
let val_low = frac << (32 - s);
|
||||
var i = 0;
|
||||
if val_high != 0 {
|
||||
parts[0] = val_high;
|
||||
i += 1;
|
||||
} else {
|
||||
dec -= 1;
|
||||
}
|
||||
if val_low != 0 {
|
||||
parts[i] = val_low;
|
||||
}
|
||||
}
|
||||
if parts[0] == 0 && parts[1] == 0 {
|
||||
dec = 0;
|
||||
}
|
||||
return FixedDec(
|
||||
u32(value < 0.0),
|
||||
dec,
|
||||
parts,
|
||||
);
|
||||
}
|
||||
|
||||
fn zero() -> FixedDec {
|
||||
return FixedDec(0, 0, array<u32, LEN>());
|
||||
}
|
||||
|
||||
fn div_euclid(x: i32, y: i32) -> i32 {
|
||||
if x < 0 {
|
||||
return -((-x - 1) / y) - 1;
|
||||
}
|
||||
return x / y;
|
||||
}
|
||||
|
||||
fn rem_euclid(x: i32, y: i32) -> i32 {
|
||||
return x - div_euclid(x, y) * y;
|
||||
}
|
||||
123
src/client/render/compute/layout.rs
Normal file
123
src/client/render/compute/layout.rs
Normal file
@@ -0,0 +1,123 @@
|
||||
use wgpu::{PipelineCompilationOptions, ShaderStages};
|
||||
|
||||
use crate::client::render::util::ArrayBuffer;
|
||||
|
||||
use super::{
|
||||
util::{Storage, Texture},
|
||||
ComputeView,
|
||||
};
|
||||
|
||||
pub struct Layout {
|
||||
bind_layout: wgpu::BindGroupLayout,
|
||||
pipeline_layout: wgpu::PipelineLayout,
|
||||
format: wgpu::TextureFormat,
|
||||
pub output: Texture,
|
||||
pub view: Storage,
|
||||
pub work: ArrayBuffer<u32>,
|
||||
}
|
||||
|
||||
impl Layout {
|
||||
pub fn init(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration, len: usize) -> Self {
|
||||
let view = Storage::init_with(device, "view", ComputeView::default().bytes());
|
||||
let work = ArrayBuffer::init_with(
|
||||
device,
|
||||
"test",
|
||||
wgpu::BufferUsages::STORAGE,
|
||||
&work_vec(config.width, config.height, len),
|
||||
);
|
||||
|
||||
let desc = wgpu::TextureDescriptor {
|
||||
label: Some("compute output"),
|
||||
size: wgpu::Extent3d {
|
||||
width: config.width,
|
||||
height: config.height,
|
||||
depth_or_array_layers: 1,
|
||||
},
|
||||
mip_level_count: 1,
|
||||
sample_count: 1,
|
||||
dimension: wgpu::TextureDimension::D2,
|
||||
format: wgpu::TextureFormat::Rgba8Unorm,
|
||||
usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
|
||||
view_formats: &[],
|
||||
};
|
||||
let output = Texture::init(
|
||||
device,
|
||||
desc,
|
||||
wgpu::TextureViewDescriptor::default(),
|
||||
wgpu::SamplerDescriptor::default(),
|
||||
);
|
||||
|
||||
let bind_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
|
||||
entries: &[
|
||||
view.bind_group_layout_entry(0, true, wgpu::ShaderStages::COMPUTE),
|
||||
work.bind_group_layout_entry(
|
||||
1,
|
||||
wgpu::BufferBindingType::Storage { read_only: false },
|
||||
wgpu::ShaderStages::COMPUTE,
|
||||
),
|
||||
wgpu::BindGroupLayoutEntry {
|
||||
binding: 2,
|
||||
visibility: ShaderStages::COMPUTE,
|
||||
ty: wgpu::BindingType::StorageTexture {
|
||||
access: wgpu::StorageTextureAccess::WriteOnly,
|
||||
format: output.format(),
|
||||
view_dimension: wgpu::TextureViewDimension::D2,
|
||||
},
|
||||
count: None,
|
||||
},
|
||||
],
|
||||
label: Some("compute"),
|
||||
});
|
||||
|
||||
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
|
||||
label: Some("Tile Pipeline Layout"),
|
||||
bind_group_layouts: &[&bind_layout],
|
||||
push_constant_ranges: &[],
|
||||
});
|
||||
|
||||
Self {
|
||||
view,
|
||||
output,
|
||||
bind_layout,
|
||||
pipeline_layout,
|
||||
work,
|
||||
format: config.format,
|
||||
}
|
||||
}
|
||||
|
||||
pub fn bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
|
||||
device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
layout: &self.bind_layout,
|
||||
entries: &[
|
||||
self.view.bind_group_entry(0),
|
||||
self.work.bind_group_entry(1),
|
||||
self.output.view_bind_group_entry(2),
|
||||
],
|
||||
label: Some("voxel render"),
|
||||
})
|
||||
}
|
||||
|
||||
pub fn pipeline(
|
||||
&self,
|
||||
device: &wgpu::Device,
|
||||
shader: &wgpu::ShaderModule,
|
||||
) -> wgpu::ComputePipeline {
|
||||
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
|
||||
label: Some("Voxel Pipeline"),
|
||||
layout: Some(&self.pipeline_layout),
|
||||
entry_point: Some("main"),
|
||||
module: shader,
|
||||
cache: None,
|
||||
compilation_options: PipelineCompilationOptions::default(),
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
pub fn work_size(width: u32, height: u32, len: usize) -> usize {
|
||||
let varwidth = (2 + len) * 2;
|
||||
(width * height) as usize * (varwidth * 2 + 1)
|
||||
}
|
||||
|
||||
pub fn work_vec(width: u32, height: u32, len: usize) -> Vec<u32> {
|
||||
vec![0u32; work_size(width, height, len)]
|
||||
}
|
||||
110
src/client/render/compute/mod.rs
Normal file
110
src/client/render/compute/mod.rs
Normal file
@@ -0,0 +1,110 @@
|
||||
use std::ops::{Deref, DerefMut};
|
||||
|
||||
mod data;
|
||||
mod layout;
|
||||
|
||||
use super::*;
|
||||
pub use data::*;
|
||||
use layout::*;
|
||||
|
||||
pub struct ComputePipeline {
|
||||
layout: Layout,
|
||||
pipeline: wgpu::ComputePipeline,
|
||||
bind_group: wgpu::BindGroup,
|
||||
old_view: ComputeView,
|
||||
old_len: usize,
|
||||
}
|
||||
|
||||
const FIXED_SHADER: &str = include_str!("fixed.wgsl");
|
||||
const SHADER: &str = include_str!("shader.wgsl");
|
||||
|
||||
impl ComputePipeline {
|
||||
pub fn init(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration, len: usize) -> Self {
|
||||
let layout = Layout::init(device, config, len);
|
||||
Self {
|
||||
pipeline: layout.pipeline(device, &Self::shader(device, len)),
|
||||
bind_group: layout.bind_group(device),
|
||||
layout,
|
||||
old_view: ComputeView::default(),
|
||||
old_len: len,
|
||||
}
|
||||
}
|
||||
|
||||
pub fn update(
|
||||
&mut self,
|
||||
device: &wgpu::Device,
|
||||
encoder: &mut wgpu::CommandEncoder,
|
||||
belt: &mut wgpu::util::StagingBelt,
|
||||
camera: &Camera,
|
||||
size: &Vector2<u32>,
|
||||
len: usize,
|
||||
) {
|
||||
let mut view = ComputeView::from_camera_size(camera, size, false, len);
|
||||
if view != self.old_view {
|
||||
for (i, b) in 1u32.to_le_bytes().iter().enumerate() {
|
||||
view.bytes[i] = *b;
|
||||
}
|
||||
}
|
||||
if len != self.old_len {
|
||||
self.old_len = len;
|
||||
self.pipeline = self.pipeline(device, &Self::shader(device, len));
|
||||
self.work.set(work_vec(size.x, size.y, len));
|
||||
}
|
||||
let updated = self.work.update(device, encoder, belt)
|
||||
| self.view.update(device, encoder, belt, view.bytes());
|
||||
if updated {
|
||||
self.bind_group = self.layout.bind_group(device);
|
||||
}
|
||||
self.old_view = view;
|
||||
}
|
||||
|
||||
pub fn run(&self, encoder: &mut wgpu::CommandEncoder) {
|
||||
let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor::default());
|
||||
pass.set_pipeline(&self.pipeline);
|
||||
pass.set_bind_group(0, &self.bind_group, &[]);
|
||||
pass.dispatch_workgroups(240, 135, 1);
|
||||
}
|
||||
|
||||
pub fn resize(
|
||||
&mut self,
|
||||
device: &wgpu::Device,
|
||||
encoder: &mut wgpu::CommandEncoder,
|
||||
belt: &mut wgpu::util::StagingBelt,
|
||||
size: Vector2<u32>,
|
||||
len: usize,
|
||||
) {
|
||||
self.work.set(work_vec(size.x, size.y, len));
|
||||
self.old_len = len;
|
||||
self.output.resize(
|
||||
device,
|
||||
wgpu::Extent3d {
|
||||
width: size.x,
|
||||
height: size.y,
|
||||
depth_or_array_layers: 1,
|
||||
},
|
||||
);
|
||||
self.bind_group = self.layout.bind_group(device);
|
||||
}
|
||||
|
||||
pub fn shader(device: &wgpu::Device, len: usize) -> wgpu::ShaderModule {
|
||||
let string = FIXED_SHADER.to_string() + &SHADER.replace("REPLACE_LEN", &format!("{}", len));
|
||||
device.create_shader_module(wgpu::ShaderModuleDescriptor {
|
||||
label: Some("compute"),
|
||||
source: wgpu::ShaderSource::Wgsl(string.into()),
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
impl Deref for ComputePipeline {
|
||||
type Target = Layout;
|
||||
|
||||
fn deref(&self) -> &Self::Target {
|
||||
&self.layout
|
||||
}
|
||||
}
|
||||
|
||||
impl DerefMut for ComputePipeline {
|
||||
fn deref_mut(&mut self) -> &mut Self::Target {
|
||||
&mut self.layout
|
||||
}
|
||||
}
|
||||
102
src/client/render/compute/shader.wgsl
Normal file
102
src/client/render/compute/shader.wgsl
Normal file
@@ -0,0 +1,102 @@
|
||||
override CHUNK_POW: u32 = 10;
|
||||
const LEN: u32 = REPLACE_LENu;
|
||||
const ILEN: i32 = i32(LEN);
|
||||
const LEN2: u32 = LEN * 2;
|
||||
override WGX: u32 = 8;
|
||||
override WGY: u32 = 8;
|
||||
|
||||
struct View {
|
||||
reset: u32,
|
||||
level: i32,
|
||||
dims: vec2<u32>,
|
||||
scale: FixedDec,
|
||||
corner_x: FixedDec,
|
||||
corner_y: FixedDec,
|
||||
}
|
||||
|
||||
@group(0) @binding(0)
|
||||
var<storage> view: View;
|
||||
@group(0) @binding(1)
|
||||
var<storage, read_write> work: array<u32>;
|
||||
@group(0) @binding(2)
|
||||
var output: texture_storage_2d<rgba8unorm, write>;
|
||||
|
||||
@compute @workgroup_size(WGX, WGY, 1)
|
||||
fn main(
|
||||
@builtin(global_invocation_id) id: vec3<u32>
|
||||
) {
|
||||
if id.x > view.dims.x - 1 || id.y > view.dims.y - 1 {
|
||||
return;
|
||||
}
|
||||
// TODO: actually use width
|
||||
let varwidth = LEN + 2;
|
||||
let workwidth = varwidth * 2 + 1;
|
||||
let worki = (id.x * view.dims.y + id.y) * workwidth;
|
||||
let xidx = worki + 1;
|
||||
let yidx = xidx + varwidth;
|
||||
|
||||
// let dec = view.corner_x.dec;
|
||||
// var rel_x = FixedDec(POS, dec, array<u32, LEN>());
|
||||
// rel_x.parts[0] = id.x;
|
||||
// rel_x = shr(rel_x, view.level);
|
||||
// var rel_y = FixedDec(POS, dec, array<u32, LEN>());
|
||||
// rel_y.parts[0] = id.y;
|
||||
// rel_y = shr(rel_y, view.level);
|
||||
// let cx = add(view.corner_x, rel_x);
|
||||
// let cy = add(view.corner_y, rel_y);
|
||||
let fdims = vec2<f32>(view.dims);
|
||||
var stretch: vec2<f32>;
|
||||
if fdims.x < fdims.y {
|
||||
stretch = vec2(fdims.x / fdims.y, 1.0);
|
||||
} else {
|
||||
stretch = vec2(1.0, fdims.y / fdims.x);
|
||||
}
|
||||
let fpos = (vec2<f32>(id.xy) / fdims - 0.5) * stretch;
|
||||
let cx = add(mul(from_f32(fpos.x), view.scale), view.corner_x);
|
||||
let cy = add(mul(from_f32(fpos.y), view.scale), view.corner_y);
|
||||
var x: FixedDec;
|
||||
var y: FixedDec;
|
||||
var i = work[worki];
|
||||
if bool(view.reset) {
|
||||
x = zero();
|
||||
y = zero();
|
||||
i = 0;
|
||||
} else {
|
||||
x = FixedDec(work[xidx + 0], bitcast<i32>(work[xidx + 1]), array<u32, LEN>());
|
||||
y = FixedDec(work[yidx + 0], bitcast<i32>(work[yidx + 1]), array<u32, LEN>());
|
||||
for (var j = 0u; j < LEN; j += 1u) {
|
||||
x.parts[j] = work[xidx + 2 + j];
|
||||
y.parts[j] = work[yidx + 2 + j];
|
||||
}
|
||||
}
|
||||
let max = i + 1;
|
||||
let thresh = from_f32(2.0 * 2.0);
|
||||
loop {
|
||||
let x2 = mul(x, x);
|
||||
let y2 = mul(y, y);
|
||||
if gt(add(x2, y2), thresh) || i >= max {
|
||||
break;
|
||||
}
|
||||
let xy = mul(x, y);
|
||||
y = add(add(xy, xy), cy);
|
||||
x = add(sub(x2, y2), cx);
|
||||
i += 1u;
|
||||
}
|
||||
work[worki] = i;
|
||||
work[xidx + 0] = x.sign; work[xidx + 1] = bitcast<u32>(x.dec);
|
||||
work[yidx + 0] = y.sign; work[yidx + 1] = bitcast<u32>(y.dec);
|
||||
for (var j = 0u; j < LEN; j += 1u) {
|
||||
work[xidx + 2 + j] = x.parts[j];
|
||||
work[yidx + 2 + j] = y.parts[j];
|
||||
}
|
||||
var color = vec3<f32>(0.0, 0.0, 0.0);
|
||||
if i != max {
|
||||
let pi = 3.1415;
|
||||
let hue = f32(i) / 30.0;
|
||||
color.r = cos(hue);
|
||||
color.g = cos(hue - 2.0 * pi / 3.0);
|
||||
color.b = cos(hue - 4.0 * pi / 3.0);
|
||||
}
|
||||
textureStore(output, id.xy, vec4(color, 1.0));
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user