diff --git a/Cargo.lock b/Cargo.lock index d171dda27..f0f91dfab 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -2810,6 +2810,16 @@ dependencies = [ "num-traits", ] +[[package]] +name = "naga-pixelbender" +version = "0.1.0" +dependencies = [ + "anyhow", + "bitflags 2.3.1", + "naga", + "ruffle_render", +] + [[package]] name = "naga_oil" version = "0.7.0" @@ -3793,6 +3803,7 @@ name = "ruffle_render" version = "0.1.0" dependencies = [ "approx", + "byteorder", "clap", "downcast-rs", "enum-map", @@ -3803,6 +3814,8 @@ dependencies = [ "jpeg-decoder", "lru", "lyon", + "num-derive", + "num-traits", "png", "ruffle_wstr", "serde", @@ -3859,9 +3872,11 @@ dependencies = [ "futures", "gc-arena", "image", + "indexmap", "lru", "naga", "naga-agal", + "naga-pixelbender", "naga_oil", "ouroboros", "profiling", diff --git a/Cargo.toml b/Cargo.toml index b94aea89d..4141dc1f4 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -38,6 +38,7 @@ version = "0.1.0" gc-arena = { git = "https://github.com/kyren/gc-arena", rev = "63dab12871321e0e5ada10ff1f1de8f4cf1764f9" } tracing = "0.1.37" tracing-subscriber = { version = "0.3.17", features = ["env-filter"] } +naga = { version = "0.12.2", features = ["validate", "wgsl-out"] } # Don't optimize build scripts and macros. [profile.release.build-override] diff --git a/core/src/avm2/globals/flash/display.rs b/core/src/avm2/globals/flash/display.rs index 850218e2b..b266b1b80 100644 --- a/core/src/avm2/globals/flash/display.rs +++ b/core/src/avm2/globals/flash/display.rs @@ -11,6 +11,7 @@ pub mod loader_info; pub mod morph_shape; pub mod movie_clip; pub mod shader_data; +pub mod shader_job; pub mod shader_parameter; pub mod shape; pub mod simple_button; diff --git a/core/src/avm2/globals/flash/display/ShaderData.as b/core/src/avm2/globals/flash/display/ShaderData.as index cb02a501a..da4c0430c 100644 --- a/core/src/avm2/globals/flash/display/ShaderData.as +++ b/core/src/avm2/globals/flash/display/ShaderData.as @@ -2,6 +2,7 @@ package flash.display { import flash.utils.ByteArray; import __ruffle__.stub_constructor; + [Ruffle(InstanceAllocator)] public final dynamic class ShaderData { public function ShaderData(bytecode:ByteArray) { this.init(bytecode); diff --git a/core/src/avm2/globals/flash/display/ShaderJob.as b/core/src/avm2/globals/flash/display/ShaderJob.as index 2608287e0..36513306a 100644 --- a/core/src/avm2/globals/flash/display/ShaderJob.as +++ b/core/src/avm2/globals/flash/display/ShaderJob.as @@ -6,8 +6,13 @@ package flash.display { import flash.events.EventDispatcher; public class ShaderJob extends EventDispatcher { + + private var _shader:Shader; + private var _target:Object; public function ShaderJob(shader:Shader = null, target:Object = null, width:int = 0, height:int = 0) { + this._shader = shader; + this._target = target; stub_constructor("flash.display.ShaderJob"); } @@ -15,9 +20,7 @@ package flash.display { stub_method("flash.display.ShaderJob", "cancel") } - public function start(waitForCompletion:Boolean = false):void { - stub_method("flash.display.ShaderJob", "start") - } + public native function start(waitForCompletion:Boolean = false):void; public function get height():int { stub_getter("flash.display.ShaderJob", "height"); @@ -34,12 +37,19 @@ package flash.display { } public function get shader():Shader { - stub_getter("flash.display.ShaderJob", "shader"); - return null; + return this._shader; } public function set shader(value:Shader):void { - stub_setter("flash.display.ShaderJob", "shader"); + this._shader = value; + } + + public function get target():Object { + return this._target; + } + + public function set target(value:Object):void { + this._target = value; } } } \ No newline at end of file diff --git a/core/src/avm2/globals/flash/display/shader_data.rs b/core/src/avm2/globals/flash/display/shader_data.rs index 0825e8ac6..385625dd2 100644 --- a/core/src/avm2/globals/flash/display/shader_data.rs +++ b/core/src/avm2/globals/flash/display/shader_data.rs @@ -1,12 +1,18 @@ +use ruffle_render::pixel_bender::{ + parse_shader, PixelBenderParam, PixelBenderParamQualifier, OUT_COORD_NAME, +}; + use crate::{ avm2::{ parameters::ParametersExt, string::AvmString, Activation, Error, Object, TObject, Value, }, - pixel_bender::{PixelBenderParam, PixelBenderParamQualifier}, + pixel_bender::PixelBenderTypeExt, }; use super::shader_parameter::make_shader_parameter; +pub use crate::avm2::object::shader_data_allocator; + /// Implements `ShaderData.init`, which is called from the constructor pub fn init<'gc>( activation: &mut Activation<'_, 'gc>, @@ -16,11 +22,11 @@ pub fn init<'gc>( let mut this = this.unwrap(); let bytecode = args.get_object(activation, 0, "bytecode")?; let bytecode = bytecode.as_bytearray().unwrap(); - let shader = crate::pixel_bender::parse_shader(bytecode.bytes()); + let shader = parse_shader(bytecode.bytes()).expect("Failed to parse PixelBender"); - for meta in shader.metadata { + for meta in &shader.metadata { let name = AvmString::new_utf8(activation.context.gc_context, &meta.key); - let value = meta.value.into_avm2_value(activation)?; + let value = meta.value.as_avm2_value(activation)?; this.set_public_property(name, value, activation)?; } this.set_public_property( @@ -29,13 +35,14 @@ pub fn init<'gc>( activation, )?; - for (index, param) in shader.params.into_iter().enumerate() { + for (index, param) in shader.params.iter().enumerate() { let name = match ¶m { PixelBenderParam::Normal { name, qualifier, .. } => { // Neither of these show up in Flash Player - if name == "_OutCoord" || matches!(qualifier, PixelBenderParamQualifier::Output) { + if name == OUT_COORD_NAME || matches!(qualifier, PixelBenderParamQualifier::Output) + { continue; } name @@ -47,5 +54,15 @@ pub fn init<'gc>( let param_obj = make_shader_parameter(activation, param, index)?; this.set_public_property(name, param_obj, activation)?; } + + let shader_handle = activation + .context + .renderer + .compile_pixelbender_shader(shader) + .expect("Failed to compile PixelBender shader"); + + this.as_shader_data() + .unwrap() + .set_pixel_bender_shader(shader_handle, activation.context.gc_context); Ok(Value::Undefined) } diff --git a/core/src/avm2/globals/flash/display/shader_job.rs b/core/src/avm2/globals/flash/display/shader_job.rs new file mode 100644 index 000000000..e4be03053 --- /dev/null +++ b/core/src/avm2/globals/flash/display/shader_job.rs @@ -0,0 +1,167 @@ +use ruffle_render::{ + bitmap::PixelRegion, + pixel_bender::{ + PixelBenderParam, PixelBenderParamQualifier, PixelBenderShaderArgument, PixelBenderType, + OUT_COORD_NAME, + }, +}; + +use crate::{ + avm2::{string::AvmString, Activation, Error, Object, TObject, Value}, + avm2_stub_method, + pixel_bender::PixelBenderTypeExt, +}; + +/// Implements `ShaderJob.start`. +pub fn start<'gc>( + activation: &mut Activation<'_, 'gc>, + this: Option>, + _args: &[Value<'gc>], +) -> Result, Error<'gc>> { + let this = this.unwrap(); + + avm2_stub_method!( + activation, + "flash.display.ShaderJob", + "start", + "async execution and non-BitmapData inputs" + ); + + // FIXME - determine what errors Flash Player throws here + // instead of using `expect` + let shader = this + .get_public_property("shader", activation)? + .as_object() + .expect("Missing Shader object"); + let shader_data = shader + .get_public_property("data", activation)? + .as_object() + .expect("Missing ShaderData object") + .as_shader_data() + .expect("ShaderData object is not a ShaderData instance"); + + let shader_handle = shader_data.pixel_bender_shader(); + let shader_handle = shader_handle + .as_ref() + .expect("ShaderData object has no shader"); + let shader = shader_handle.0.parsed_shader(); + + let arguments: Vec<_> = shader + .params + .iter() + .enumerate() + .flat_map(|(index, param)| { + match param { + PixelBenderParam::Normal { + qualifier, + param_type, + name, + .. + } => { + if matches!(qualifier, PixelBenderParamQualifier::Output) { + return None; + } + + if name == OUT_COORD_NAME { + // Pass in a dummy value - this will be ignored in favor of the actual pixel coordinate + return Some(PixelBenderShaderArgument::ValueInput { + index: index as u8, + value: PixelBenderType::TFloat2(f32::NAN, f32::NAN), + }); + } + let shader_param = shader_data + .get_public_property( + AvmString::new_utf8(activation.context.gc_context, name), + activation, + ) + .expect("Missing normal property"); + + let shader_param = shader_param + .as_object() + .expect("Shader property is not an object"); + + let value = shader_param + .get_public_property("value", activation) + .expect("Missing value property"); + let pb_val = PixelBenderType::from_avm2_value(activation, value, param_type) + .expect("Failed to convert AVM2 value to PixelBenderType"); + + Some(PixelBenderShaderArgument::ValueInput { + index: index as u8, + value: pb_val, + }) + } + PixelBenderParam::Texture { + index, + channels, + name, + } => { + let shader_input = shader_data + .get_public_property( + AvmString::new_utf8(activation.context.gc_context, name), + activation, + ) + .expect("Missing property") + .as_object() + .expect("Shader input is not an object"); + + let input = shader_input + .get_public_property("input", activation) + .expect("Missing input property"); + let input = input + .as_object() + .expect("ShaderInput.input is not an object"); + + let bitmap = input.as_bitmap_data().expect( + "ShaderInput.input is not a BitmapData (FIXE - support other types)", + ); + + // FIXME - this really only needs to be a CPU->GPU sync + let bitmap = bitmap.sync(); + let mut bitmap_data = bitmap.write(activation.context.gc_context); + bitmap_data.update_dirty_texture(activation.context.renderer); + + Some(PixelBenderShaderArgument::ImageInput { + index: *index, + channels: *channels, + name: name.clone(), + texture: bitmap_data + .bitmap_handle(activation.context.renderer) + .expect("Missing input BitmapHandle"), + }) + } + } + }) + .collect(); + + let target = this + .get_public_property("target", activation)? + .as_object() + .expect("ShaderJob.target is not an object"); + + let target_bitmap = target + .as_bitmap_data() + .expect("ShaderJob.target is not a BitmapData (FIXE - support other types)") + .sync(); + + // Perform both a GPU->CPU and CPU->GPU sync before writing to it. + // FIXME - are both necessary? + let mut target_bitmap_data = target_bitmap.write(activation.context.gc_context); + target_bitmap_data.update_dirty_texture(activation.context.renderer); + + let target_handle = target_bitmap_data + .bitmap_handle(activation.context.renderer) + .expect("Missing handle"); + + let sync_handle = activation + .context + .renderer + .run_pixelbender_shader(shader_handle.clone(), &arguments, target_handle) + .expect("Failed to run shader"); + + let width = target_bitmap_data.width(); + let height = target_bitmap_data.height(); + target_bitmap_data.set_gpu_dirty(sync_handle, PixelRegion::for_whole_size(width, height)); + + Ok(Value::Undefined) +} diff --git a/core/src/avm2/globals/flash/display/shader_parameter.rs b/core/src/avm2/globals/flash/display/shader_parameter.rs index 15f98bc70..d196762b2 100644 --- a/core/src/avm2/globals/flash/display/shader_parameter.rs +++ b/core/src/avm2/globals/flash/display/shader_parameter.rs @@ -1,11 +1,13 @@ +use ruffle_render::pixel_bender::PixelBenderParam; + use crate::{ avm2::{string::AvmString, Activation, Error, Multiname, TObject, Value}, - pixel_bender::PixelBenderParam, + pixel_bender::PixelBenderTypeExt, }; pub fn make_shader_parameter<'gc>( activation: &mut Activation<'_, 'gc>, - param: PixelBenderParam, + param: &PixelBenderParam, index: usize, ) -> Result, Error<'gc>> { let ns = activation.avm2().flash_display_internal; @@ -29,7 +31,7 @@ pub fn make_shader_parameter<'gc>( obj.set_property(&Multiname::new(ns, "_type"), type_name.into(), activation)?; for meta in metadata { let name = AvmString::new_utf8(activation.context.gc_context, &meta.key); - let value = meta.value.clone().into_avm2_value(activation)?; + let value = meta.value.clone().as_avm2_value(activation)?; obj.set_public_property(name, value, activation)?; } obj.set_public_property( @@ -47,7 +49,7 @@ pub fn make_shader_parameter<'gc>( .construct(activation, &[])?; obj.set_property( &Multiname::new(ns, "_channels"), - channels.into(), + (*channels).into(), activation, )?; obj.set_property(&Multiname::new(ns, "_index"), index.into(), activation)?; diff --git a/core/src/avm2/object.rs b/core/src/avm2/object.rs index 9d850dbf5..ddb49c035 100644 --- a/core/src/avm2/object.rs +++ b/core/src/avm2/object.rs @@ -51,6 +51,7 @@ mod proxy_object; mod qname_object; mod regexp_object; mod script_object; +mod shader_data_object; mod sound_object; mod soundchannel_object; mod stage3d_object; @@ -104,6 +105,9 @@ pub use crate::avm2::object::proxy_object::{proxy_allocator, ProxyObject, ProxyO pub use crate::avm2::object::qname_object::{q_name_allocator, QNameObject, QNameObjectWeak}; pub use crate::avm2::object::regexp_object::{reg_exp_allocator, RegExpObject, RegExpObjectWeak}; pub use crate::avm2::object::script_object::{ScriptObject, ScriptObjectData, ScriptObjectWeak}; +pub use crate::avm2::object::shader_data_object::{ + shader_data_allocator, ShaderDataObject, ShaderDataObjectWeak, +}; pub use crate::avm2::object::sound_object::{ sound_allocator, QueuedPlay, SoundData, SoundObject, SoundObjectWeak, }; @@ -166,6 +170,7 @@ pub use crate::avm2::object::xml_object::{xml_allocator, XmlObject, XmlObjectWea TextureObject(TextureObject<'gc>), Program3DObject(Program3DObject<'gc>), NetStreamObject(NetStreamObject<'gc>), + ShaderDataObject(ShaderDataObject<'gc>), } )] pub trait TObject<'gc>: 'gc + Collect + Debug + Into> + Clone + Copy { @@ -1281,6 +1286,10 @@ pub trait TObject<'gc>: 'gc + Collect + Debug + Into> + Clone + Copy None } + fn as_shader_data(&self) -> Option> { + None + } + /// Initialize the bitmap data in this object, if it's capable of /// supporting said data. /// @@ -1394,6 +1403,7 @@ impl<'gc> Object<'gc> { Self::TextureObject(o) => WeakObject::TextureObject(TextureObjectWeak(GcCell::downgrade(o.0))), Self::Program3DObject(o) => WeakObject::Program3DObject(Program3DObjectWeak(GcCell::downgrade(o.0))), Self::NetStreamObject(o) => WeakObject::NetStreamObject(NetStreamObjectWeak(GcCell::downgrade(o.0))), + Self::ShaderDataObject(o) => WeakObject::ShaderDataObject(ShaderDataObjectWeak(GcCell::downgrade(o.0))), } } } @@ -1448,6 +1458,7 @@ pub enum WeakObject<'gc> { TextureObject(TextureObjectWeak<'gc>), Program3DObject(Program3DObjectWeak<'gc>), NetStreamObject(NetStreamObjectWeak<'gc>), + ShaderDataObject(ShaderDataObjectWeak<'gc>), } impl<'gc> WeakObject<'gc> { @@ -1485,6 +1496,7 @@ impl<'gc> WeakObject<'gc> { Self::TextureObject(o) => TextureObject(o.0.upgrade(mc)?).into(), Self::Program3DObject(o) => Program3DObject(o.0.upgrade(mc)?).into(), Self::NetStreamObject(o) => NetStreamObject(o.0.upgrade(mc)?).into(), + Self::ShaderDataObject(o) => ShaderDataObject(o.0.upgrade(mc)?).into(), }) } } diff --git a/core/src/avm2/object/shader_data_object.rs b/core/src/avm2/object/shader_data_object.rs new file mode 100644 index 000000000..fe262f537 --- /dev/null +++ b/core/src/avm2/object/shader_data_object.rs @@ -0,0 +1,87 @@ +//! Object representation for `ShaderData` + +use crate::avm2::activation::Activation; +use crate::avm2::object::script_object::ScriptObjectData; +use crate::avm2::object::{ClassObject, Object, ObjectPtr, TObject}; +use crate::avm2::value::Value; +use crate::avm2::Error; +use core::fmt; +use gc_arena::{Collect, GcCell, GcWeakCell, MutationContext}; +use ruffle_render::pixel_bender::PixelBenderShaderHandle; +use std::cell::{Ref, RefMut}; + +/// A class instance allocator that allocates ShaderData objects. +pub fn shader_data_allocator<'gc>( + class: ClassObject<'gc>, + activation: &mut Activation<'_, 'gc>, +) -> Result, Error<'gc>> { + let base = ScriptObjectData::new(class); + + Ok(ShaderDataObject(GcCell::allocate( + activation.context.gc_context, + ShaderDataObjectData { base, shader: None }, + )) + .into()) +} + +#[derive(Clone, Collect, Copy)] +#[collect(no_drop)] +pub struct ShaderDataObject<'gc>(pub GcCell<'gc, ShaderDataObjectData<'gc>>); + +#[derive(Clone, Collect, Copy, Debug)] +#[collect(no_drop)] +pub struct ShaderDataObjectWeak<'gc>(pub GcWeakCell<'gc, ShaderDataObjectData<'gc>>); + +impl fmt::Debug for ShaderDataObject<'_> { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + f.debug_struct("ShaderDataObject") + .field("ptr", &self.0.as_ptr()) + .finish() + } +} + +impl<'gc> ShaderDataObject<'gc> { + pub fn pixel_bender_shader(&self) -> Ref<'_, Option> { + Ref::map(self.0.read(), |read| &read.shader) + } + + pub fn set_pixel_bender_shader( + &self, + shader: PixelBenderShaderHandle, + mc: MutationContext<'gc, '_>, + ) { + self.0.write(mc).shader = Some(shader); + } +} + +#[derive(Collect)] +#[collect(no_drop)] +pub struct ShaderDataObjectData<'gc> { + /// Base script object + base: ScriptObjectData<'gc>, + + #[collect(require_static)] + shader: Option, +} + +impl<'gc> TObject<'gc> for ShaderDataObject<'gc> { + fn base(&self) -> Ref> { + Ref::map(self.0.read(), |read| &read.base) + } + + fn base_mut(&self, mc: MutationContext<'gc, '_>) -> RefMut> { + RefMut::map(self.0.write(mc), |write| &mut write.base) + } + + fn as_ptr(&self) -> *const ObjectPtr { + self.0.as_ptr() as *const ObjectPtr + } + + fn value_of(&self, _mc: MutationContext<'gc, '_>) -> Result, Error<'gc>> { + Ok(Value::Object(Object::from(*self))) + } + + fn as_shader_data(&self) -> Option> { + Some(*self) + } +} diff --git a/core/src/lib.rs b/core/src/lib.rs index ce6cca383..d202daf90 100644 --- a/core/src/lib.rs +++ b/core/src/lib.rs @@ -34,7 +34,7 @@ mod library; pub mod limits; pub mod loader; mod locale; -mod pixel_bender; +pub mod pixel_bender; mod player; mod prelude; mod streams; diff --git a/core/src/pixel_bender.rs b/core/src/pixel_bender.rs index 682a2d41c..71960b0c7 100644 --- a/core/src/pixel_bender.rs +++ b/core/src/pixel_bender.rs @@ -1,535 +1,142 @@ -//! Pixel bender bytecode parsing code. -//! This is heavling based on https://github.com/jamesward/pbjas and https://github.com/HaxeFoundation/format/tree/master/format/pbj - -#[cfg(test)] -mod tests; - -use byteorder::{BigEndian, LittleEndian, ReadBytesExt}; -use num_traits::FromPrimitive; -use std::{ - fmt::{Display, Formatter}, - io::Read, -}; +use ruffle_render::pixel_bender::{PixelBenderType, PixelBenderTypeOpcode}; use crate::{ - avm2::{Activation, ArrayObject, ArrayStorage, Error, Value}, + avm2::{Activation, ArrayObject, ArrayStorage, Error, TObject, Value}, ecma_conversions::f64_to_wrapping_i32, string::AvmString, }; -#[repr(u8)] -#[derive(Debug, Clone, PartialEq)] -pub enum PixelBenderType { - TFloat(f32) = 0x1, - TFloat2(f32, f32) = 0x2, - TFloat3(f32, f32, f32) = 0x3, - TFloat4(f32, f32, f32, f32) = 0x4, - TFloat2x2([f32; 4]) = 0x5, - TFloat3x3([f32; 9]) = 0x6, - TFloat4x4([f32; 16]) = 0x7, - TInt(i16) = 0x8, - TInt2(i16, i16) = 0x9, - TInt3(i16, i16, i16) = 0xA, - TInt4(i16, i16, i16, i16) = 0xB, - TString(String) = 0xC, +pub trait PixelBenderTypeExt { + fn from_avm2_value<'gc>( + activation: &mut Activation<'_, 'gc>, + value: Value<'gc>, + kind: &PixelBenderTypeOpcode, + ) -> Result> + where + Self: Sized; + + fn as_avm2_value<'gc>( + &self, + activation: &mut Activation<'_, 'gc>, + ) -> Result, Error<'gc>>; } -impl PixelBenderType { - pub fn into_avm2_value<'gc>( - self, +impl PixelBenderTypeExt for PixelBenderType { + fn from_avm2_value<'gc>( + activation: &mut Activation<'_, 'gc>, + value: Value<'gc>, + kind: &PixelBenderTypeOpcode, + ) -> Result> + where + Self: Sized, + { + let is_float = matches!( + kind, + PixelBenderTypeOpcode::TFloat + | PixelBenderTypeOpcode::TFloat2 + | PixelBenderTypeOpcode::TFloat3 + | PixelBenderTypeOpcode::TFloat4 + ); + + match value { + Value::String(s) => Ok(PixelBenderType::TString(s.to_string())), + Value::Number(n) => Ok(PixelBenderType::TFloat(n as f32)), + Value::Integer(i) => Ok(PixelBenderType::TInt(i as i16)), + Value::Object(o) => { + if let Some(array) = o.as_array_storage() { + if is_float { + let mut vals = array.iter().map(|val| { + val.expect("Array with hole") + .coerce_to_number(activation) + .unwrap() as f32 + }); + match kind { + PixelBenderTypeOpcode::TFloat => { + Ok(PixelBenderType::TFloat(vals.next().unwrap())) + } + PixelBenderTypeOpcode::TFloat2 => Ok(PixelBenderType::TFloat2( + vals.next().unwrap(), + vals.next().unwrap(), + )), + PixelBenderTypeOpcode::TFloat3 => Ok(PixelBenderType::TFloat3( + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + )), + PixelBenderTypeOpcode::TFloat4 => Ok(PixelBenderType::TFloat4( + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + )), + _ => unreachable!("Unexpected float kind {kind:?}"), + } + } else { + let mut vals = array.iter().map(|val| { + val.expect("Array with hole") + .coerce_to_i32(activation) + .unwrap() as i16 + }); + match kind { + PixelBenderTypeOpcode::TInt => { + Ok(PixelBenderType::TInt(vals.next().unwrap())) + } + PixelBenderTypeOpcode::TInt2 => Ok(PixelBenderType::TInt2( + vals.next().unwrap(), + vals.next().unwrap(), + )), + PixelBenderTypeOpcode::TInt3 => Ok(PixelBenderType::TInt3( + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + )), + PixelBenderTypeOpcode::TInt4 => Ok(PixelBenderType::TInt4( + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + vals.next().unwrap(), + )), + _ => unreachable!("Unexpected int kind {kind:?}"), + } + } + } else { + panic!("Unexpected object {o:?}") + } + } + _ => panic!("Unexpected value {value:?}"), + } + } + fn as_avm2_value<'gc>( + &self, activation: &mut Activation<'_, 'gc>, ) -> Result, Error<'gc>> { // Flash appears to use a uint/int if the float has no fractional part - let cv = |f: f32| -> Value<'gc> { + let cv = |f: &f32| -> Value<'gc> { if f.fract() == 0.0 { - f64_to_wrapping_i32(f as f64).into() + f64_to_wrapping_i32(*f as f64).into() } else { - f.into() + (*f).into() } }; let vals: Vec> = match self { PixelBenderType::TString(string) => { return Ok(AvmString::new_utf8(activation.context.gc_context, string).into()); } - PixelBenderType::TInt(i) => return Ok(i.into()), + PixelBenderType::TInt(i) => return Ok((*i).into()), PixelBenderType::TFloat(f) => vec![cv(f)], PixelBenderType::TFloat2(f1, f2) => vec![cv(f1), cv(f2)], PixelBenderType::TFloat3(f1, f2, f3) => vec![cv(f1), cv(f2), cv(f3)], PixelBenderType::TFloat4(f1, f2, f3, f4) => vec![cv(f1), cv(f2), cv(f3), cv(f4)], - PixelBenderType::TFloat2x2(floats) => floats.iter().map(|f| cv(*f)).collect(), - PixelBenderType::TFloat3x3(floats) => floats.iter().map(|f| cv(*f)).collect(), - PixelBenderType::TFloat4x4(floats) => floats.iter().map(|f| cv(*f)).collect(), - PixelBenderType::TInt2(i1, i2) => vec![i1.into(), i2.into()], - PixelBenderType::TInt3(i1, i2, i3) => vec![i1.into(), i2.into(), i3.into()], + PixelBenderType::TFloat2x2(floats) => floats.iter().map(|f| cv(f)).collect(), + PixelBenderType::TFloat3x3(floats) => floats.iter().map(|f| cv(f)).collect(), + PixelBenderType::TFloat4x4(floats) => floats.iter().map(|f| cv(f)).collect(), + PixelBenderType::TInt2(i1, i2) => vec![(*i1).into(), (*i2).into()], + PixelBenderType::TInt3(i1, i2, i3) => vec![(*i1).into(), (*i2).into(), (*i3).into()], PixelBenderType::TInt4(i1, i2, i3, i4) => { - vec![i1.into(), i2.into(), i3.into(), i4.into()] + vec![(*i1).into(), (*i2).into(), (*i3).into(), (*i4).into()] } }; let storage = ArrayStorage::from_args(&vals); Ok(ArrayObject::from_storage(activation, storage)?.into()) } } - -// FIXME - come up with a way to reduce duplication here -#[derive(num_derive::FromPrimitive, Debug, PartialEq)] -pub enum PixelBenderTypeOpcode { - TFloat = 0x1, - TFloat2 = 0x2, - TFloat3 = 0x3, - TFloat4 = 0x4, - TFloat2x2 = 0x5, - TFloat3x3 = 0x6, - TFloat4x4 = 0x7, - TInt = 0x8, - TInt2 = 0x9, - TInt3 = 0xA, - TInt4 = 0xB, - TString = 0xC, -} - -#[derive(num_derive::FromPrimitive, Debug, PartialEq)] -pub enum PixelBenderParamQualifier { - Input = 1, - Output = 2, -} - -impl Display for PixelBenderTypeOpcode { - fn fmt(&self, f: &mut Formatter<'_>) -> std::fmt::Result { - write!( - f, - "{}", - match self { - PixelBenderTypeOpcode::TFloat => "float", - PixelBenderTypeOpcode::TFloat2 => "float2", - PixelBenderTypeOpcode::TFloat3 => "float3", - PixelBenderTypeOpcode::TFloat4 => "float4", - PixelBenderTypeOpcode::TFloat2x2 => "matrix2x2", - PixelBenderTypeOpcode::TFloat3x3 => "matrix3x3", - PixelBenderTypeOpcode::TFloat4x4 => "matrix4x4", - PixelBenderTypeOpcode::TInt => "int", - PixelBenderTypeOpcode::TInt2 => "int2", - PixelBenderTypeOpcode::TInt3 => "int3", - PixelBenderTypeOpcode::TInt4 => "int4", - PixelBenderTypeOpcode::TString => "string", - } - ) - } -} - -#[derive(num_derive::FromPrimitive, Debug, PartialEq)] -pub enum Opcode { - Nop = 0x0, - Add = 0x1, - Sub = 0x2, - Mul = 0x3, - Rcp = 0x4, - Div = 0x5, - Atan2 = 0x6, - Pow = 0x7, - Mod = 0x8, - Min = 0x9, - Max = 0xA, - Step = 0xB, - Sin = 0xC, - Cos = 0xD, - Tan = 0xE, - Asin = 0xF, - Acos = 0x10, - Atan = 0x11, - Exp = 0x12, - Exp2 = 0x13, - Log = 0x14, - Log2 = 0x15, - Sqrt = 0x16, - RSqrt = 0x17, - Abs = 0x18, - Sign = 0x19, - Floor = 0x1A, - Ceil = 0x1B, - Fract = 0x1C, - Mov = 0x1D, - FloatToInt = 0x1E, - IntToFloat = 0x1F, - MatMatMul = 0x20, - VecMatMul = 0x21, - MatVecMul = 0x22, - Normalize = 0x23, - Length = 0x24, - Distance = 0x25, - DotProduct = 0x26, - CrossProduct = 0x27, - Equal = 0x28, - NotEqual = 0x29, - LessThan = 0x2A, - LessThanEqual = 0x2B, - LogicalNot = 0x2C, - LogicalAnd = 0x2D, - LogicalOr = 0x2E, - LogicalXor = 0x2F, - SampleNearest = 0x30, - SampleLinear = 0x31, - LoadIntOrFloat = 0x32, - Loop = 0x33, - If = 0x34, - Else = 0x35, - EndIf = 0x36, - FloatToBool = 0x37, - BoolToFloat = 0x38, - IntToBool = 0x39, - BoolToInt = 0x3A, - VectorEqual = 0x3B, - VectorNotEqual = 0x3C, - BoolAny = 0x3D, - BoolAll = 0x3E, - PBJMeta1 = 0xA0, - PBJParam = 0xA1, - PBJMeta2 = 0xA2, - PBJParamTexture = 0xA3, - Name = 0xA4, - Version = 0xA5, -} - -#[derive(Debug, PartialEq)] -pub enum Operation { - Nop, - Normal { - opcode: Opcode, - dst: u16, - mask: u8, - src: u32, - other: u8, - }, - LoadInt { - dst: u16, - mask: u8, - val: i32, - }, - LoadFloat { - dst: u16, - mask: u8, - val: f32, - }, - If { - src: u32, - }, - SampleNearest { - dst: u16, - src: u32, - mask: u8, - tf: u8, - }, - SampleLinear { - dst: u16, - src: u32, - mask: u8, - tf: u8, - }, - Else, - EndIf, -} - -#[derive(Debug, PartialEq)] -pub struct PixelBenderShader { - pub name: String, - pub version: i32, - pub params: Vec, - pub metadata: Vec, - pub operations: Vec, -} - -#[derive(Debug, PartialEq)] -pub enum PixelBenderParam { - Normal { - qualifier: PixelBenderParamQualifier, - param_type: PixelBenderTypeOpcode, - reg: u16, - mask: u8, - name: String, - metadata: Vec, - }, - Texture { - index: u8, - channels: u8, - name: String, - }, -} - -#[derive(Debug, Clone, PartialEq)] -pub struct PixelBenderMetadata { - pub key: String, - pub value: PixelBenderType, -} - -/// Parses PixelBender bytecode -pub fn parse_shader(mut data: &[u8]) -> PixelBenderShader { - let mut shader = PixelBenderShader { - name: String::new(), - version: 0, - params: Vec::new(), - metadata: Vec::new(), - operations: Vec::new(), - }; - let data = &mut data; - let mut metadata = Vec::new(); - while !data.is_empty() { - read_op(data, &mut shader, &mut metadata).unwrap(); - } - // Any metadata left in the vec is associated with our final parameter. - apply_metadata(&mut shader, &mut metadata); - shader -} - -fn read_op( - data: &mut R, - shader: &mut PixelBenderShader, - metadata: &mut Vec, -) -> Result<(), Box> { - let raw = data.read_u8()?; - let opcode = Opcode::from_u8(raw).expect("Unknown opcode"); - match opcode { - Opcode::Nop => { - assert_eq!(data.read_u32::()?, 0); - assert_eq!(data.read_u16::()?, 0); - shader.operations.push(Operation::Nop); - } - Opcode::PBJMeta1 | Opcode::PBJMeta2 => { - let meta_type = data.read_u8()?; - let meta_key = read_string(data)?; - let meta_value = read_value(data, PixelBenderTypeOpcode::from_u8(meta_type).unwrap())?; - metadata.push(PixelBenderMetadata { - key: meta_key, - value: meta_value, - }); - } - Opcode::PBJParam => { - let qualifier = data.read_u8()?; - let param_type = data.read_u8()?; - let reg = data.read_u16::()?; - let mask = data.read_u8()?; - let name = read_string(data)?; - - let param_type = PixelBenderTypeOpcode::from_u8(param_type).unwrap_or_else(|| { - panic!("Unexpected param type {param_type}"); - }); - let qualifier = PixelBenderParamQualifier::from_u8(qualifier) - .unwrap_or_else(|| panic!("Unexpected param qualifier {qualifier:?}")); - apply_metadata(shader, metadata); - - shader.params.push(PixelBenderParam::Normal { - qualifier, - param_type, - reg, - mask, - name, - metadata: Vec::new(), - }) - } - Opcode::PBJParamTexture => { - let index = data.read_u8()?; - let channels = data.read_u8()?; - let name = read_string(data)?; - apply_metadata(shader, metadata); - - shader.params.push(PixelBenderParam::Texture { - index, - channels, - name, - }); - } - Opcode::Name => { - let len = data.read_u16::()?; - let mut string_bytes = vec![0; len as usize]; - data.read_exact(&mut string_bytes)?; - shader.name = String::from_utf8(string_bytes)?; - } - Opcode::Version => { - shader.version = data.read_i32::()?; - } - Opcode::If => { - assert_eq!(read_uint24(data)?, 0); - let src = read_uint24(data)?; - assert_eq!(data.read_u8()?, 0); - shader.operations.push(Operation::If { src }); - } - Opcode::Else => { - assert_eq!(data.read_u32::()?, 0); - assert_eq!(read_uint24(data)?, 0); - shader.operations.push(Operation::Else); - } - Opcode::EndIf => { - assert_eq!(data.read_u32::()?, 0); - assert_eq!(read_uint24(data)?, 0); - shader.operations.push(Operation::EndIf); - } - Opcode::LoadIntOrFloat => { - let dst = data.read_u16::()?; - let mask = data.read_u8()?; - assert_eq!(mask & 0xF, 0); - if dst & 0x8000 != 0 { - let val = data.read_i32::()?; - shader.operations.push(Operation::LoadInt { - dst: dst - 0x8000, - mask, - val, - }) - } else { - let val = read_float(data)?; - shader - .operations - .push(Operation::LoadFloat { dst, mask, val }) - } - } - Opcode::SampleNearest | Opcode::SampleLinear => { - let dst = data.read_u16::()?; - let mask = data.read_u8()?; - let src = read_uint24(data)?; - let tf = data.read_u8()?; - match opcode { - Opcode::SampleNearest => { - shader - .operations - .push(Operation::SampleNearest { dst, mask, src, tf }) - } - Opcode::SampleLinear => { - shader - .operations - .push(Operation::SampleLinear { dst, mask, src, tf }) - } - _ => unreachable!(), - } - } - _ => { - let dst = data.read_u16::()?; - let mask = data.read_u8()?; - let src = read_uint24(data)?; - assert_eq!(data.read_u8()?, 0, "Unexpected u8 for opcode {opcode:?}"); - shader.operations.push(Operation::Normal { - opcode, - dst, - mask, - src, - other: 0, - }) - } - }; - Ok(()) -} - -fn read_string(data: &mut R) -> Result> { - let mut string = String::new(); - let mut b = data.read_u8()?; - while b != 0 { - string.push(b as char); - b = data.read_u8()?; - } - Ok(string) -} - -fn read_float(data: &mut R) -> Result> { - Ok(data.read_f32::()?) -} - -fn read_value( - data: &mut R, - opcode: PixelBenderTypeOpcode, -) -> Result> { - match opcode { - PixelBenderTypeOpcode::TFloat => Ok(PixelBenderType::TFloat(read_float(data)?)), - PixelBenderTypeOpcode::TFloat2 => Ok(PixelBenderType::TFloat2( - read_float(data)?, - read_float(data)?, - )), - PixelBenderTypeOpcode::TFloat3 => Ok(PixelBenderType::TFloat3( - read_float(data)?, - read_float(data)?, - read_float(data)?, - )), - PixelBenderTypeOpcode::TFloat4 => Ok(PixelBenderType::TFloat4( - read_float(data)?, - read_float(data)?, - read_float(data)?, - read_float(data)?, - )), - PixelBenderTypeOpcode::TFloat2x2 => Ok(PixelBenderType::TFloat2x2([ - read_float(data)?, - read_float(data)?, - read_float(data)?, - read_float(data)?, - ])), - PixelBenderTypeOpcode::TFloat3x3 => { - let mut floats: [f32; 9] = [0.0; 9]; - for float in &mut floats { - *float = read_float(data)?; - } - Ok(PixelBenderType::TFloat3x3(floats)) - } - PixelBenderTypeOpcode::TFloat4x4 => { - let mut floats: [f32; 16] = [0.0; 16]; - for float in &mut floats { - *float = read_float(data)?; - } - Ok(PixelBenderType::TFloat4x4(floats)) - } - PixelBenderTypeOpcode::TInt => Ok(PixelBenderType::TInt(data.read_i16::()?)), - PixelBenderTypeOpcode::TInt2 => Ok(PixelBenderType::TInt2( - data.read_i16::()?, - data.read_i16::()?, - )), - PixelBenderTypeOpcode::TInt3 => Ok(PixelBenderType::TInt3( - data.read_i16::()?, - data.read_i16::()?, - data.read_i16::()?, - )), - PixelBenderTypeOpcode::TInt4 => Ok(PixelBenderType::TInt4( - data.read_i16::()?, - data.read_i16::()?, - data.read_i16::()?, - data.read_i16::()?, - )), - PixelBenderTypeOpcode::TString => Ok(PixelBenderType::TString(read_string(data)?)), - } -} - -fn read_uint24(data: &mut R) -> Result> { - let mut src = data.read_u16::()? as u32; - src += data.read_u8()? as u32; - Ok(src) -} - -// The opcodes are laid out like this: -// -// ``` -// PBJMeta1 (for overall program) -// PBJMeta1 (for overall program) -// PBJParam (param 1) -// ... -// PBJMeta1 (for param 1) -// PBJMeta1 (for param 1) -// ... -// PBJParam (param 2) -// ,,, -// PBJMeta2 (for param 2) -// ``` -// -// The metadata associated with parameter is determined by all of the metadata opcodes -// that come after it and before the next parameter opcode. The metadata opcodes -// that come before all params are associated with the overall program. - -fn apply_metadata(shader: &mut PixelBenderShader, metadata: &mut Vec) { - // Reset the accumulated metadata Vec - we will start accumulating metadata for the next param - let metadata = std::mem::take(metadata); - if shader.params.is_empty() { - shader.metadata = metadata; - } else { - match shader.params.last_mut().unwrap() { - PixelBenderParam::Normal { metadata: meta, .. } => { - *meta = metadata; - } - param => { - if !metadata.is_empty() { - panic!("Tried to apply metadata to texture parameter {param:?}") - } - } - } - } -} diff --git a/desktop/src/main.rs b/desktop/src/main.rs index 81b7b24c2..eab3f676e 100644 --- a/desktop/src/main.rs +++ b/desktop/src/main.rs @@ -67,6 +67,7 @@ fn init() { let subscriber = tracing_subscriber::fmt::Subscriber::builder() .with_env_filter(tracing_subscriber::EnvFilter::from_default_env()) .finish(); + #[cfg(feature = "tracy")] let subscriber = { use tracing_subscriber::layer::SubscriberExt; diff --git a/render/Cargo.toml b/render/Cargo.toml index 8676482f9..a5d2216e3 100644 --- a/render/Cargo.toml +++ b/render/Cargo.toml @@ -25,6 +25,9 @@ serde = { version = "1.0.164", features = ["derive"] } clap = { version = "4.3.3", features = ["derive"], optional = true } h263-rs-yuv = { git = "https://github.com/ruffle-rs/h263-rs", rev = "d5d78eb251c1ce1f1da57c63db14f0fdc77a4b36"} lru = "0.10.0" +num-traits = "0.2" +num-derive = "0.3" +byteorder = "1.4" [dependencies.jpeg-decoder] version = "0.3.0" diff --git a/render/canvas/src/lib.rs b/render/canvas/src/lib.rs index 7463051ad..412af30a5 100644 --- a/render/canvas/src/lib.rs +++ b/render/canvas/src/lib.rs @@ -487,6 +487,22 @@ impl RenderBackend for WebCanvasRenderBackend { } fn set_quality(&mut self, _quality: StageQuality) {} + + fn compile_pixelbender_shader( + &mut self, + _shader: ruffle_render::pixel_bender::PixelBenderShader, + ) -> Result { + Err(Error::Unimplemented("compile_pixelbender_shader".into())) + } + + fn run_pixelbender_shader( + &mut self, + _handle: ruffle_render::pixel_bender::PixelBenderShaderHandle, + _arguments: &[ruffle_render::pixel_bender::PixelBenderShaderArgument], + _target: BitmapHandle, + ) -> Result, Error> { + Err(Error::Unimplemented("run_pixelbender_shader".into())) + } } impl CommandHandler for WebCanvasRenderBackend { diff --git a/render/naga-agal/Cargo.toml b/render/naga-agal/Cargo.toml index c4976a7c6..57e6acc7f 100644 --- a/render/naga-agal/Cargo.toml +++ b/render/naga-agal/Cargo.toml @@ -9,10 +9,10 @@ version.workspace = true [dependencies] bitflags = "2.3.1" -naga = "0.12.2" +naga = { workspace = true } num-derive = "0.3.3" num-traits = "0.2.15" [dev-dependencies] insta = "1.29.0" -naga = { version = "0.12.2", features = ["wgsl-out", "validate"] } +naga = { workspace = true, features = ["wgsl-out", "validate"] } diff --git a/render/naga-pixelbender/Cargo.toml b/render/naga-pixelbender/Cargo.toml new file mode 100644 index 000000000..7550f3e4c --- /dev/null +++ b/render/naga-pixelbender/Cargo.toml @@ -0,0 +1,15 @@ +[package] +name = "naga-pixelbender" +authors.workspace = true +edition.workspace = true +homepage.workspace = true +license.workspace = true +repository.workspace = true +version.workspace = true + +[dependencies] +ruffle_render = { path = "../" } +naga = { workspace = true } +anyhow = "1.0.71" +bitflags = "2.3.1" + diff --git a/render/naga-pixelbender/src/lib.rs b/render/naga-pixelbender/src/lib.rs new file mode 100644 index 000000000..9f4d198e2 --- /dev/null +++ b/render/naga-pixelbender/src/lib.rs @@ -0,0 +1,1077 @@ +use std::sync::OnceLock; + +use anyhow::Result; +use naga::{ + valid::{Capabilities, ValidationFlags, Validator}, + AddressSpace, ArraySize, BinaryOperator, Binding, Block, BuiltIn, Constant, ConstantInner, + EntryPoint, Expression, Function, FunctionArgument, FunctionResult, GlobalVariable, Handle, + ImageClass, ImageDimension, ImageQuery, LocalVariable, MathFunction, Module, ResourceBinding, + ScalarKind, ScalarValue, ShaderStage, Span, Statement, SwizzleComponent, Type, TypeInner, +}; +use ruffle_render::pixel_bender::{ + Opcode, Operation, PixelBenderParam, PixelBenderParamQualifier, PixelBenderReg, + PixelBenderRegChannel, PixelBenderRegKind, PixelBenderShader, PixelBenderTypeOpcode, + OUT_COORD_NAME, +}; + +/// The entrypoint name for the vertex and fragment shaders. +pub const SHADER_ENTRYPOINT: &str = "main"; + +pub struct NagaModules { + pub vertex: naga::Module, + pub fragment: naga::Module, + + pub float_parameters_buffer_size: u64, + pub int_parameters_buffer_size: u64, +} + +pub struct ShaderBuilder<'a> { + module: Module, + func: Function, + shader: &'a PixelBenderShader, + + vec2f: Handle, + vec4f: Handle, + vec4i: Handle, + image2d: Handle, + sampler: Handle, + + clamp_nearest: Handle, + clamp_linear: Handle, + // FIXME - implement the corresponding opcode 'Sample' + #[allow(dead_code)] + clamp_bilinear: Handle, + + textures: Vec>>, + + // Whenever we read from a particular register + // for the first time, we create a new local variable + // and store an expression here. All registers are of type vec4f + // for simplicity. When we write to a destination register, we only + // update the components specified in the destination write mask + float_registers: Vec>>, + + /// Like float_registesr but with vec4i + int_registers: Vec>>, + + // A stack of if/else blocks, using to push statements + // into the correct block. + blocks: Vec, +} + +/// Handles 'if' and 'else' blocks in PixelBender bytecode. +/// When we encounter an 'OpIf' opcode, we push an `IfElse` entry onto the block stack. +/// Any subsequent opcodes will be added to the `after_if` block. +/// When we encounter an 'OpElse' opcode, we switch to adding opcodes to the `after_else` block +/// by setting `in_after_if` to false. +/// When we encouter an `OpEndIf` opcode, we pop our `IfElse` entry from the stack, and emit +/// a `Statement::If` with the `after_if` and `after_else` blocks. +#[derive(Debug)] +enum BlockStackEntry { + Normal(Block), + IfElse { + after_if: Block, + after_else: Block, + in_after_if: bool, + condition: Handle, + }, +} + +const TEXTURE_SAMPLER_START_BIND_INDEX: u32 = 0; + +// FIXME - this shouldn't actually be clamp - https://www.mcjones.org/paul/PixelBenderReference.pdf +// says that coordinates outside the range are 'transparent black' +pub const SAMPLER_CLAMP_NEAREST: u32 = 0; +pub const SAMPLER_CLAMP_LINEAR: u32 = 1; +pub const SAMPLER_CLAMP_BILINEAR: u32 = 2; + +pub const SHADER_FLOAT_PARAMETERS_INDEX: u32 = 3; +// This covers ints and bool parameters +pub const SHADER_INT_PARAMETERS_INDEX: u32 = 4; + +pub const TEXTURE_START_BIND_INDEX: u32 = 5; + +const VERTEX_SHADER_WGSL: &str = r#" +struct VertexOutput { + @builtin(position) position: vec4, +}; + +@vertex +fn main( + @location(0) position: vec2, +) -> VertexOutput { + // Map coordinates from [0, 1] to [-1, 1] + return VertexOutput(vec4((position * vec2(2.0, 2.0)) - vec2(1.0, 1.0), 0.0, 1.0)); +}; +"#; + +impl<'a> ShaderBuilder<'a> { + pub fn build(shader: &PixelBenderShader) -> Result { + let mut module = Module::default(); + + static VERTEX_SHADER: OnceLock = OnceLock::new(); + let vertex_shader = VERTEX_SHADER + .get_or_init(|| { + naga::front::wgsl::Frontend::new() + .parse(VERTEX_SHADER_WGSL) + .expect("Failed to parse vertex shader") + }) + .clone(); + + let vec2f = module.types.insert( + Type { + name: None, + inner: TypeInner::Vector { + size: naga::VectorSize::Bi, + kind: ScalarKind::Float, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let vec4f = module.types.insert( + Type { + name: None, + inner: TypeInner::Vector { + size: naga::VectorSize::Quad, + kind: ScalarKind::Float, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let vec4i = module.types.insert( + Type { + name: None, + inner: TypeInner::Vector { + size: naga::VectorSize::Quad, + kind: ScalarKind::Sint, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let image2d = module.types.insert( + Type { + name: None, + inner: TypeInner::Image { + dim: ImageDimension::D2, + arrayed: false, + class: ImageClass::Sampled { + kind: ScalarKind::Float, + multi: false, + }, + }, + }, + Span::UNDEFINED, + ); + + let sampler = module.types.insert( + Type { + name: None, + inner: TypeInner::Sampler { comparison: false }, + }, + Span::UNDEFINED, + ); + + let mut func = Function::default(); + func.arguments.push(FunctionArgument { + name: None, + ty: vec4f, + binding: Some(Binding::BuiltIn(BuiltIn::Position { invariant: false })), + }); + + func.result = Some(FunctionResult { + ty: vec4f, + binding: Some(Binding::Location { + location: 0, + interpolation: None, + sampling: None, + }), + }); + + let samplers = (0..3) + .map(|i| { + let var = module.global_variables.append( + GlobalVariable { + name: Some(format!("sampler{}", i)), + space: naga::AddressSpace::Handle, + binding: Some(naga::ResourceBinding { + group: 0, + binding: TEXTURE_SAMPLER_START_BIND_INDEX + i, + }), + ty: module.types.insert( + Type { + name: None, + inner: TypeInner::Sampler { comparison: false }, + }, + Span::UNDEFINED, + ), + init: None, + }, + Span::UNDEFINED, + ); + func.expressions + .append(Expression::GlobalVariable(var), Span::UNDEFINED) + }) + .collect::>(); + + let mut builder = ShaderBuilder { + module, + func, + vec2f, + vec4f, + vec4i, + image2d, + sampler, + clamp_nearest: samplers[SAMPLER_CLAMP_NEAREST as usize], + clamp_linear: samplers[SAMPLER_CLAMP_LINEAR as usize], + clamp_bilinear: samplers[SAMPLER_CLAMP_BILINEAR as usize], + + shader, + textures: Vec::new(), + float_registers: Vec::new(), + int_registers: Vec::new(), + blocks: vec![BlockStackEntry::Normal(Block::new())], + }; + + let wrapper_func = builder.make_sampler_wrapper(); + + let (float_parameters_buffer_size, int_parameters_buffer_size) = builder.add_arguments()?; + builder.process_opcodes(wrapper_func)?; + + let dst = shader + .params + .iter() + .find_map(|p| { + if let PixelBenderParam::Normal { + qualifier: PixelBenderParamQualifier::Output, + reg, + .. + } = p + { + Some(reg) + } else { + None + } + }) + .expect("Missing destination register!"); + assert_eq!( + dst.channels, + PixelBenderRegChannel::RGBA, + "Invalid 'dest' parameter register {dst:?}" + ); + + // We've emitted all of the opcodes into the function body, so we can now load + // from the destination register and return it from the function. + let dst_load = builder.load_src_register(dst)?; + builder.push_statement(Statement::Return { + value: Some(dst_load), + }); + + let block = match builder.blocks.pop().unwrap() { + BlockStackEntry::Normal(block) => block, + block => panic!("Unfinished if statement: {:?}", block), + }; + + if !builder.blocks.is_empty() { + panic!("Unbalanced blocks: {:?}", builder.blocks); + } + if !builder.func.body.is_empty() { + panic!( + "Incorrectly wrote to function body: {:?}", + builder.func.body + ); + } + builder.func.body = block; + + builder.module.entry_points.push(EntryPoint { + name: "main".to_string(), + stage: ShaderStage::Fragment, + early_depth_test: None, + workgroup_size: [0; 3], + function: builder.func, + }); + + Ok(NagaModules { + vertex: vertex_shader, + fragment: builder.module, + float_parameters_buffer_size, + int_parameters_buffer_size, + }) + } + + fn add_arguments(&mut self) -> Result<(u64, u64)> { + let mut num_floats = 0; + let mut num_ints = 0; + + let mut param_offsets = Vec::new(); + + let mut out_coord = None; + + for param in &self.shader.params { + match param { + PixelBenderParam::Normal { + qualifier: PixelBenderParamQualifier::Input, + param_type, + reg, + name, + metadata: _, + } => { + if name == OUT_COORD_NAME { + // This is passed in through a builtin, not a uniform + out_coord = Some(reg); + continue; + } + + let float_offset = num_floats; + let int_offset = num_ints; + + // To meet alignment requirements, each parameter is stored as a vec4 in the constants array. + // Smaller types (e.g. Float, Float2, Float3) are padded with zeros. + let (offset, is_float) = match param_type { + PixelBenderTypeOpcode::TFloat + | PixelBenderTypeOpcode::TFloat2 + | PixelBenderTypeOpcode::TFloat3 + | PixelBenderTypeOpcode::TFloat4 => { + num_floats += 1; + (float_offset, true) + } + PixelBenderTypeOpcode::TInt + | PixelBenderTypeOpcode::TInt2 + | PixelBenderTypeOpcode::TInt3 + | PixelBenderTypeOpcode::TInt4 => { + num_ints += 1; + (int_offset, false) + } + PixelBenderTypeOpcode::TString => continue, + _ => unimplemented!("Unsupported parameter type {:?}", param_type), + }; + + param_offsets.push((reg, offset, is_float)); + } + PixelBenderParam::Texture { + index, + channels: _, + name: _, + } => { + let index = *index as usize; + let global_var = self.module.global_variables.append( + GlobalVariable { + name: Some(format!("texture{}", index)), + space: AddressSpace::Handle, + binding: Some(ResourceBinding { + group: 0, + binding: TEXTURE_START_BIND_INDEX + index as u32, + }), + ty: self.image2d, + init: None, + }, + Span::UNDEFINED, + ); + + if index >= self.textures.len() { + self.textures.resize(index + 1, None); + } + self.textures[index] = Some( + self.func + .expressions + .append(Expression::GlobalVariable(global_var), Span::UNDEFINED), + ); + } + _ => {} + } + } + + // These globals must have at least one entry in the array to satisfy naga, + // even if we don't have any parameters of that type. + + let num_floats_constant = self.module.constants.append( + Constant { + name: None, + specialization: None, + inner: naga::ConstantInner::Scalar { + width: 4, + value: naga::ScalarValue::Uint(num_floats.max(1) as u64), + }, + }, + Span::UNDEFINED, + ); + + let num_ints_constant = self.module.constants.append( + Constant { + name: None, + specialization: None, + inner: naga::ConstantInner::Scalar { + width: 4, + value: naga::ScalarValue::Uint(num_ints.max(1) as u64), + }, + }, + Span::UNDEFINED, + ); + + let shader_float_parameters = self.module.global_variables.append( + GlobalVariable { + name: Some("shader_float_parameters".to_string()), + space: naga::AddressSpace::Uniform, + binding: Some(naga::ResourceBinding { + group: 0, + binding: SHADER_FLOAT_PARAMETERS_INDEX, + }), + ty: self.module.types.insert( + Type { + name: None, + inner: TypeInner::Array { + base: self.vec4f, + size: ArraySize::Constant(num_floats_constant), + stride: std::mem::size_of::() as u32 * 4, + }, + }, + Span::UNDEFINED, + ), + init: None, + }, + Span::UNDEFINED, + ); + + let shader_int_parameters = self.module.global_variables.append( + GlobalVariable { + name: Some("shader_int_parameters".to_string()), + space: naga::AddressSpace::Uniform, + binding: Some(naga::ResourceBinding { + group: 0, + binding: SHADER_INT_PARAMETERS_INDEX, + }), + ty: self.module.types.insert( + Type { + name: None, + inner: TypeInner::Array { + base: self.vec4i, + size: ArraySize::Constant(num_ints_constant), + stride: std::mem::size_of::() as u32 * 4, + }, + }, + Span::UNDEFINED, + ), + init: None, + }, + Span::UNDEFINED, + ); + + for (reg, offset, is_float) in param_offsets { + let global = if is_float { + shader_float_parameters + } else { + shader_int_parameters + }; + + let global_base = self + .func + .expressions + .append(Expression::GlobalVariable(global), Span::UNDEFINED); + + let src_ptr = self.evaluate_expr(Expression::AccessIndex { + base: global_base, + index: offset, + }); + + let src = self.evaluate_expr(Expression::Load { pointer: src_ptr }); + + self.emit_dest_store(src, reg); + } + + // Emit this after all other registers have been initialized + // (it may use te same register as another parameter, but with different components) + + if let Some(coord_reg) = out_coord { + let coord_val = self + .func + .expressions + .append(Expression::FunctionArgument(0), Span::UNDEFINED); + self.emit_dest_store(coord_val, coord_reg); + } + + Ok(( + num_floats.max(1) as u64 * 4 * std::mem::size_of::() as u64, + num_ints.max(1) as u64 * 4 * std::mem::size_of::() as u64, + )) + } + + // Works around wgpu requiring naga's strict level of uniformity analysis + // See https://github.com/gpuweb/gpuweb/issues/3479#issuecomment-1519140312 + fn make_sampler_wrapper(&mut self) -> Handle { + let mut func = Function { + name: Some("sampler_wrapper".to_string()), + arguments: vec![ + FunctionArgument { + name: Some("image".to_string()), + ty: self.image2d, + binding: None, + }, + FunctionArgument { + name: Some("sampler".to_string()), + ty: self.sampler, + binding: None, + }, + FunctionArgument { + name: Some("coord".to_string()), + ty: self.vec2f, + binding: None, + }, + ], + result: Some(FunctionResult { + ty: self.vec4f, + binding: None, + }), + ..Default::default() + }; + + let image = func + .expressions + .append(Expression::FunctionArgument(0), Span::UNDEFINED); + let sampler = func + .expressions + .append(Expression::FunctionArgument(1), Span::UNDEFINED); + let coordinate = func + .expressions + .append(Expression::FunctionArgument(2), Span::UNDEFINED); + + let sample = func.expressions.append( + Expression::ImageSample { + image, + sampler, + coordinate, + array_index: None, + offset: None, + level: naga::SampleLevel::Auto, + depth_ref: None, + gather: None, + }, + Span::UNDEFINED, + ); + + func.body.push( + Statement::Emit(func.expressions.range_from(func.expressions.len() - 1)), + Span::UNDEFINED, + ); + + func.body.push( + Statement::Return { + value: Some(sample), + }, + Span::UNDEFINED, + ); + self.module.functions.append(func, Span::UNDEFINED) + } + + fn process_opcodes(&mut self, sample_wrapper_func: Handle) -> Result<()> { + for op in &self.shader.operations { + match op { + Operation::Normal { opcode, dst, src } => { + let src = self.load_src_register(src)?; + let mut dst = dst.clone(); + let evaluated = match opcode { + Opcode::Mov => src, + Opcode::Rcp => { + let const_one = self.module.constants.append( + Constant { + name: None, + specialization: None, + inner: naga::ConstantInner::Scalar { + width: 4, + value: naga::ScalarValue::Float(1.0), + }, + }, + Span::UNDEFINED, + ); + let expr_one = self + .func + .expressions + .append(Expression::Constant(const_one), Span::UNDEFINED); + + let vec_one = self.evaluate_expr(Expression::Splat { + size: naga::VectorSize::Quad, + value: expr_one, + }); + + // Perform 'vec4(1.0, 1.0, 1.0. 1.0) / src' + self.evaluate_expr(Expression::Binary { + op: BinaryOperator::Divide, + left: vec_one, + right: src, + }) + } + Opcode::Sub | Opcode::Add | Opcode::Mul => { + // The destiation is also used as the first operand: 'dst = dst src' + let left = self.load_src_register(&dst)?; + + let op = match opcode { + Opcode::Sub => BinaryOperator::Subtract, + Opcode::Add => BinaryOperator::Add, + Opcode::Mul => BinaryOperator::Multiply, + _ => unreachable!(), + }; + + self.evaluate_expr(Expression::Binary { + op, + left, + right: src, + }) + } + Opcode::LessThan => { + let left = self.load_src_register(&dst)?; + let res = self.evaluate_expr(Expression::Binary { + op: BinaryOperator::Less, + left, + right: src, + }); + + // The 'LessThan' opcodes appears to compare the src and dst, and then + // write the result to the 'R' component of int register 0 + dst = PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::R], + kind: PixelBenderRegKind::Int, + }; + // We get back a vec of bools from BinaryOperator::Less, so convert it to a vec of floats + self.evaluate_expr(Expression::As { + expr: res, + kind: ScalarKind::Float, + convert: Some(4), + }) + } + Opcode::LogicalOr => { + // The destiation is also used as the first operand: 'dst = dst - src' + let left = self.load_src_register(&dst)?; + let left_bool = self.evaluate_expr(Expression::As { + expr: left, + kind: ScalarKind::Bool, + convert: Some(1), + }); + let right_bool = self.evaluate_expr(Expression::As { + expr: src, + kind: ScalarKind::Bool, + convert: Some(1), + }); + + // Note - this should just be a `LogicalOr` between two vectors. + // However, Naga currently handles this incorrectly - see https://github.com/gfx-rs/naga/issues/1931 + // For now, work around this by manually applying it component-wise. + + let source_components: Vec<_> = (0..4) + .map(|index| { + self.evaluate_expr(Expression::AccessIndex { + base: left_bool, + index, + }) + }) + .collect(); + + let dest_components: Vec<_> = (0..4) + .map(|index| { + self.evaluate_expr(Expression::AccessIndex { + base: right_bool, + index, + }) + }) + .collect(); + + let res_components = (0..4) + .map(|index| { + let component_or = self.evaluate_expr(Expression::Binary { + op: BinaryOperator::LogicalOr, + left: source_components[index], + right: dest_components[index], + }); + + // We get back a bool from BinaryOperator::LogicalOr, so convert it to a float + self.evaluate_expr(Expression::As { + expr: component_or, + kind: ScalarKind::Float, + convert: Some(4), + }) + }) + .collect(); + + self.evaluate_expr(Expression::Compose { + ty: self.vec4f, + components: res_components, + }) + } + Opcode::Floor => self.evaluate_expr(Expression::Math { + fun: MathFunction::Floor, + arg: src, + arg1: None, + arg2: None, + arg3: None, + }), + Opcode::Length => { + let length = self.evaluate_expr(Expression::Math { + fun: MathFunction::Length, + arg: src, + arg1: None, + arg2: None, + arg3: None, + }); + self.evaluate_expr(Expression::Splat { + size: naga::VectorSize::Quad, + value: length, + }) + } + _ => { + unimplemented!("Unimplemented opcode {opcode:?}"); + } + }; + self.emit_dest_store(evaluated, &dst); + } + Operation::SampleLinear { dst, src, tf } + | Operation::SampleNearest { dst, src, tf } => { + let mut coord = self.load_src_register(src)?; + coord = self.evaluate_expr(Expression::Swizzle { + size: naga::VectorSize::Bi, + vector: coord, + // Only the first two components matter + pattern: [ + SwizzleComponent::X, + SwizzleComponent::Y, + SwizzleComponent::W, + SwizzleComponent::W, + ], + }); + + let size_vec = self.evaluate_expr(Expression::ImageQuery { + image: self.textures[*tf as usize].unwrap(), + query: ImageQuery::Size { level: None }, + }); + + let size_vec_float = self.evaluate_expr(Expression::As { + kind: crate::ScalarKind::Float, + expr: size_vec, + convert: Some(4), + }); + + let normalized_coord = self.evaluate_expr(Expression::Binary { + op: BinaryOperator::Divide, + left: coord, + right: size_vec_float, + }); + + let image = self.textures[*tf as usize].unwrap(); + + let sampler = match op { + Operation::SampleNearest { .. } => self.clamp_nearest, + Operation::SampleLinear { .. } => self.clamp_linear, + _ => unreachable!(), + }; + + // Don't evaluate this expression - it gets evaluated by Statement::Call + let result = self + .func + .expressions + .append(Expression::CallResult(sample_wrapper_func), Span::UNDEFINED); + + // Call our helper function, which just calls 'Expression::ImageSample' with + // the provided parameters. This works around a uniformity analysis issue + // with wgpu/naga + self.push_statement(Statement::Call { + function: sample_wrapper_func, + arguments: vec![image, sampler, normalized_coord], + result: Some(result), + }); + + self.emit_dest_store(result, dst); + } + Operation::LoadFloat { dst, val } => { + let const_val = self.module.constants.append( + crate::Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: ScalarValue::Float(*val as f64), + }, + }, + Span::UNDEFINED, + ); + let const_expr = self + .func + .expressions + .append(Expression::Constant(const_val), Span::UNDEFINED); + + let const_vec = self.evaluate_expr(Expression::Splat { + size: naga::VectorSize::Quad, + value: const_expr, + }); + self.emit_dest_store(const_vec, dst); + } + Operation::LoadInt { dst, val } => { + let const_val = self.module.constants.append( + crate::Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: ScalarValue::Sint(*val as i64), + }, + }, + Span::UNDEFINED, + ); + let const_expr = self + .func + .expressions + .append(Expression::Constant(const_val), Span::UNDEFINED); + + let const_vec = self.evaluate_expr(Expression::Splat { + size: naga::VectorSize::Quad, + value: const_expr, + }); + self.emit_dest_store(const_vec, dst); + } + Operation::If { src } => { + let scalar_zero = match src.kind { + PixelBenderRegKind::Float => ScalarValue::Float(0.0), + PixelBenderRegKind::Int => ScalarValue::Sint(0), + }; + if src.channels.len() != 1 { + panic!("If condition must be a scalar: {src:?}"); + } + + // FIXME - `load_src_register` always gives us a vec4 - ideally, we would + // have a flag to avoid this pointless splat-and-extract. + let src = self.load_src_register(src)?; + let first_component = self.evaluate_expr(Expression::AccessIndex { + base: src, + index: 0, + }); + + let const_zero = self.module.constants.append( + Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: scalar_zero, + }, + }, + Span::UNDEFINED, + ); + + let expr_zero = self + .func + .expressions + .append(Expression::Constant(const_zero), Span::UNDEFINED); + + let is_true = self.evaluate_expr(Expression::Binary { + op: BinaryOperator::NotEqual, + left: first_component, + right: expr_zero, + }); + + self.blocks.push(BlockStackEntry::IfElse { + after_if: Block::new(), + after_else: Block::new(), + in_after_if: true, + condition: is_true, + }) + } + Operation::Else => { + if let BlockStackEntry::IfElse { + after_if: _, + after_else: _, + in_after_if, + condition: _, + } = self.blocks.last_mut().unwrap() + { + if !*in_after_if { + panic!("Multiple' els' opcodes for single 'if' opcode"); + } + *in_after_if = false; + } else { + unreachable!() + } + } + Operation::EndIf => { + let block = self.blocks.pop().unwrap(); + + match block { + BlockStackEntry::IfElse { + after_if, + after_else, + in_after_if: _, + condition, + } => { + self.push_statement(Statement::If { + condition, + // The opcodes occurig directly after the 'if' opcode + // get run if the condition is true + accept: after_if, + // The opcodes occurring directly after the 'els' opcode + // get run if the condition is false + reject: after_else, + }); + } + BlockStackEntry::Normal(block) => { + panic!("Eif opcode without matching 'if': {:?}", block) + } + } + } + _ => unimplemented!("Operation {op:?} not yet implemented"), + } + } + Ok(()) + } + + /// Gets a pointer to the given register - this does *not* perform a load, so it can + /// be used with both `Expression::Load` and `Statement::Store` + fn register_pointer(&mut self, reg: &PixelBenderReg) -> Result> { + let index = reg.index as usize; + + let (ty, registers, register_kind_name) = match reg.kind { + PixelBenderRegKind::Float => (self.vec4f, &mut self.float_registers, "float"), + PixelBenderRegKind::Int => (self.vec4i, &mut self.int_registers, "int"), + }; + + if index >= registers.len() { + registers.resize(index + 1, None); + } + + if registers[index].is_none() { + let local = self.func.local_variables.append( + LocalVariable { + name: Some(format!("local_{register_kind_name}_reg_{index}")), + ty, + init: None, + }, + Span::UNDEFINED, + ); + + let expr = self + .func + .expressions + .append(Expression::LocalVariable(local), Span::UNDEFINED); + registers[index] = Some(expr); + } + Ok(registers[index].unwrap()) + } + + /// Loads a vec4f/vec4i from the given register. Note that all registers are 4-component + /// vectors - if the `PixelBenderReg` requests fewer components then that, then the extra + /// components will be meaningless. This greatly simplifies the code, since we don't need + /// to track whether or not we have a scalar or a vector everywhere. + fn load_src_register(&mut self, reg: &PixelBenderReg) -> Result> { + let reg_ptr = self.register_pointer(reg)?; + let reg_value = self.evaluate_expr(Expression::Load { pointer: reg_ptr }); + + let mut swizzle_components = reg + .channels + .iter() + .map(|c| match c { + PixelBenderRegChannel::R => SwizzleComponent::X, + PixelBenderRegChannel::G => SwizzleComponent::Y, + PixelBenderRegChannel::B => SwizzleComponent::Z, + PixelBenderRegChannel::A => SwizzleComponent::W, + }) + .collect::>(); + + if swizzle_components.len() < 4 { + // Pad with W - these components will be ignored, since whatever uses + // the result will only use the components corresponding to 'reg.channels' + swizzle_components.resize(4, SwizzleComponent::W); + } + + Ok(self.evaluate_expr(Expression::Swizzle { + size: naga::VectorSize::Quad, + vector: reg_value, + pattern: swizzle_components.try_into().unwrap(), + })) + } + + /// Creates a `Statement::Emit` covering `expr` + fn evaluate_expr(&mut self, expr: Expression) -> Handle { + let prev_len = self.func.expressions.len(); + let expr = self.func.expressions.append(expr, Span::UNDEFINED); + let range = self.func.expressions.range_from(prev_len); + self.push_statement(Statement::Emit(range)); + expr + } + + // Emits a store of `expr` to the destination register, taking into account the store mask. + fn emit_dest_store(&mut self, expr: Handle, dst: &PixelBenderReg) { + let dst_register = self.register_pointer(dst).unwrap(); + for (dst_channel, src_channel) in + dst.channels.iter().zip(PixelBenderRegChannel::RGBA.iter()) + { + // Write each channel of the source to the channel specified by the destiation mask + let src_component_index = *src_channel as u32; + let dst_component_index = *dst_channel as u32; + let src_component = self.evaluate_expr(Expression::AccessIndex { + base: expr, + index: src_component_index, + }); + + let dst_component = self.evaluate_expr(Expression::AccessIndex { + base: dst_register, + index: dst_component_index, + }); + + let scalar_kind = match dst.kind { + PixelBenderRegKind::Float => ScalarKind::Float, + PixelBenderRegKind::Int => ScalarKind::Sint, + }; + + let src_cast = self.evaluate_expr(Expression::As { + kind: scalar_kind, + expr: src_component, + convert: Some(4), + }); + + self.push_statement(Statement::Store { + pointer: dst_component, + value: src_cast, + }) + } + } + + /// Pushes a statement, taking into account our current 'if' block. + /// Use this instead of `self.func.body.push` + fn push_statement(&mut self, stmt: Statement) { + let block = match self.blocks.last_mut().unwrap() { + BlockStackEntry::Normal(block) => block, + BlockStackEntry::IfElse { + after_if, + after_else, + in_after_if, + condition: _, + } => { + if *in_after_if { + after_if + } else { + after_else + } + } + }; + block.push(stmt, Span::UNDEFINED); + } +} + +#[allow(dead_code)] +fn to_wgsl(module: &naga::Module) -> String { + let mut out = String::new(); + + let mut validator = Validator::new( + ValidationFlags::all() - ValidationFlags::CONTROL_FLOW_UNIFORMITY, + Capabilities::all(), + ); + let module_info = validator + .validate(module) + .unwrap_or_else(|e| panic!("Validation failed: {:#?}", e)); + + let mut writer = + naga::back::wgsl::Writer::new(&mut out, naga::back::wgsl::WriterFlags::EXPLICIT_TYPES); + + writer.write(module, &module_info).expect("Writing failed"); + out +} diff --git a/render/src/backend.rs b/render/src/backend.rs index f03385672..9febca569 100644 --- a/render/src/backend.rs +++ b/render/src/backend.rs @@ -4,6 +4,7 @@ use crate::bitmap::{Bitmap, BitmapHandle, BitmapSource, PixelRegion, SyncHandle} use crate::commands::CommandList; use crate::error::Error; use crate::filters::Filter; +use crate::pixel_bender::{PixelBenderShader, PixelBenderShaderArgument, PixelBenderShaderHandle}; use crate::quality::StageQuality; use crate::shape_utils::DistilledShape; use downcast_rs::{impl_downcast, Downcast}; @@ -75,6 +76,18 @@ pub trait RenderBackend: Downcast { fn name(&self) -> &'static str; fn set_quality(&mut self, quality: StageQuality); + + fn compile_pixelbender_shader( + &mut self, + shader: PixelBenderShader, + ) -> Result; + + fn run_pixelbender_shader( + &mut self, + handle: PixelBenderShaderHandle, + arguments: &[PixelBenderShaderArgument], + target: BitmapHandle, + ) -> Result, Error>; } impl_downcast!(RenderBackend); diff --git a/render/src/backend/null.rs b/render/src/backend/null.rs index 629529084..7cbc1212f 100644 --- a/render/src/backend/null.rs +++ b/render/src/backend/null.rs @@ -7,6 +7,7 @@ use crate::bitmap::{ }; use crate::commands::CommandList; use crate::error::Error; +use crate::pixel_bender::{PixelBenderShader, PixelBenderShaderArgument, PixelBenderShaderHandle}; use crate::quality::StageQuality; use crate::shape_utils::DistilledShape; use swf::Color; @@ -98,4 +99,22 @@ impl RenderBackend for NullRenderer { } fn set_quality(&mut self, _quality: StageQuality) {} + + fn run_pixelbender_shader( + &mut self, + _shader: PixelBenderShaderHandle, + _arguments: &[PixelBenderShaderArgument], + _target: BitmapHandle, + ) -> Result, Error> { + Err(Error::Unimplemented("Pixel bender shader".into())) + } + + fn compile_pixelbender_shader( + &mut self, + _shader: PixelBenderShader, + ) -> Result { + Err(Error::Unimplemented( + "Pixel bender shader compilation".into(), + )) + } } diff --git a/render/src/lib.rs b/render/src/lib.rs index b695eec44..be6818b52 100644 --- a/render/src/lib.rs +++ b/render/src/lib.rs @@ -5,6 +5,7 @@ pub mod bitmap; pub mod error; pub mod filters; pub mod matrix; +pub mod pixel_bender; pub mod shape_utils; pub mod transform; pub mod utils; diff --git a/render/src/pixel_bender.rs b/render/src/pixel_bender.rs new file mode 100644 index 000000000..0bb296a26 --- /dev/null +++ b/render/src/pixel_bender.rs @@ -0,0 +1,636 @@ +//! Pixel bender bytecode parsing code. +//! This is heavily based on https://github.com/jamesward/pbjas and https://github.com/HaxeFoundation/format/tree/master/format/pbj + +#[cfg(test)] +mod tests; + +use byteorder::{BigEndian, LittleEndian, ReadBytesExt}; +use downcast_rs::{impl_downcast, Downcast}; +use gc_arena::Collect; +use num_traits::FromPrimitive; +use std::{ + fmt::{Debug, Display, Formatter}, + io::Read, + sync::Arc, +}; + +use crate::bitmap::BitmapHandle; + +/// The name of a special parameter, which gets automatically filled in with the coordinates +/// of the pixel being processed. +pub const OUT_COORD_NAME: &str = "_OutCoord"; + +#[derive(Clone, Debug, Collect)] +#[collect(require_static)] +pub struct PixelBenderShaderHandle(pub Arc); + +pub trait PixelBenderShaderImpl: Downcast + Debug { + fn parsed_shader(&self) -> &PixelBenderShader; +} +impl_downcast!(PixelBenderShaderImpl); + +#[repr(u8)] +#[derive(Debug, Clone, PartialEq)] +pub enum PixelBenderType { + TFloat(f32) = 0x1, + TFloat2(f32, f32) = 0x2, + TFloat3(f32, f32, f32) = 0x3, + TFloat4(f32, f32, f32, f32) = 0x4, + TFloat2x2([f32; 4]) = 0x5, + TFloat3x3([f32; 9]) = 0x6, + TFloat4x4([f32; 16]) = 0x7, + TInt(i16) = 0x8, + TInt2(i16, i16) = 0x9, + TInt3(i16, i16, i16) = 0xA, + TInt4(i16, i16, i16, i16) = 0xB, + TString(String) = 0xC, +} + +// FIXME - come up with a way to reduce duplication here +#[derive(num_derive::FromPrimitive, Debug, PartialEq, Clone, Copy)] +pub enum PixelBenderTypeOpcode { + TFloat = 0x1, + TFloat2 = 0x2, + TFloat3 = 0x3, + TFloat4 = 0x4, + TFloat2x2 = 0x5, + TFloat3x3 = 0x6, + TFloat4x4 = 0x7, + TInt = 0x8, + TInt2 = 0x9, + TInt3 = 0xA, + TInt4 = 0xB, + TString = 0xC, +} + +#[derive(Debug, PartialEq, Copy, Clone)] +pub enum PixelBenderRegChannel { + R = 0, + G = 1, + B = 2, + A = 3, +} + +impl PixelBenderRegChannel { + pub const RGBA: [PixelBenderRegChannel; 4] = [ + PixelBenderRegChannel::R, + PixelBenderRegChannel::G, + PixelBenderRegChannel::B, + PixelBenderRegChannel::A, + ]; +} + +#[derive(Debug, PartialEq, Clone)] +pub struct PixelBenderReg { + pub index: u32, + pub channels: Vec, + pub kind: PixelBenderRegKind, +} + +#[derive(Debug, PartialEq, Clone, Copy)] +pub enum PixelBenderRegKind { + Float, + Int, +} + +#[derive(num_derive::FromPrimitive, Debug, PartialEq, Clone, Copy)] +pub enum PixelBenderParamQualifier { + Input = 1, + Output = 2, +} + +impl Display for PixelBenderTypeOpcode { + fn fmt(&self, f: &mut Formatter<'_>) -> std::fmt::Result { + write!( + f, + "{}", + match self { + PixelBenderTypeOpcode::TFloat => "float", + PixelBenderTypeOpcode::TFloat2 => "float2", + PixelBenderTypeOpcode::TFloat3 => "float3", + PixelBenderTypeOpcode::TFloat4 => "float4", + PixelBenderTypeOpcode::TFloat2x2 => "matrix2x2", + PixelBenderTypeOpcode::TFloat3x3 => "matrix3x3", + PixelBenderTypeOpcode::TFloat4x4 => "matrix4x4", + PixelBenderTypeOpcode::TInt => "int", + PixelBenderTypeOpcode::TInt2 => "int2", + PixelBenderTypeOpcode::TInt3 => "int3", + PixelBenderTypeOpcode::TInt4 => "int4", + PixelBenderTypeOpcode::TString => "string", + } + ) + } +} + +#[derive(num_derive::FromPrimitive, Debug, PartialEq, Clone, Copy)] +pub enum Opcode { + Nop = 0x0, + Add = 0x1, + Sub = 0x2, + Mul = 0x3, + Rcp = 0x4, + Div = 0x5, + Atan2 = 0x6, + Pow = 0x7, + Mod = 0x8, + Min = 0x9, + Max = 0xA, + Step = 0xB, + Sin = 0xC, + Cos = 0xD, + Tan = 0xE, + Asin = 0xF, + Acos = 0x10, + Atan = 0x11, + Exp = 0x12, + Exp2 = 0x13, + Log = 0x14, + Log2 = 0x15, + Sqrt = 0x16, + RSqrt = 0x17, + Abs = 0x18, + Sign = 0x19, + Floor = 0x1A, + Ceil = 0x1B, + Fract = 0x1C, + Mov = 0x1D, + FloatToInt = 0x1E, + IntToFloat = 0x1F, + MatMatMul = 0x20, + VecMatMul = 0x21, + MatVecMul = 0x22, + Normalize = 0x23, + Length = 0x24, + Distance = 0x25, + DotProduct = 0x26, + CrossProduct = 0x27, + Equal = 0x28, + NotEqual = 0x29, + LessThan = 0x2A, + LessThanEqual = 0x2B, + LogicalNot = 0x2C, + LogicalAnd = 0x2D, + LogicalOr = 0x2E, + LogicalXor = 0x2F, + SampleNearest = 0x30, + SampleLinear = 0x31, + LoadIntOrFloat = 0x32, + Loop = 0x33, + If = 0x34, + Else = 0x35, + EndIf = 0x36, + FloatToBool = 0x37, + BoolToFloat = 0x38, + IntToBool = 0x39, + BoolToInt = 0x3A, + VectorEqual = 0x3B, + VectorNotEqual = 0x3C, + BoolAny = 0x3D, + BoolAll = 0x3E, + PBJMeta1 = 0xA0, + PBJParam = 0xA1, + PBJMeta2 = 0xA2, + PBJParamTexture = 0xA3, + Name = 0xA4, + Version = 0xA5, +} + +#[derive(Debug, PartialEq, Clone)] +pub enum Operation { + Nop, + Normal { + opcode: Opcode, + dst: PixelBenderReg, + src: PixelBenderReg, + }, + LoadInt { + dst: PixelBenderReg, + val: i32, + }, + LoadFloat { + dst: PixelBenderReg, + val: f32, + }, + If { + src: PixelBenderReg, + }, + SampleNearest { + dst: PixelBenderReg, + src: PixelBenderReg, + tf: u8, + }, + SampleLinear { + dst: PixelBenderReg, + src: PixelBenderReg, + tf: u8, + }, + Else, + EndIf, +} + +#[derive(Debug, Clone)] +pub enum PixelBenderShaderArgument { + ImageInput { + index: u8, + channels: u8, + name: String, + texture: BitmapHandle, + }, + ValueInput { + index: u8, + value: PixelBenderType, + }, +} + +#[derive(Debug, PartialEq, Clone)] +pub struct PixelBenderShader { + pub name: String, + pub version: i32, + pub params: Vec, + pub metadata: Vec, + pub operations: Vec, +} + +#[derive(Debug, PartialEq, Clone)] +pub enum PixelBenderParam { + Normal { + qualifier: PixelBenderParamQualifier, + param_type: PixelBenderTypeOpcode, + reg: PixelBenderReg, + name: String, + metadata: Vec, + }, + Texture { + index: u8, + channels: u8, + name: String, + }, +} + +#[derive(Debug, Clone, PartialEq)] +pub struct PixelBenderMetadata { + pub key: String, + pub value: PixelBenderType, +} + +/// Parses PixelBender bytecode +pub fn parse_shader(mut data: &[u8]) -> Result> { + let mut shader = PixelBenderShader { + name: String::new(), + version: 0, + params: Vec::new(), + metadata: Vec::new(), + operations: Vec::new(), + }; + let data = &mut data; + let mut metadata = Vec::new(); + while !data.is_empty() { + read_op(data, &mut shader, &mut metadata)?; + } + // Any metadata left in the vec is associated with our final parameter. + apply_metadata(&mut shader, &mut metadata); + Ok(shader) +} + +fn read_src_reg(val: u32, size: u8) -> Result> { + const CHANNELS: [PixelBenderRegChannel; 4] = [ + PixelBenderRegChannel::R, + PixelBenderRegChannel::G, + PixelBenderRegChannel::B, + PixelBenderRegChannel::A, + ]; + + let swizzle = val >> 16; + let mut channels = Vec::new(); + for i in 0..size { + channels.push(CHANNELS[(swizzle >> (6 - i * 2) & 3) as usize]) + } + + let kind = if val & 0x8000 != 0 { + PixelBenderRegKind::Int + } else { + PixelBenderRegKind::Float + }; + + Ok(PixelBenderReg { + // Mask off the 0x8000 bit + index: val & 0x7FFF, + channels, + kind, + }) +} + +fn read_dst_reg(val: u16, mask: u8) -> Result> { + let mut channels = Vec::new(); + if mask & 0x8 != 0 { + channels.push(PixelBenderRegChannel::R); + } + if mask & 0x4 != 0 { + channels.push(PixelBenderRegChannel::G); + } + if mask & 0x2 != 0 { + channels.push(PixelBenderRegChannel::B); + } + if mask & 0x1 != 0 { + channels.push(PixelBenderRegChannel::A); + } + + let kind = if val & 0x8000 != 0 { + PixelBenderRegKind::Int + } else { + PixelBenderRegKind::Float + }; + + Ok(PixelBenderReg { + // Mask off the 0x8000 bit + index: (val & 0x7FFF) as u32, + channels, + kind, + }) +} + +fn read_op( + data: &mut R, + shader: &mut PixelBenderShader, + metadata: &mut Vec, +) -> Result<(), Box> { + let raw = data.read_u8()?; + let opcode = Opcode::from_u8(raw).expect("Unknown opcode"); + match opcode { + Opcode::Nop => { + assert_eq!(data.read_u32::()?, 0); + assert_eq!(data.read_u16::()?, 0); + shader.operations.push(Operation::Nop); + } + Opcode::PBJMeta1 | Opcode::PBJMeta2 => { + let meta_type = data.read_u8()?; + let meta_key = read_string(data)?; + let meta_value = read_value( + data, + PixelBenderTypeOpcode::from_u8(meta_type) + .unwrap_or_else(|| panic!("Unexpected meta type {meta_type}")), + )?; + metadata.push(PixelBenderMetadata { + key: meta_key, + value: meta_value, + }); + } + Opcode::PBJParam => { + let qualifier = data.read_u8()?; + let param_type = data.read_u8()?; + let reg = data.read_u16::()?; + let mask = data.read_u8()?; + let name = read_string(data)?; + + let param_type = PixelBenderTypeOpcode::from_u8(param_type).unwrap_or_else(|| { + panic!("Unexpected param type {param_type}"); + }); + let qualifier = PixelBenderParamQualifier::from_u8(qualifier) + .unwrap_or_else(|| panic!("Unexpected param qualifier {qualifier:?}")); + apply_metadata(shader, metadata); + + match param_type { + PixelBenderTypeOpcode::TFloat2x2 + | PixelBenderTypeOpcode::TFloat3x3 + | PixelBenderTypeOpcode::TFloat4x4 => { + panic!("Unsupported param type {param_type:?}"); + } + _ => {} + } + + let dst_reg = read_dst_reg(reg, mask)?; + + shader.params.push(PixelBenderParam::Normal { + qualifier, + param_type, + reg: dst_reg, + name, + metadata: Vec::new(), + }) + } + Opcode::PBJParamTexture => { + let index = data.read_u8()?; + let channels = data.read_u8()?; + let name = read_string(data)?; + apply_metadata(shader, metadata); + + shader.params.push(PixelBenderParam::Texture { + index, + channels, + name, + }); + } + Opcode::Name => { + let len = data.read_u16::()?; + let mut string_bytes = vec![0; len as usize]; + data.read_exact(&mut string_bytes)?; + shader.name = String::from_utf8(string_bytes)?; + } + Opcode::Version => { + shader.version = data.read_i32::()?; + } + Opcode::If => { + assert_eq!(read_uint24(data)?, 0); + let src = read_uint24(data)?; + assert_eq!(data.read_u8()?, 0); + let src_reg = read_src_reg(src, 1)?; + shader.operations.push(Operation::If { src: src_reg }); + } + Opcode::Else => { + assert_eq!(data.read_u32::()?, 0); + assert_eq!(read_uint24(data)?, 0); + shader.operations.push(Operation::Else); + } + Opcode::EndIf => { + assert_eq!(data.read_u32::()?, 0); + assert_eq!(read_uint24(data)?, 0); + shader.operations.push(Operation::EndIf); + } + Opcode::LoadIntOrFloat => { + let dst = data.read_u16::()?; + let mask = data.read_u8()?; + assert_eq!(mask & 0xF, 0); + let dst_reg = read_dst_reg(dst, mask >> 4)?; + match dst_reg.kind { + PixelBenderRegKind::Float => { + let val = read_float(data)?; + shader + .operations + .push(Operation::LoadFloat { dst: dst_reg, val }) + } + PixelBenderRegKind::Int => { + let val = data.read_i32::()?; + shader + .operations + .push(Operation::LoadInt { dst: dst_reg, val }) + } + } + } + Opcode::SampleNearest | Opcode::SampleLinear => { + let dst = data.read_u16::()?; + let mask = data.read_u8()?; + let src = read_uint24(data)?; + let tf = data.read_u8()?; + + let dst_reg = read_dst_reg(dst, mask >> 4)?; + let src_reg = read_src_reg(src, 2)?; + + match opcode { + Opcode::SampleNearest => shader.operations.push(Operation::SampleNearest { + dst: dst_reg, + src: src_reg, + tf, + }), + Opcode::SampleLinear => shader.operations.push(Operation::SampleLinear { + dst: dst_reg, + src: src_reg, + tf, + }), + _ => unreachable!(), + } + } + _ => { + let dst = data.read_u16::()?; + let mut mask = data.read_u8()?; + let size = (mask & 0x3) + 1; + let matrix = (mask >> 2) & 3; + let src = read_uint24(data)?; + assert_eq!(data.read_u8()?, 0, "Unexpected u8 for opcode {opcode:?}"); + mask >>= 4; + + let src_reg = read_src_reg(src, size)?; + let dst_reg = if matrix != 0 { + assert_eq!(src >> 16, 0); + assert_eq!(size, 1); + panic!("Matrix with mask {mask:b} matrix {matrix:b}"); + } else { + read_dst_reg(dst, mask)? + }; + shader.operations.push(Operation::Normal { + opcode, + dst: dst_reg, + src: src_reg, + }) + } + }; + Ok(()) +} + +fn read_string(data: &mut R) -> Result> { + let mut string = String::new(); + let mut b = data.read_u8()?; + while b != 0 { + string.push(b as char); + b = data.read_u8()?; + } + Ok(string) +} + +fn read_float(data: &mut R) -> Result> { + Ok(data.read_f32::()?) +} + +fn read_value( + data: &mut R, + opcode: PixelBenderTypeOpcode, +) -> Result> { + match opcode { + PixelBenderTypeOpcode::TFloat => Ok(PixelBenderType::TFloat(read_float(data)?)), + PixelBenderTypeOpcode::TFloat2 => Ok(PixelBenderType::TFloat2( + read_float(data)?, + read_float(data)?, + )), + PixelBenderTypeOpcode::TFloat3 => Ok(PixelBenderType::TFloat3( + read_float(data)?, + read_float(data)?, + read_float(data)?, + )), + PixelBenderTypeOpcode::TFloat4 => Ok(PixelBenderType::TFloat4( + read_float(data)?, + read_float(data)?, + read_float(data)?, + read_float(data)?, + )), + PixelBenderTypeOpcode::TFloat2x2 => Ok(PixelBenderType::TFloat2x2([ + read_float(data)?, + read_float(data)?, + read_float(data)?, + read_float(data)?, + ])), + PixelBenderTypeOpcode::TFloat3x3 => { + let mut floats: [f32; 9] = [0.0; 9]; + for float in &mut floats { + *float = read_float(data)?; + } + Ok(PixelBenderType::TFloat3x3(floats)) + } + PixelBenderTypeOpcode::TFloat4x4 => { + let mut floats: [f32; 16] = [0.0; 16]; + for float in &mut floats { + *float = read_float(data)?; + } + Ok(PixelBenderType::TFloat4x4(floats)) + } + PixelBenderTypeOpcode::TInt => Ok(PixelBenderType::TInt(data.read_i16::()?)), + PixelBenderTypeOpcode::TInt2 => Ok(PixelBenderType::TInt2( + data.read_i16::()?, + data.read_i16::()?, + )), + PixelBenderTypeOpcode::TInt3 => Ok(PixelBenderType::TInt3( + data.read_i16::()?, + data.read_i16::()?, + data.read_i16::()?, + )), + PixelBenderTypeOpcode::TInt4 => Ok(PixelBenderType::TInt4( + data.read_i16::()?, + data.read_i16::()?, + data.read_i16::()?, + data.read_i16::()?, + )), + PixelBenderTypeOpcode::TString => Ok(PixelBenderType::TString(read_string(data)?)), + } +} + +fn read_uint24(data: &mut R) -> Result> { + let ch1 = data.read_u8()? as u32; + let ch2 = data.read_u8()? as u32; + let ch3 = data.read_u8()? as u32; + Ok(ch1 | (ch2 << 8) | (ch3 << 16)) +} + +// The opcodes are laid out like this: +// +// ``` +// PBJMeta1 (for overall program) +// PBJMeta1 (for overall program) +// PBJParam (param 1) +// ... +// PBJMeta1 (for param 1) +// PBJMeta1 (for param 1) +// ... +// PBJParam (param 2) +// ,,, +// PBJMeta2 (for param 2) +// ``` +// +// The metadata associated with parameter is determined by all of the metadata opcodes +// that come after it and before the next parameter opcode. The metadata opcodes +// that come before all params are associated with the overall program. + +fn apply_metadata(shader: &mut PixelBenderShader, metadata: &mut Vec) { + // Reset the accumulated metadata Vec - we will start accumulating metadata for the next param + let metadata = std::mem::take(metadata); + match shader.params.last_mut() { + Some(PixelBenderParam::Normal { metadata: meta, .. }) => { + *meta = metadata; + } + Some(param) => { + if !metadata.is_empty() { + panic!("Tried to apply metadata to texture parameter {param:?}") + } + } + None => { + shader.metadata = metadata; + } + } +} diff --git a/core/src/pixel_bender/tests.rs b/render/src/pixel_bender/tests.rs similarity index 51% rename from core/src/pixel_bender/tests.rs rename to render/src/pixel_bender/tests.rs index efaa33139..657ca94b5 100644 --- a/core/src/pixel_bender/tests.rs +++ b/render/src/pixel_bender/tests.rs @@ -1,6 +1,7 @@ use crate::pixel_bender::{ Opcode, Operation, PixelBenderMetadata, PixelBenderParam, PixelBenderParamQualifier, - PixelBenderShader, PixelBenderType, PixelBenderTypeOpcode, + PixelBenderReg, PixelBenderRegChannel, PixelBenderRegKind, PixelBenderShader, PixelBenderType, + PixelBenderTypeOpcode, }; use super::parse_shader; @@ -42,8 +43,11 @@ fn simple_shader() { PixelBenderParam::Normal { qualifier: PixelBenderParamQualifier::Input, param_type: PixelBenderTypeOpcode::TFloat2, - reg: 0, - mask: 12, + reg: PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::R, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, name: "_OutCoord".to_string(), metadata: vec![], }, @@ -55,16 +59,22 @@ fn simple_shader() { PixelBenderParam::Normal { qualifier: PixelBenderParamQualifier::Output, param_type: PixelBenderTypeOpcode::TFloat4, - reg: 1, - mask: 15, + reg: PixelBenderReg { + index: 1, + channels: PixelBenderRegChannel::RGBA.to_vec(), + kind: PixelBenderRegKind::Float, + }, name: "dst".to_string(), metadata: vec![], }, PixelBenderParam::Normal { qualifier: PixelBenderParamQualifier::Input, param_type: PixelBenderTypeOpcode::TFloat2, - reg: 0, - mask: 3, + reg: PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, name: "size".to_string(), metadata: vec![ PixelBenderMetadata { @@ -90,8 +100,11 @@ fn simple_shader() { PixelBenderParam::Normal { qualifier: PixelBenderParamQualifier::Input, param_type: PixelBenderTypeOpcode::TFloat, - reg: 2, - mask: 8, + reg: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::R], + kind: PixelBenderRegKind::Float, + }, name: "radius".to_string(), metadata: vec![ PixelBenderMetadata { @@ -136,103 +149,197 @@ fn simple_shader() { operations: vec![ Operation::Normal { opcode: Opcode::Rcp, - dst: 2, - mask: 64, - src: 2, - other: 0, + dst: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::R], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mul, - dst: 2, - mask: 64, - src: 2, - other: 0, + dst: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::R], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Rcp, - dst: 2, - mask: 49, - src: 176, - other: 0, + dst: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mul, - dst: 2, - mask: 49, - src: 176, - other: 0, + dst: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mov, - dst: 3, - mask: 193, - src: 82, - other: 0, + dst: PixelBenderReg { + index: 3, + channels: vec![PixelBenderRegChannel::R, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::G, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mul, - dst: 3, - mask: 193, - src: 178, - other: 0, + dst: PixelBenderReg { + index: 3, + channels: vec![PixelBenderRegChannel::R, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::B, PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mov, - dst: 2, - mask: 97, - src: 19, - other: 0, + dst: PixelBenderReg { + index: 2, + channels: vec![PixelBenderRegChannel::G, PixelBenderRegChannel::B], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 3, + channels: vec![PixelBenderRegChannel::R, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, }, Operation::SampleNearest { - dst: 3, - mask: 241, - src: 16, + dst: PixelBenderReg { + index: 3, + channels: PixelBenderRegChannel::RGBA.to_vec(), + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 0, + channels: vec![PixelBenderRegChannel::R, PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, tf: 0, }, Operation::LoadFloat { - dst: 4, - mask: 128, + dst: PixelBenderReg { + index: 4, + channels: vec![PixelBenderRegChannel::R], + kind: PixelBenderRegKind::Float, + }, val: 100.0, }, Operation::LoadFloat { - dst: 4, - mask: 64, + dst: PixelBenderReg { + index: 4, + channels: vec![PixelBenderRegChannel::G], + kind: PixelBenderRegKind::Float, + }, val: 0.0, }, Operation::LoadFloat { - dst: 4, - mask: 32, + dst: PixelBenderReg { + index: 4, + channels: vec![PixelBenderRegChannel::B], + kind: PixelBenderRegKind::Float, + }, val: 100.0, }, Operation::LoadFloat { - dst: 4, - mask: 16, + dst: PixelBenderReg { + index: 4, + channels: vec![PixelBenderRegChannel::A], + kind: PixelBenderRegKind::Float, + }, val: 1.0, }, Operation::Normal { opcode: Opcode::Mov, - dst: 5, - mask: 243, - src: 30, - other: 0, + dst: PixelBenderReg { + index: 5, + channels: vec![ + PixelBenderRegChannel::R, + PixelBenderRegChannel::G, + PixelBenderRegChannel::B, + PixelBenderRegChannel::A, + ], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 3, + channels: vec![ + PixelBenderRegChannel::R, + PixelBenderRegChannel::G, + PixelBenderRegChannel::B, + PixelBenderRegChannel::A, + ], + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Add, - dst: 5, - mask: 243, - src: 31, - other: 0, + dst: PixelBenderReg { + index: 5, + channels: vec![ + PixelBenderRegChannel::R, + PixelBenderRegChannel::G, + PixelBenderRegChannel::B, + PixelBenderRegChannel::A, + ], + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 4, + channels: PixelBenderRegChannel::RGBA.to_vec(), + kind: PixelBenderRegKind::Float, + }, }, Operation::Normal { opcode: Opcode::Mov, - dst: 1, - mask: 243, - src: 32, - other: 0, + dst: PixelBenderReg { + index: 1, + channels: PixelBenderRegChannel::RGBA.to_vec(), + kind: PixelBenderRegKind::Float, + }, + src: PixelBenderReg { + index: 5, + channels: PixelBenderRegChannel::RGBA.to_vec(), + kind: PixelBenderRegKind::Float, + }, }, ], }; - let shader = parse_shader(shader); + let shader = parse_shader(shader).expect("Failed to parse shader"); assert_eq!(shader, expected, "Shader parsed incorrectly!"); } diff --git a/render/webgl/src/lib.rs b/render/webgl/src/lib.rs index 824ded50b..4386b81e6 100644 --- a/render/webgl/src/lib.rs +++ b/render/webgl/src/lib.rs @@ -1099,6 +1099,24 @@ impl RenderBackend for WebGlRenderBackend { } fn set_quality(&mut self, _quality: StageQuality) {} + + fn compile_pixelbender_shader( + &mut self, + _shader: ruffle_render::pixel_bender::PixelBenderShader, + ) -> Result { + Err(BitmapError::Unimplemented( + "compile_pixelbender_shader".into(), + )) + } + + fn run_pixelbender_shader( + &mut self, + _handle: ruffle_render::pixel_bender::PixelBenderShaderHandle, + _arguments: &[ruffle_render::pixel_bender::PixelBenderShaderArgument], + _target: BitmapHandle, + ) -> Result, BitmapError> { + Err(BitmapError::Unimplemented("run_pixelbender_shader".into())) + } } impl CommandHandler for WebGlRenderBackend { diff --git a/render/wgpu/Cargo.toml b/render/wgpu/Cargo.toml index 12454b04b..ea09393b7 100644 --- a/render/wgpu/Cargo.toml +++ b/render/wgpu/Cargo.toml @@ -23,10 +23,12 @@ ouroboros = "0.15.6" typed-arena = "2.0.2" gc-arena = { workspace = true } naga-agal = { path = "../naga-agal" } +naga-pixelbender = { path = "../naga-pixelbender" } downcast-rs = "1.2.0" profiling = { version = "1.0", default-features = false, optional = true } -naga = { version = "0.12.2", features = ["validate", "wgsl-out"] } lru = "0.10.0" +naga = { workspace = true } +indexmap = "1.9.3" # desktop [target.'cfg(not(target_family = "wasm"))'.dependencies.futures] diff --git a/render/wgpu/src/backend.rs b/render/wgpu/src/backend.rs index a9291d399..4820320e8 100644 --- a/render/wgpu/src/backend.rs +++ b/render/wgpu/src/backend.rs @@ -20,6 +20,9 @@ use ruffle_render::bitmap::{ use ruffle_render::commands::CommandList; use ruffle_render::error::Error as BitmapError; use ruffle_render::filters::Filter; +use ruffle_render::pixel_bender::{ + PixelBenderShader, PixelBenderShaderArgument, PixelBenderShaderHandle, +}; use ruffle_render::quality::StageQuality; use ruffle_render::shape_utils::DistilledShape; use ruffle_render::tessellator::ShapeTessellator; @@ -36,7 +39,7 @@ use tracing::instrument; const TEXTURE_READS_BEFORE_PROMOTION: u8 = 5; pub struct WgpuRenderBackend { - descriptors: Arc, + pub(crate) descriptors: Arc, uniform_buffers_storage: BufferStorage, color_buffers_storage: BufferStorage, target: T, @@ -48,7 +51,7 @@ pub struct WgpuRenderBackend { viewport_scale_factor: f64, texture_pool: TexturePool, offscreen_texture_pool: TexturePool, - offscreen_buffer_pool: Arc>, + pub(crate) offscreen_buffer_pool: Arc>, } impl WgpuRenderBackend { @@ -758,6 +761,22 @@ impl RenderBackend for WgpuRenderBackend { }) => unreachable!("Buffer must be Borrowed as it was set to be Borrowed earlier"), } } + + fn compile_pixelbender_shader( + &mut self, + shader: PixelBenderShader, + ) -> Result { + self.compile_pixelbender_shader_impl(shader) + } + + fn run_pixelbender_shader( + &mut self, + shader: PixelBenderShaderHandle, + arguments: &[PixelBenderShaderArgument], + target_handle: BitmapHandle, + ) -> Result, BitmapError> { + self.run_pixelbender_shader_impl(shader, arguments, target_handle) + } } pub async fn request_adapter_and_device( diff --git a/render/wgpu/src/lib.rs b/render/wgpu/src/lib.rs index de9e30c5d..229b4e4de 100644 --- a/render/wgpu/src/lib.rs +++ b/render/wgpu/src/lib.rs @@ -31,6 +31,7 @@ mod bitmaps; mod context3d; mod globals; mod pipelines; +mod pixel_bender; pub mod target; mod uniform_buffer; diff --git a/render/wgpu/src/pixel_bender.rs b/render/wgpu/src/pixel_bender.rs new file mode 100644 index 000000000..9b1e82781 --- /dev/null +++ b/render/wgpu/src/pixel_bender.rs @@ -0,0 +1,492 @@ +use std::cell::RefCell; +use std::num::NonZeroU64; +use std::{borrow::Cow, cell::Cell, sync::Arc}; + +use indexmap::IndexMap; +use ruffle_render::error::Error as BitmapError; +use ruffle_render::pixel_bender::{ + PixelBenderShaderHandle, PixelBenderShaderImpl, PixelBenderType, OUT_COORD_NAME, +}; +use ruffle_render::{ + bitmap::{BitmapHandle, PixelRegion, SyncHandle}, + pixel_bender::{PixelBenderParam, PixelBenderShader, PixelBenderShaderArgument}, +}; +use wgpu::util::StagingBelt; +use wgpu::{ + BindGroupEntry, BindingResource, BlendComponent, BufferDescriptor, BufferUsages, + ColorTargetState, ColorWrites, FrontFace, ImageCopyTexture, RenderPipelineDescriptor, + SamplerBindingType, ShaderModuleDescriptor, TextureDescriptor, TextureFormat, TextureView, + VertexState, +}; + +use crate::{ + as_texture, + backend::WgpuRenderBackend, + descriptors::Descriptors, + pipelines::VERTEX_BUFFERS_DESCRIPTION_POS, + target::{RenderTarget, RenderTargetFrame, TextureTarget}, + QueueSyncHandle, Texture, +}; + +#[derive(Debug)] +pub struct PixelBenderWgpuShader { + bind_group_layout: wgpu::BindGroupLayout, + pipeline: wgpu::RenderPipeline, + shader: PixelBenderShader, + float_parameters_buffer: wgpu::Buffer, + float_parameters_buffer_size: u64, + int_parameters_buffer: wgpu::Buffer, + int_parameters_buffer_size: u64, + staging_belt: RefCell, +} + +impl PixelBenderShaderImpl for PixelBenderWgpuShader { + fn parsed_shader(&self) -> &PixelBenderShader { + &self.shader + } +} + +pub fn as_cache_holder(handle: &PixelBenderShaderHandle) -> &PixelBenderWgpuShader { + ::downcast_ref(&*handle.0).unwrap() +} + +impl PixelBenderWgpuShader { + pub fn new(descriptors: &Descriptors, shader: PixelBenderShader) -> PixelBenderWgpuShader { + let mut layout_entries = vec![ + // One sampler per filter/wrapping combination - see BitmapFilters + // An AGAL shader can use any of these samplers, so + // we need to bind them all. + wgpu::BindGroupLayoutEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_NEAREST, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(SamplerBindingType::Filtering), + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_LINEAR, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(SamplerBindingType::Filtering), + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_BILINEAR, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(SamplerBindingType::Filtering), + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: naga_pixelbender::SHADER_FLOAT_PARAMETERS_INDEX, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: naga_pixelbender::SHADER_INT_PARAMETERS_INDEX, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ]; + + for param in &shader.params { + if let PixelBenderParam::Texture { index, .. } = param { + let binding = naga_pixelbender::TEXTURE_START_BIND_INDEX + *index as u32; + layout_entries.push(wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: None, + }); + } + } + + let globals_layout_label = + create_debug_label!("PixelBender bind group layout for {:?}", shader.name); + let bind_group_layout = + descriptors + .device + .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: globals_layout_label.as_deref(), + entries: &layout_entries, + }); + + let pipeline_layout_label = + create_debug_label!("PixelBender pipeline layout for {:?}", shader.name); + let pipeline_layout = + descriptors + .device + .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: pipeline_layout_label.as_deref(), + bind_group_layouts: &[&bind_group_layout], + push_constant_ranges: &[], + }); + + let shaders = + naga_pixelbender::ShaderBuilder::build(&shader).expect("Failed to compile shader"); + + let float_label = + create_debug_label!("PixelBender float parameters buffer for {:?}", shader.name); + + let float_parameters_buffer = descriptors.device.create_buffer(&BufferDescriptor { + label: float_label.as_deref(), + size: shaders.float_parameters_buffer_size, + usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + let int_label = + create_debug_label!("PixelBender int parameters buffer for {:?}", shader.name); + + let int_parameters_buffer = descriptors.device.create_buffer(&BufferDescriptor { + label: int_label.as_deref(), + size: shaders.int_parameters_buffer_size, + usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + let vertex_shader = descriptors + .device + .create_shader_module(ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Naga(Cow::Owned(shaders.vertex)), + }); + + let fragment_shader = descriptors + .device + .create_shader_module(ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Naga(Cow::Owned(shaders.fragment)), + }); + + let pipeline = descriptors + .device + .create_render_pipeline(&RenderPipelineDescriptor { + label: create_debug_label!("RenderPipeline").as_deref(), + layout: Some(&pipeline_layout), + vertex: VertexState { + module: &vertex_shader, + entry_point: naga_pixelbender::SHADER_ENTRYPOINT, + buffers: &VERTEX_BUFFERS_DESCRIPTION_POS, + }, + fragment: Some(wgpu::FragmentState { + module: &fragment_shader, + entry_point: naga_pixelbender::SHADER_ENTRYPOINT, + targets: &[Some(ColorTargetState { + format: TextureFormat::Rgba8Unorm, + // FIXME - what should this be? + blend: Some(wgpu::BlendState { + color: BlendComponent::OVER, + alpha: BlendComponent::OVER, + }), + write_mask: ColorWrites::all(), + })], + }), + primitive: wgpu::PrimitiveState { + front_face: FrontFace::Ccw, + cull_mode: None, + ..Default::default() + }, + depth_stencil: None, + multisample: wgpu::MultisampleState { + count: 1, + mask: !0, + alpha_to_coverage_enabled: false, + }, + multiview: Default::default(), + }); + + PixelBenderWgpuShader { + bind_group_layout, + pipeline, + shader, + float_parameters_buffer, + float_parameters_buffer_size: shaders.float_parameters_buffer_size, + int_parameters_buffer, + int_parameters_buffer_size: shaders.int_parameters_buffer_size, + // FIXME - come up with a good chunk size + staging_belt: RefCell::new(StagingBelt::new(8)), + } + } +} + +impl WgpuRenderBackend { + pub(super) fn compile_pixelbender_shader_impl( + &mut self, + shader: PixelBenderShader, + ) -> Result { + let handle = PixelBenderWgpuShader::new(&self.descriptors, shader); + Ok(PixelBenderShaderHandle(Arc::new(handle))) + } + + pub(super) fn run_pixelbender_shader_impl( + &mut self, + shader: PixelBenderShaderHandle, + arguments: &[PixelBenderShaderArgument], + target_handle: BitmapHandle, + ) -> Result, BitmapError> { + let compiled_shader = &as_cache_holder(&shader); + let mut staging_belt = compiled_shader.staging_belt.borrow_mut(); + + let mut arguments = arguments.to_vec(); + + let target = as_texture(&target_handle); + let extent = wgpu::Extent3d { + width: target.width, + height: target.height, + depth_or_array_layers: 1, + }; + + let mut texture_target = TextureTarget { + size: extent, + texture: target.texture.clone(), + format: wgpu::TextureFormat::Rgba8Unorm, + buffer: None, + }; + + let frame_output = texture_target + .get_next_texture() + .expect("TextureTargetFrame.get_next_texture is infallible"); + + let mut bind_group_entries = vec![ + BindGroupEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_NEAREST, + resource: BindingResource::Sampler(&self.descriptors.bitmap_samplers.clamp_nearest), + }, + BindGroupEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_LINEAR, + resource: BindingResource::Sampler(&self.descriptors.bitmap_samplers.clamp_linear), + }, + BindGroupEntry { + binding: naga_pixelbender::SAMPLER_CLAMP_BILINEAR, + // FIXME - create bilinear sampler + resource: BindingResource::Sampler(&self.descriptors.bitmap_samplers.clamp_linear), + }, + BindGroupEntry { + binding: naga_pixelbender::SHADER_FLOAT_PARAMETERS_INDEX, + resource: BindingResource::Buffer(wgpu::BufferBinding { + buffer: &compiled_shader.float_parameters_buffer, + offset: 0, + size: Some( + NonZeroU64::new(compiled_shader.float_parameters_buffer_size).unwrap(), + ), + }), + }, + BindGroupEntry { + binding: naga_pixelbender::SHADER_INT_PARAMETERS_INDEX, + resource: BindingResource::Buffer(wgpu::BufferBinding { + buffer: &compiled_shader.int_parameters_buffer, + offset: 0, + size: Some( + NonZeroU64::new(compiled_shader.int_parameters_buffer_size).unwrap(), + ), + }), + }, + ]; + + let mut texture_views: IndexMap = Default::default(); + + let mut render_command_encoder = + self.descriptors + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: create_debug_label!("Render command encoder").as_deref(), + }); + + let mut target_clone = None; + + let mut float_offset = 0; + let mut int_offset = 0; + + for input in &mut arguments { + match input { + PixelBenderShaderArgument::ImageInput { index, texture, .. } => { + // The input is the same as the output - we need to clone the input. + // We will write to the original output, and use a clone of the input as a texture input binding + if std::ptr::eq( + Arc::as_ptr(&texture.0) as *const (), + Arc::as_ptr(&target_handle.0) as *const (), + ) { + let cached_fresh_handle = target_clone.get_or_insert_with(|| { + let extent = wgpu::Extent3d { + width: target.width, + height: target.height, + depth_or_array_layers: 1, + }; + let fresh_texture = + self.descriptors.device.create_texture(&TextureDescriptor { + label: Some("PixelBenderShader target clone"), + size: extent, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rgba8Unorm, + usage: wgpu::TextureUsages::COPY_DST + | wgpu::TextureUsages::TEXTURE_BINDING, + view_formats: &[wgpu::TextureFormat::Rgba8Unorm], + }); + render_command_encoder.copy_texture_to_texture( + ImageCopyTexture { + texture: &target.texture, + mip_level: 0, + origin: Default::default(), + aspect: Default::default(), + }, + ImageCopyTexture { + texture: &fresh_texture, + mip_level: 0, + origin: Default::default(), + aspect: Default::default(), + }, + extent, + ); + + BitmapHandle(Arc::new(Texture { + texture: Arc::new(fresh_texture), + bind_linear: Default::default(), + bind_nearest: Default::default(), + width: extent.width, + height: extent.height, + copy_count: Cell::new(0), + })) + }); + *texture = cached_fresh_handle.clone(); + } + texture_views.insert( + *index, + as_texture(texture) + .texture + .create_view(&wgpu::TextureViewDescriptor::default()), + ); + } + PixelBenderShaderArgument::ValueInput { index, value } => { + let param = &compiled_shader.shader.params[*index as usize]; + + let name = match param { + PixelBenderParam::Normal { name, .. } => name, + _ => unreachable!(), + }; + + if name == OUT_COORD_NAME { + continue; + } + + let (value_vec, is_float): ([f32; 4], bool) = match value { + PixelBenderType::TFloat(f1) => ([*f1, 0.0, 0.0, 0.0], true), + PixelBenderType::TFloat2(f1, f2) => ([*f1, *f2, 0.0, 0.0], true), + PixelBenderType::TFloat3(f1, f2, f3) => ([*f1, *f2, *f3, 0.0], true), + PixelBenderType::TFloat4(f1, f2, f3, f4) => ([*f1, *f2, *f3, *f4], true), + PixelBenderType::TInt(i1) => ([*i1 as f32, 0.0, 0.0, 0.0], false), + PixelBenderType::TInt2(i1, i2) => { + ([*i1 as f32, *i2 as f32, 0.0, 0.0], false) + } + PixelBenderType::TInt3(i1, i2, i3) => { + ([*i1 as f32, *i2 as f32, *i3 as f32, 0.0], false) + } + PixelBenderType::TInt4(i1, i2, i3, i4) => { + ([*i1 as f32, *i2 as f32, *i3 as f32, *i4 as f32], false) + } + _ => unreachable!("Unimplemented value {value:?}"), + }; + + // Both float32 and int are 4 bytes + let component_size_bytes = 4; + + let (buffer, vec4_count) = if is_float { + let res = (&compiled_shader.float_parameters_buffer, float_offset); + float_offset += 1; + res + } else { + let res = (&compiled_shader.int_parameters_buffer, int_offset); + int_offset += 1; + res + }; + + let mut buffer_slice = staging_belt.write_buffer( + &mut render_command_encoder, + buffer, + vec4_count * 4 * component_size_bytes, + NonZeroU64::new(4 * component_size_bytes).unwrap(), + &self.descriptors.device, + ); + buffer_slice.copy_from_slice(bytemuck::cast_slice(&value_vec)); + } + } + } + + // This needs to be a separate loop, so that we can get references into `texture_views` + for input in &arguments { + match input { + PixelBenderShaderArgument::ImageInput { index, .. } => { + let binding = naga_pixelbender::TEXTURE_START_BIND_INDEX + *index as u32; + bind_group_entries.push(BindGroupEntry { + binding, + resource: BindingResource::TextureView(&texture_views[index]), + }); + } + PixelBenderShaderArgument::ValueInput { .. } => {} + } + } + + let bind_group = self + .descriptors + .device + .create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &compiled_shader.bind_group_layout, + entries: &bind_group_entries, + }); + + staging_belt.finish(); + + let mut render_pass = + render_command_encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("PixelBender render pass"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: frame_output.view(), + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), + store: true, + }, + })], + depth_stencil_attachment: None, + }); + render_pass.set_bind_group(0, &bind_group, &[]); + render_pass.set_pipeline(&compiled_shader.pipeline); + + render_pass.set_vertex_buffer(0, self.descriptors.quad.vertices_pos.slice(..)); + render_pass.set_index_buffer( + self.descriptors.quad.indices.slice(..), + wgpu::IndexFormat::Uint32, + ); + + render_pass.draw_indexed(0..6, 0, 0..1); + + drop(render_pass); + + self.descriptors + .queue + .submit(Some(render_command_encoder.finish())); + + staging_belt.recall(); + + Ok(Box::new(QueueSyncHandle::NotCopied { + handle: target_handle, + copy_area: PixelRegion::for_whole_size(extent.width, extent.height), + descriptors: self.descriptors.clone(), + pool: self.offscreen_buffer_pool.clone(), + })) + } +} diff --git a/tests/tests/swfs/avm2/pixelbender_images/Test.as b/tests/tests/swfs/avm2/pixelbender_images/Test.as new file mode 100755 index 000000000..c87101492 --- /dev/null +++ b/tests/tests/swfs/avm2/pixelbender_images/Test.as @@ -0,0 +1,35 @@ +package { + import flash.display.BitmapData; + import flash.display.ShaderJob; + import flash.display.Shader; + import flash.display.Bitmap; + import flash.display.MovieClip; + + public class Test { + + [Embed(source = "mandelbrot.png")] + public static var MANDELBROT: Class; + + // Shader from https://github.com/8bitavenue/Adobe-Pixel-Bender-Effects/blob/master/Donut%20Shader.cpp + [Embed(source = "donut.pbj", mimeType="application/octet-stream")] + public static var DONUT_BYTES: Class; + + public function Test(main: MovieClip) { + var mandelbrot: Bitmap = new MANDELBROT(); + main.addChild(new Bitmap(donut(mandelbrot.bitmapData.clone()))); + } + + private function donut(input: BitmapData): BitmapData { + var shader = new ShaderJob(new Shader(new DONUT_BYTES()), input); + shader.shader.data.BlockCount.value = [56.5]; + shader.shader.data.Min.value = [0.29]; + shader.shader.data.Max.value = [0.51]; + shader.shader.data.Width.value = [100.0]; + shader.shader.data.Height.value = [100.0]; + shader.shader.data.color.value = [0.34, 0.1, 0.2, 1]; + shader.shader.data.src.input = input; + shader.start(true); + return input + } + } +} \ No newline at end of file diff --git a/tests/tests/swfs/avm2/pixelbender_images/donut.pbj b/tests/tests/swfs/avm2/pixelbender_images/donut.pbj new file mode 100755 index 000000000..a6dc9b648 Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_images/donut.pbj differ diff --git a/tests/tests/swfs/avm2/pixelbender_images/donut.pbk b/tests/tests/swfs/avm2/pixelbender_images/donut.pbk new file mode 100755 index 000000000..a8c72ce74 --- /dev/null +++ b/tests/tests/swfs/avm2/pixelbender_images/donut.pbk @@ -0,0 +1,98 @@ + +kernel Donut +< + namespace : "8bitavenue"; + vendor : "8bitavenue"; + version : 1; +> + +{ + //Donut density + parameter float BlockCount + < + minValue: 1.0; + maxValue: 100.0; + defaultValue: 5.0; + >; + + //Inner circle + parameter float Min + < + minValue: 0.0; + maxValue: 1.0; + defaultValue: 0.25; + >; + + //Outer circle + parameter float Max + < + minValue: 0.0; + maxValue: 1.0; + defaultValue: 0.45; + >; + + //Scale width + parameter float Width + < + minValue: 1.0; + maxValue: 1000.0; + defaultValue: 100.0; + >; + + //Scale height + parameter float Height + < + minValue: 1.0; + maxValue: 1000.0; + defaultValue: 100.0; + >; + + //Background color + parameter pixel4 color + < + minValue: float4(0.0,0.0,0.0,0.0); + maxValue: float4(1.0,1.0,1.0,1.0); + defaultValue: float4(0.2, 0.2, 0.2, 1.0); + >; + + //Input image + input image4 src; + + //Output image + output pixel4 dst; + + //Apply this filter + void evaluatePixel() + { + //Calculate block size + float myblockcount = BlockCount/5.0; + float BlockSize = 1.0/myblockcount; + + float2 temp = outCoord(); + temp.x = temp.x/Width; + temp.y = temp.y/Height; + + //Calculate block position and center + float2 blockPos = floor(temp * myblockcount); + float2 blockCenter = blockPos * BlockSize + BlockSize * 0.5; + + //Pixel distance from center + float dist = length(temp - blockCenter) * myblockcount; + + //If pixel is inside inner circle + //or outside outer circle then color + //it with background color + //otherwise color it with the color + //of the pixel at the center + if(dist < Min || dist > Max) + { + dst = color; + } + else + { + blockCenter.x = blockCenter.x * Width; + blockCenter.y = blockCenter.y * Height; + dst = sampleNearest(src, blockCenter); + } + } +} \ No newline at end of file diff --git a/tests/tests/swfs/avm2/pixelbender_images/expected.png b/tests/tests/swfs/avm2/pixelbender_images/expected.png new file mode 100644 index 000000000..c0aac90fb Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_images/expected.png differ diff --git a/tests/tests/swfs/avm2/pixelbender_images/mandelbrot.png b/tests/tests/swfs/avm2/pixelbender_images/mandelbrot.png new file mode 100755 index 000000000..7f169b437 Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_images/mandelbrot.png differ diff --git a/tests/tests/swfs/avm2/pixelbender_images/output.txt b/tests/tests/swfs/avm2/pixelbender_images/output.txt new file mode 100644 index 000000000..e69de29bb diff --git a/tests/tests/swfs/avm2/pixelbender_images/test.fla b/tests/tests/swfs/avm2/pixelbender_images/test.fla new file mode 100755 index 000000000..e95ca578e Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_images/test.fla differ diff --git a/tests/tests/swfs/avm2/pixelbender_images/test.swf b/tests/tests/swfs/avm2/pixelbender_images/test.swf new file mode 100755 index 000000000..cab75ce7f Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_images/test.swf differ diff --git a/tests/tests/swfs/avm2/pixelbender_images/test.toml b/tests/tests/swfs/avm2/pixelbender_images/test.toml new file mode 100644 index 000000000..d8a3f915a --- /dev/null +++ b/tests/tests/swfs/avm2/pixelbender_images/test.toml @@ -0,0 +1,7 @@ +num_frames = 1 + +[image_comparison] +tolerance = 1 + +[player_options] +with_renderer = { optional = false, sample_count = 1 } diff --git a/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbj b/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbj new file mode 100755 index 000000000..e09439e29 Binary files /dev/null and b/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbj differ diff --git a/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbk b/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbk new file mode 100755 index 000000000..e71ba430a --- /dev/null +++ b/tests/tests/swfs/avm2/pixelbender_shaderdata/simple_shader.pbk @@ -0,0 +1,36 @@ + + + kernel DoNothing + < + namespace: "Adobe::Example"; + vendor: "Adobe examples"; + version: 1; + description: "A shader that does nothing, but does it well."; + > + { + + output pixel4 dst; + + parameter float radius + < + description: "The radius of the effect"; + minValue: 0.0; + maxValue: 50.0; + defaultValue: 25.0; + >; + + parameter float otherParam + < + description: "Other param"; + minValue: 0.0; + maxValue: 255.0; + defaultValue: 25.0; + >; + + input image4 src; + + void evaluatePixel() + { + dst = float4((otherParam + radius) / 255.0, 0.0, 0.0, 1.0); + } + } diff --git a/tests/tests/swfs/avm2/pixelbender_shaderdata/test.toml b/tests/tests/swfs/avm2/pixelbender_shaderdata/test.toml index dbee897f5..1a5d49917 100644 --- a/tests/tests/swfs/avm2/pixelbender_shaderdata/test.toml +++ b/tests/tests/swfs/avm2/pixelbender_shaderdata/test.toml @@ -1 +1,4 @@ num_frames = 1 + +[player_options] +with_renderer = { optional = false, sample_count = 1 }