From d0230a2beaa99dd067b94cf34d484f10545f3210 Mon Sep 17 00:00:00 2001 From: Aaron Hill Date: Thu, 22 Sep 2022 19:44:49 -0500 Subject: [PATCH] render: Add 'naga-agal' crate to compile AGAL shaders to Naga This is the first part of the Stage3D implementation, and can be reviewed independently. Stage3D shaders use the Adobe Graphics Assembly Language (AGAL), which is a binary shader format. It supports vertex attributes, varying registers, program constants (uniforms), and texture sampling. This PR only implements a few parts of AGAL: * The 'mov' and 'm44' opcodes * Vertex attributes, varying registers, program constants, and 'output' registers (position or color, depending on shader type) This is sufficient to get a non-trivial Stage3D program running (the rotating cube demo from the Adobe docs). The output of `naga-agal` is a `naga::Module`. This can be passed directly to wgpu, or compiled into a shader language using a Naga backend (glsl, wgsl, SPIR-V, etc). The test suite output WGSL files, and uses the 'insta' crate to compare against saved files on disk. Currently, the only real way to write AGAL bytecode is using the Adobe-provided 'AGALMiniAssembler' flash class. This class assembles the textual reprentation of AGAL into the binary format. To make writing tests easier, I've added a 'agal_compiler' test, which can easily be modified to add more Agal textual assembly. --- Cargo.lock | 45 + Cargo.toml | 1 + render/naga-agal/Cargo.toml | 16 + render/naga-agal/src/builder.rs | 761 ++++++++++++++++ render/naga-agal/src/lib.rs | 93 ++ render/naga-agal/src/types.rs | 195 +++++ .../tests/snapshots/wgsl__shaders-2.snap | 25 + .../tests/snapshots/wgsl__shaders-3.snap | 21 + .../tests/snapshots/wgsl__shaders.snap | 29 + render/naga-agal/tests/wgsl.rs | 83 ++ tests/tests/regression_tests.rs | 1 + .../avm2/agal_compiler/AGALMiniAssembler.as | 819 ++++++++++++++++++ tests/tests/swfs/avm2/agal_compiler/Test.as | 54 ++ .../tests/swfs/avm2/agal_compiler/output.txt | 14 + tests/tests/swfs/avm2/agal_compiler/test.fla | Bin 0 -> 3815 bytes tests/tests/swfs/avm2/agal_compiler/test.swf | Bin 0 -> 10580 bytes 16 files changed, 2157 insertions(+) create mode 100644 render/naga-agal/Cargo.toml create mode 100644 render/naga-agal/src/builder.rs create mode 100644 render/naga-agal/src/lib.rs create mode 100644 render/naga-agal/src/types.rs create mode 100644 render/naga-agal/tests/snapshots/wgsl__shaders-2.snap create mode 100644 render/naga-agal/tests/snapshots/wgsl__shaders-3.snap create mode 100644 render/naga-agal/tests/snapshots/wgsl__shaders.snap create mode 100644 render/naga-agal/tests/wgsl.rs create mode 100644 tests/tests/swfs/avm2/agal_compiler/AGALMiniAssembler.as create mode 100644 tests/tests/swfs/avm2/agal_compiler/Test.as create mode 100644 tests/tests/swfs/avm2/agal_compiler/output.txt create mode 100644 tests/tests/swfs/avm2/agal_compiler/test.fla create mode 100644 tests/tests/swfs/avm2/agal_compiler/test.swf diff --git a/Cargo.lock b/Cargo.lock index 76dda7c59..199727ea0 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1975,6 +1975,19 @@ dependencies = [ "unicode-width", ] +[[package]] +name = "insta" +version = "1.21.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "581d4e3314cae4536e5d22ffd23189d4a374696c5ef733eadafae0ed273fd303" +dependencies = [ + "console", + "lazy_static", + "linked-hash-map", + "similar", + "yaml-rust", +] + [[package]] name = "instant" version = "0.1.12" @@ -2182,6 +2195,12 @@ dependencies = [ "cc", ] +[[package]] +name = "linked-hash-map" +version = "0.5.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0717cef1bc8b636c6e1c1bbdefc09e6322da8a9321966e8928ef80d20f7f770f" + [[package]] name = "lock_api" version = "0.4.9" @@ -2401,6 +2420,17 @@ dependencies = [ "unicode-xid", ] +[[package]] +name = "naga-agal" +version = "0.1.0" +dependencies = [ + "bitflags", + "insta", + "naga", + "num-derive", + "num-traits", +] + [[package]] name = "nanorand" version = "0.7.0" @@ -3679,6 +3709,12 @@ version = "1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "43b2853a4d09f215c24cc5489c992ce46052d359b5109343cbafbf26bc62f8a3" +[[package]] +name = "similar" +version = "2.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62ac7f900db32bf3fd12e0117dd3dc4da74bc52ebaac97f39668446d89694803" + [[package]] name = "simple_asn1" version = "0.6.2" @@ -4779,6 +4815,15 @@ version = "0.8.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d2d7d3948613f75c98fd9328cfdcc45acc4d360655289d0a7d4ec931392200a3" +[[package]] +name = "yaml-rust" +version = "0.4.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "56c1936c4cc7a1c9ab21a1ebb602eb942ba868cbd44a99cb7cdc5892335e1c85" +dependencies = [ + "linked-hash-map", +] + [[package]] name = "yansi" version = "0.5.1" diff --git a/Cargo.toml b/Cargo.toml index 9c85a2960..0bb4b8d91 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -13,6 +13,7 @@ members = [ "render", "render/canvas", + "render/naga-agal", "render/wgpu", "render/webgl", diff --git a/render/naga-agal/Cargo.toml b/render/naga-agal/Cargo.toml new file mode 100644 index 000000000..dd4853146 --- /dev/null +++ b/render/naga-agal/Cargo.toml @@ -0,0 +1,16 @@ +[package] +name = "naga-agal" +version = "0.1.0" +edition = "2021" + +# See more keys and their definitions at https://doc.rust-lang.org/cargo/reference/manifest.html + +[dependencies] +bitflags = "1.3.2" +naga = { version = "0.10.0" } +num-derive = "0.3.3" +num-traits = "0.2.15" + +[dev-dependencies] +insta = "1.20.0" +naga = { version = "0.10.0", features = ["wgsl-out", "validate"] } diff --git a/render/naga-agal/src/builder.rs b/render/naga-agal/src/builder.rs new file mode 100644 index 000000000..806b30de3 --- /dev/null +++ b/render/naga-agal/src/builder.rs @@ -0,0 +1,761 @@ +use std::io::Read; + +use naga::{ + ArraySize, BuiltIn, Constant, ConstantInner, EntryPoint, FunctionArgument, FunctionResult, + GlobalVariable, Interpolation, ScalarValue, ShaderStage, StructMember, SwizzleComponent, +}; +use naga::{BinaryOperator, MathFunction}; +use naga::{ + Binding, Expression, Function, Handle, LocalVariable, Module, ScalarKind, Span, Statement, + Type, TypeInner, VectorSize, +}; +use num_traits::FromPrimitive; + +use crate::{ + types::*, Error, ShaderType, VertexAttributeFormat, ENTRY_POINT, MAX_VERTEX_ATTRIBUTES, +}; + +const VERTEX_PROGRAM_CONTANTS: u64 = 128; +const FRAGMENT_PROGRAM_CONSTANTS: u64 = 28; + +pub type Result = std::result::Result; + +pub(crate) struct NagaBuilder<'a> { + module: Module, + func: Function, + + // This evaluate to a Pointer to the temporary 'main' destiation location + // (the output position for a vertex shader, or the output color for a fragment shader) + // which can be used with Expression::Load and Expression::Store + // This is needed because an AGAL shader can write to the output register + // multiple times. + dest: Handle, + + shader_config: ShaderConfig<'a>, + // Whenever we read from a varying register in a fragment shader, + // we create a new argument binding for it. + argument_expressions: Vec>>, + + varying_pointers: Vec>>, + + // An `Expression::GlobalVariables` for the uniform buffer + // that stores all of the program constants. + constant_registers: Handle, + + // The function return type being built up. Each time a vertex + // shader writes to a varying register, we add a new member to this + return_type: Type, + + // The Naga representation of 'vec4f' + vec4f: Handle, + // The Naga representation of 'mat4x4f' + matrix4x4f: Handle, +} + +impl VertexAttributeFormat { + fn to_naga_type(self, module: &mut Module) -> Handle { + if let VertexAttributeFormat::Float1 = self { + return module.types.insert( + Type { + name: None, + inner: TypeInner::Scalar { + kind: ScalarKind::Float, + width: 4, + }, + }, + Span::UNDEFINED, + ); + } + let (size, width, kind) = match self { + VertexAttributeFormat::Float1 => unreachable!(), + VertexAttributeFormat::Float2 => (VectorSize::Bi, 4, ScalarKind::Float), + VertexAttributeFormat::Float3 => (VectorSize::Tri, 4, ScalarKind::Float), + VertexAttributeFormat::Float4 => (VectorSize::Quad, 4, ScalarKind::Float), + VertexAttributeFormat::Bytes4 => (VectorSize::Quad, 1, ScalarKind::Uint), + }; + + module.types.insert( + Type { + name: None, + inner: TypeInner::Vector { size, kind, width }, + }, + Span::UNDEFINED, + ) + } + + fn extend_to_float4( + &self, + base_expr: Handle, + builder: &mut NagaBuilder, + ) -> Result> { + Ok(match self { + // This does 'vec4f(my_vec3, 1.0)' + VertexAttributeFormat::Float3 => { + let expr = base_expr; + let mut components = vec![]; + for i in 0..3 { + components.push(builder.evaluate_expr(Expression::AccessIndex { + base: expr, + index: i, + })); + } + components.push(builder.func.expressions.append( + Expression::Constant(builder.module.constants.append( + Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: ScalarValue::Float(1.0), + }, + }, + Span::UNDEFINED, + )), + Span::UNDEFINED, + )); + builder.evaluate_expr(Expression::Compose { + ty: builder.vec4f, + components, + }) + } + VertexAttributeFormat::Float4 => base_expr, + _ => { + return Err(Error::Unimplemented(format!( + "Unsupported conversion from {:?} to float4", + self + ))) + } + }) + } +} + +#[derive(Debug)] +pub struct ShaderConfig<'a> { + pub shader_type: ShaderType, + pub vertex_attributes: &'a [Option; 8], +} + +impl<'a> NagaBuilder<'a> { + pub fn process_agal( + mut agal: &[u8], + vertex_attributes: &[Option; MAX_VERTEX_ATTRIBUTES], + ) -> Result { + let data = &mut agal; + + let mut header = [0; 7]; + data.read_exact(&mut header)?; + + if header[0..6] != [0xA0, 0x01, 0x00, 0x00, 0x00, 0xA1] { + return Err(Error::InvalidHeader); + } + + let shader_type = match header[6] { + 0x00 => ShaderType::Vertex, + 0x01 => ShaderType::Fragment, + _ => return Err(Error::InvalidShaderType(header[6])), + }; + + let mut builder = NagaBuilder::new(ShaderConfig { + shader_type, + vertex_attributes, + }); + + while !data.is_empty() { + let mut token = [0; 24]; + data.read_exact(&mut token)?; + let raw_opcode = u32::from_le_bytes(token[0..4].try_into().unwrap()); + + // FIXME - this is a clippy false-positive + #[allow(clippy::or_fun_call)] + let opcode = Opcode::from_u32(raw_opcode).ok_or(Error::InvalidOpcode(raw_opcode))?; + + let dest = DestField::parse(u32::from_le_bytes(token[4..8].try_into().unwrap()))?; + let source1 = SourceField::parse(u64::from_le_bytes(token[8..16].try_into().unwrap()))?; + + let source2 = if let Opcode::Tex = opcode { + Source2::Sampler(SamplerField::parse(u64::from_le_bytes( + token[16..24].try_into().unwrap(), + ))?) + } else { + Source2::SourceField(SourceField::parse(u64::from_le_bytes( + token[16..24].try_into().unwrap(), + ))?) + }; + + builder.process_opcode(&opcode, &dest, &source1, &source2)?; + } + builder.finish() + } + + fn new(shader_config: ShaderConfig<'a>) -> Self { + let mut module = Module::default(); + let mut func = Function::default(); + + let vec4f = VertexAttributeFormat::Float4.to_naga_type(&mut module); + + let matrix4x4f = module.types.insert( + Type { + name: None, + inner: TypeInner::Matrix { + columns: VectorSize::Quad, + rows: VectorSize::Quad, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + // The return type always has at least one component - the vec4f that's the 'main' + // output of our shader (the position for the vertex shader, and the color for the fragment shader) + let return_type = match shader_config.shader_type { + ShaderType::Vertex => Type { + name: None, + inner: TypeInner::Struct { + members: vec![StructMember { + name: None, + ty: vec4f, + binding: Some(Binding::BuiltIn(BuiltIn::Position { invariant: false })), + offset: 0, + }], + span: 16, + }, + }, + ShaderType::Fragment => Type { + name: None, + inner: TypeInner::Struct { + members: vec![StructMember { + name: None, + ty: vec4f, + binding: Some(Binding::Location { + location: 0, + interpolation: None, + sampling: None, + }), + offset: 0, + }], + span: 16, + }, + }, + }; + + match shader_config.shader_type { + ShaderType::Vertex => { + func.result = Some(FunctionResult { + ty: vec4f, + binding: Some(Binding::BuiltIn(BuiltIn::Position { invariant: false })), + }); + } + ShaderType::Fragment => { + func.result = Some(FunctionResult { + ty: vec4f, + binding: Some(Binding::Location { + location: 0, + interpolation: None, + sampling: None, + }), + }); + } + } + + // Holds the value we're going to return. + // This corresponds to RegisterType::Output + let output_temp_handle = func.local_variables.append( + LocalVariable { + name: Some("dest_temp".to_string()), + ty: vec4f, + init: None, + }, + Span::UNDEFINED, + ); + let dest = func.expressions.append( + Expression::LocalVariable(output_temp_handle), + Span::UNDEFINED, + ); + + let num_const_registers = module.constants.append( + Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: ScalarValue::Uint(match shader_config.shader_type { + ShaderType::Vertex => VERTEX_PROGRAM_CONTANTS, + ShaderType::Fragment => FRAGMENT_PROGRAM_CONSTANTS, + }), + }, + }, + Span::UNDEFINED, + ); + + let binding_num = match shader_config.shader_type { + ShaderType::Vertex => 0, + ShaderType::Fragment => 1, + }; + + let constant_registers_global = module.global_variables.append( + GlobalVariable { + name: Some("constant_registers".to_string()), + space: naga::AddressSpace::Uniform, + binding: Some(naga::ResourceBinding { + group: 0, + binding: binding_num, + }), + ty: module.types.insert( + Type { + name: None, + inner: TypeInner::Array { + base: vec4f, + size: ArraySize::Constant(num_const_registers), + stride: std::mem::size_of::() as u32 * 4, + }, + }, + Span::UNDEFINED, + ), + init: None, + }, + Span::UNDEFINED, + ); + + let constant_registers = func.expressions.append( + Expression::GlobalVariable(constant_registers_global), + Span::UNDEFINED, + ); + + NagaBuilder { + module, + func, + dest, + shader_config, + argument_expressions: vec![], + varying_pointers: vec![], + return_type, + matrix4x4f, + vec4f, + constant_registers, + } + } + + fn get_vertex_input(&mut self, index: usize) -> Result> { + if index >= self.argument_expressions.len() { + self.argument_expressions.resize(index + 1, None); + + // FIXME - this is a clippy false-positive + #[allow(clippy::or_fun_call)] + let ty = self.shader_config.vertex_attributes[index] + .as_ref() + .ok_or(Error::MissingVertexAttributeData(index))? + .to_naga_type(&mut self.module); + + self.func.arguments.push(FunctionArgument { + name: None, + ty, + binding: Some(Binding::Location { + location: index as u32, + interpolation: None, + sampling: None, + }), + }); + + // Arguments map one-to-one to vertex attributes. + let expr = self + .func + .expressions + .append(Expression::FunctionArgument(index as u32), Span::UNDEFINED); + self.argument_expressions[index] = Some(expr); + } + Ok(self.argument_expressions[index].unwrap()) + } + + fn get_varying_pointer(&mut self, index: usize) -> Result> { + if index >= self.varying_pointers.len() { + self.varying_pointers.resize(index + 1, None); + } + + if self.varying_pointers[index].is_none() { + match self.shader_config.shader_type { + ShaderType::Vertex => { + // We can write to varying variables in the vertex shader, + // and the fragment shader will receive them is input. + // Therefore, we create a local variable for each varying, + // and return them at the end of the function. + let local = self.func.local_variables.append( + LocalVariable { + name: Some(format!("varying_{}", index)), + ty: self.vec4f, + init: None, + }, + Span::UNDEFINED, + ); + + let expr = self + .func + .expressions + .append(Expression::LocalVariable(local), Span::UNDEFINED); + let _range = self + .func + .expressions + .range_from(self.func.expressions.len() - 1); + + if let TypeInner::Struct { members, .. } = &mut self.return_type.inner { + members.push(StructMember { + name: Some(format!("varying_{}", index)), + ty: self.vec4f, + binding: Some(Binding::Location { + location: index as u32, + interpolation: Some(naga::Interpolation::Perspective), + sampling: None, + }), + offset: 0, + }); + } else { + unreachable!(); + } + + self.varying_pointers[index] = Some(expr); + } + ShaderType::Fragment => { + self.func.arguments.push(FunctionArgument { + name: None, + ty: self.vec4f, + binding: Some(Binding::Location { + location: index as u32, + interpolation: Some(Interpolation::Perspective), + sampling: None, + }), + }); + + let expr = self + .func + .expressions + .append(Expression::FunctionArgument(index as u32), Span::UNDEFINED); + self.varying_pointers[index] = Some(expr); + } + }; + }; + + Ok(self.varying_pointers[index].unwrap()) + } + + fn emit_const_register_load(&mut self, index: usize) -> Result> { + let index_const = self.module.constants.append( + Constant { + name: None, + specialization: None, + inner: ConstantInner::Scalar { + width: 4, + value: ScalarValue::Uint(index as u64), + }, + }, + Span::UNDEFINED, + ); + let index_expr = self + .func + .expressions + .append(Expression::Constant(index_const), Span::UNDEFINED); + + let register_pointer = self.func.expressions.append( + Expression::Access { + base: self.constant_registers, + index: index_expr, + }, + Span::UNDEFINED, + ); + + Ok(self.evaluate_expr(Expression::Load { + pointer: register_pointer, + })) + } + + fn emit_varying_load(&mut self, index: usize) -> Result> { + // A LocalVariable evaluates to a pointer, so we need to load it + let varying_expr = self.get_varying_pointer(index)?; + Ok(match self.shader_config.shader_type { + ShaderType::Vertex => self.evaluate_expr(Expression::Load { + pointer: varying_expr, + }), + ShaderType::Fragment => varying_expr, + }) + } + + fn emit_source_field_load( + &mut self, + source: &SourceField, + extend_to_vec4: bool, + ) -> Result> { + let (mut base_expr, source_type) = match source.register_type { + // We can use a function argument directly - we don't need + // a separate Expression::Load + RegisterType::Attribute => ( + self.get_vertex_input(source.reg_num as usize)?, + // FIXME - this is a clippy false-positive + #[allow(clippy::or_fun_call)] + self.shader_config.vertex_attributes[source.reg_num as usize] + .ok_or(Error::MissingVertexAttributeData(source.reg_num as usize))?, + ), + RegisterType::Varying => ( + self.emit_varying_load(source.reg_num as usize)?, + VertexAttributeFormat::Float4, + ), + RegisterType::Constant => ( + self.emit_const_register_load(source.reg_num as usize)?, + // Constants are always a vec4 + VertexAttributeFormat::Float4, + ), + _ => { + return Err(Error::Unimplemented(format!( + "Unimplemented source reg type {:?}", + source.register_type + ))) + } + }; + + if matches!(source.direct_mode, DirectMode::Indirect) { + return Err(Error::Unimplemented( + "Indirect addressing not implemented".to_string(), + )); + } + + if extend_to_vec4 && source_type != VertexAttributeFormat::Float4 { + base_expr = source_type.extend_to_float4(base_expr, self)?; + } + + // Swizzle is 'xyzw', which is a no-op. Just return the base expression. + if source.swizzle == 0b11100100 { + return Ok(base_expr); + } + + let swizzle_flags = [ + source.swizzle & 0b11, + (source.swizzle >> 2) & 0b11, + (source.swizzle >> 4) & 0b11, + (source.swizzle >> 6) & 0b11, + ]; + + let swizzle_components: [SwizzleComponent; 4] = swizzle_flags + .into_iter() + .map(|flag| match flag { + 0b00 => SwizzleComponent::X, + 0b01 => SwizzleComponent::Y, + 0b10 => SwizzleComponent::Z, + 0b11 => SwizzleComponent::W, + _ => unreachable!(), + }) + .collect::>() + .try_into() + .unwrap(); + + Ok(self.func.expressions.append( + Expression::Swizzle { + size: VectorSize::Quad, + vector: base_expr, + pattern: swizzle_components, + }, + Span::UNDEFINED, + )) + } + + fn emit_dest_store(&mut self, dest: &DestField, expr: Handle) -> Result<()> { + let base_expr = match dest.register_type { + RegisterType::Output => self.dest, + RegisterType::Varying => self.get_varying_pointer(dest.reg_num as usize)?, + _ => { + return Err(Error::Unimplemented(format!( + "Unimplemented dest reg type: {:?}", + dest + ))) + } + }; + + // Optimization - use a Store instead of writing individual fields + // when we're writing to the entire output register. + if dest.write_mask.is_all() { + let store = Statement::Store { + pointer: base_expr, + value: expr, + }; + self.func.body.push(store, Span::UNDEFINED); + } else { + for (i, mask) in [(0, Mask::X), (1, Mask::Y), (2, Mask::Z), (3, Mask::W)] { + if dest.write_mask.contains(mask) { + self.func.body.push( + Statement::Store { + pointer: self.func.expressions.append( + Expression::AccessIndex { + base: base_expr, + index: i, + }, + Span::UNDEFINED, + ), + value: expr, + }, + Span::UNDEFINED, + ); + } + } + } + Ok(()) + } + + /// 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.func.body.push(Statement::Emit(range), Span::UNDEFINED); + expr + } + + fn process_opcode( + &mut self, + opcode: &Opcode, + dest: &DestField, + source1: &SourceField, + source2: &Source2, + ) -> Result<()> { + match opcode { + // Copy the source register to the destination register + Opcode::Mov => { + // On the ActionScript side, the user might have specified something *other* than + // vec4f. In that case, we need to extend the source to a vec4f if we're writing to + // a vec4f register. + // FIXME - do we need to do this extension in other cases? + let do_extend = matches!( + dest.register_type, + RegisterType::Output | RegisterType::Varying + ); + let source = self.emit_source_field_load(source1, do_extend)?; + self.emit_dest_store(dest, source)?; + } + // Perform 'M * v', where M is a 4x4 matrix, and 'v' is a column vector. + Opcode::M44 => { + let source2 = match source2 { + Source2::SourceField(source2) => source2, + _ => unreachable!(), + }; + + // Read each row of the matrix + let source2_row0 = self.emit_source_field_load(source2, false)?; + let source2_row1 = self.emit_source_field_load( + &SourceField { + reg_num: source2.reg_num + 1, + ..source2.clone() + }, + false, + )?; + let source2_row2 = self.emit_source_field_load( + &SourceField { + reg_num: source2.reg_num + 2, + ..source2.clone() + }, + false, + )?; + let source2_row3 = self.emit_source_field_load( + &SourceField { + reg_num: source2.reg_num + 3, + ..source2.clone() + }, + false, + )?; + + // FIXME - The naga spv backend hits an 'unreachable!' + // if we don't create a Statement::Emit for each of these, + // even though validation passes. We should investigate this + // and report it upstream. + let matrix = self.evaluate_expr(Expression::Compose { + ty: self.matrix4x4f, + components: vec![source2_row0, source2_row1, source2_row2, source2_row3], + }); + + // Naga interprets each component of the matrix as a *column*. + // However, the matrix is stored in memory as a *row*, so we need + // to transpose it. + let matrix = self.evaluate_expr(Expression::Math { + fun: MathFunction::Transpose, + arg: matrix, + arg1: None, + arg2: None, + arg3: None, + }); + + let vector = self.emit_source_field_load(source1, true)?; + + let multiply = self.evaluate_expr(Expression::Binary { + op: BinaryOperator::Multiply, + left: matrix, + right: vector, + }); + + self.emit_dest_store(dest, multiply)?; + } + _ => { + return Err(Error::Unimplemented(format!( + "Unimplemented opcode: {:?}", + opcode + ))) + } + } + Ok(()) + } + + fn finish(mut self) -> Result { + // Load the 'main' output (a position or color) from our temporary location. + let dest_load = self.evaluate_expr(Expression::Load { pointer: self.dest }); + let mut components = vec![dest_load]; + + // If the vertex shader wrote to any varying registers, we need to + // return them as well. + if let ShaderType::Vertex = self.shader_config.shader_type { + for i in 0..self.varying_pointers.len() { + if self.varying_pointers[i].is_some() { + components.push(self.emit_varying_load(i)?); + } + } + } + + // We're consuming 'self', so just store store garbage here so that we can continue + // to use methods on 'self' + let return_ty = std::mem::replace( + &mut self.return_type, + Type { + name: None, + inner: TypeInner::Scalar { + kind: ScalarKind::Float, + width: 0, + }, + }, + ); + + // Finalize the return type, and do emit the actual return + let return_ty = self.module.types.insert(return_ty, Span::UNDEFINED); + self.func.result = Some(FunctionResult { + ty: return_ty, + binding: None, + }); + + let return_expr = self.evaluate_expr(Expression::Compose { + ty: return_ty, + components, + }); + + self.func.body.push( + Statement::Return { + value: Some(return_expr), + }, + Span::UNDEFINED, + ); + + let entry_point = EntryPoint { + name: ENTRY_POINT.to_string(), + stage: match self.shader_config.shader_type { + ShaderType::Vertex => ShaderStage::Vertex, + ShaderType::Fragment => ShaderStage::Fragment, + }, + early_depth_test: None, + workgroup_size: [0; 3], + function: self.func, + }; + + self.module.entry_points.push(entry_point); + Ok(self.module) + } +} diff --git a/render/naga-agal/src/lib.rs b/render/naga-agal/src/lib.rs new file mode 100644 index 000000000..30b7e9131 --- /dev/null +++ b/render/naga-agal/src/lib.rs @@ -0,0 +1,93 @@ +use naga::Module; + +mod builder; +mod types; + +use builder::NagaBuilder; + +const ENTRY_POINT: &str = "main"; + +pub const MAX_VERTEX_ATTRIBUTES: usize = 8; + +#[derive(Copy, Clone, Debug, Eq, PartialEq)] +pub enum VertexAttributeFormat { + Float1, + Float2, + Float3, + Float4, + Bytes4, +} + +#[derive(Debug)] +pub enum Error { + InvalidHeader, + InvalidShaderType(u8), + MissingVertexAttributeData(usize), + Unimplemented(String), + ReadError(std::io::Error), + InvalidOpcode(u32), +} + +impl From for Error { + fn from(err: std::io::Error) -> Self { + Error::ReadError(err) + } +} + +#[derive(Debug)] +pub enum ShaderType { + Vertex, + Fragment, +} + +/** + * Compiles an Adobe AGAL shader to a Naga Module. + * + * The `vertex_attributes` parameter is only used when compiling + * a vertex shader. + * + * The returning Module can be passed directly to `wgpu`, + * or compiled to a particular shader language using a `naga` backend. + * + * The shader entrypoint is always named `main`. + * + * We compile an AGAL shader as follows: + * + * # Vertex Shader + * + * * Vertex attributes - AGAL supports up to 8 vertex attributes, + * stored in `va0` to `va7`. You must provide the format of each attribute + * in the corresponding entry in the `vertex_attributes` array. + * Each attribute is mapped to the corresponding binding in the Naga shader + * - for example, va3 will have binding id 3. + * + * + * * Vertex output - An AGAL vertex shader has one main output (a vec4 position), + * and 8 varying outputs. The main output is mapped to the Naga 'Position' output, + * while each *used* varying register is mapped to a corresponding field in + * the Naga output struct. For example, if a vertex shader uses varying registers + * 2 and 5, then the Naga output struct type will have two members, with binding ids 2 and 5. + * If a shader does not write to a varying register, then it is not included in the + * Naga output struct type. + * + * * Program constants - An AGAL vertex shader has access to 128 program constants. + * These are mapped to a single Naga uniform buffer, with a binding id of 0. + * Each program constant is a vec4, and are stored in increasing order of register number. + * + * # Fragment Shader + * + * * Fragment input - An AGAL fragment shader can read from the 8 varying registers + * set by the fragment shader. Each *used* varying register is mapped to a corresponding + * binding in the Naga input type. For example, if a fragment shader uses varying registers + * 2 and 5, then the Naga input type will have two members, with binding ids 2 and 5. + * + * * Program constants - An AGAL fragment shader has access to 28 program constants. + * These are mapped to a single Naga uniform buffer, with a binding id of 1. + * + */ +pub fn agal_to_naga( + agal: &[u8], + vertex_attributes: &[Option; MAX_VERTEX_ATTRIBUTES], +) -> Result { + NagaBuilder::process_agal(agal, vertex_attributes) +} diff --git a/render/naga-agal/src/types.rs b/render/naga-agal/src/types.rs new file mode 100644 index 000000000..a7cf21eda --- /dev/null +++ b/render/naga-agal/src/types.rs @@ -0,0 +1,195 @@ +use crate::Error; +use num_derive::FromPrimitive; +use num_traits::FromPrimitive; + +#[derive(num_derive::FromPrimitive, Debug)] +pub enum Opcode { + Mov = 0x00, + Add = 0x01, + Sub = 0x02, + Mul = 0x03, + Div = 0x04, + Rcp = 0x05, + Min = 0x06, + Max = 0x07, + Frc = 0x08, + Sqt = 0x09, + Rsq = 0x0a, + Pow = 0x0b, + Log = 0x0c, + Exp = 0x0d, + Nrm = 0x0e, + Sin = 0x0f, + Cos = 0x10, + Crs = 0x11, + Dp3 = 0x12, + Dp4 = 0x13, + Abs = 0x14, + Neg = 0x15, + Sat = 0x16, + M33 = 0x17, + M44 = 0x18, + M34 = 0x19, + Kil = 0x27, + Tex = 0x28, + Sge = 0x29, + Slt = 0x2a, + Seq = 0x2b, + Sne = 0x2d, + Ddx = 0x1a, + Ddy = 0x1b, + Ife = 0x1c, + Ine = 0x1d, + Ifg = 0x1e, + Ifl = 0x1f, + Els = 0x20, + Eif = 0x21, +} + +#[derive(FromPrimitive, Debug, Clone, PartialEq, Eq)] +pub enum RegisterType { + Attribute = 0, + Constant = 1, + Temporary = 2, + Output = 3, + Varying = 4, + Sampler = 5, + FragmentRegister = 6, +} + +#[derive(Debug, FromPrimitive, Clone)] +pub enum DirectMode { + Direct = 0, + Indirect = 1, +} + +#[derive(Debug)] +pub struct DestField { + pub register_type: RegisterType, + pub write_mask: Mask, + pub reg_num: u16, +} + +impl DestField { + pub fn parse(val: u32) -> Result { + let reg_num = (val & 0xFFFF) as u16; + let write_mask = Mask::from_bits(((val >> 16) & 0xF) as u8).unwrap(); + let reg_type = RegisterType::from_u16(((val >> 24) & 0xF) as u16).unwrap(); + Ok(DestField { + register_type: reg_type, + write_mask, + reg_num, + }) + } +} + +#[derive(Debug, Clone)] +#[allow(dead_code)] +pub struct SourceField { + pub direct_mode: DirectMode, + pub index_select: u8, + pub index_type: RegisterType, + pub register_type: RegisterType, + pub swizzle: u8, + pub indirect_offset: u8, + pub reg_num: u16, +} + +bitflags::bitflags! { + pub struct Mask: u8 { + const X = 0b0001; + const Y = 0b0010; + const Z = 0b0100; + const W = 0b1000; + } +} + +impl SourceField { + pub fn parse(val: u64) -> Result { + // FIXME - check that all the other bits are 0 + let reg_num = (val & 0xFFFF) as u16; + let indirect_offset = ((val >> 16) & 0xFF) as u8; + let swizzle = ((val >> 24) & 0xFF) as u8; + let register_type = RegisterType::from_u16(((val >> 32) & 0xF) as u16).unwrap(); + let index_type = RegisterType::from_u16(((val >> 40) & 0xF) as u16).unwrap(); + let index_select = ((val >> 48) & 0x3) as u8; + let direct_mode = DirectMode::from_u16(((val >> 63) & 0x1) as u16).unwrap(); + Ok(SourceField { + direct_mode, + index_select, + index_type, + register_type, + swizzle, + indirect_offset, + reg_num, + }) + } +} + +#[derive(FromPrimitive)] +pub enum Filter { + Nearest = 0, + Linear = 1, +} + +#[derive(FromPrimitive)] +pub enum Mipmap { + Disable = 0, + Nearest = 1, + Linear = 2, +} + +#[derive(FromPrimitive)] +pub enum Wrapping { + Clamp = 0, + Repeat = 1, +} + +#[derive(FromPrimitive)] +pub enum Dimension { + TwoD = 0, + Cube = 1, +} + +#[allow(dead_code)] +pub struct SamplerField { + pub filter: Filter, + pub mipmap: Mipmap, + pub wrapping: Wrapping, + pub dimension: Dimension, + /// Texture level-of-detail (LOD) bias + pub texture_lod_bias: i8, + pub reg_num: u16, + pub reg_type: RegisterType, +} + +impl SamplerField { + pub fn parse(val: u64) -> Result { + let reg_num = (val & 0xFFFF) as u16; + let load_bias = ((val >> 16) & 0xFF) as i8; + let reg_type = RegisterType::from_u64((val >> 32) & 0xF).unwrap(); + let dimension = Dimension::from_u64((val >> 44) & 0xF).unwrap(); + + // FIXME - check that the actual field is 0 + let _special = 0; + + let wrapping = Wrapping::from_u64((val >> 52) & 0xF).unwrap(); + let mipmap = Mipmap::from_u64((val >> 56) & 0xF).unwrap(); + let filter = Filter::from_u64((val >> 60) & 0xF).unwrap(); + + Ok(SamplerField { + filter, + mipmap, + wrapping, + dimension, + texture_lod_bias: load_bias, + reg_num, + reg_type, + }) + } +} + +pub enum Source2 { + SourceField(SourceField), + Sampler(SamplerField), +} diff --git a/render/naga-agal/tests/snapshots/wgsl__shaders-2.snap b/render/naga-agal/tests/snapshots/wgsl__shaders-2.snap new file mode 100644 index 000000000..84a82e23e --- /dev/null +++ b/render/naga-agal/tests/snapshots/wgsl__shaders-2.snap @@ -0,0 +1,25 @@ +--- +source: tests/wgsl.rs +assertion_line: 53 +expression: output +--- +struct VertexOutput { + @builtin(position) member: vec4, + @location(0) varying_0_: vec4, +} + +@group(0) @binding(0) +var constant_registers: array,128u>; + +@vertex +fn main(@location(0) param: vec4, @location(1) param_1: vec4) -> VertexOutput { + var dest_temp: vec4; + var varying_0_: vec4; + + dest_temp = param; + varying_0_ = param_1; + let _e5: vec4 = dest_temp; + let _e6: vec4 = varying_0_; + return VertexOutput(_e5, _e6); +} + diff --git a/render/naga-agal/tests/snapshots/wgsl__shaders-3.snap b/render/naga-agal/tests/snapshots/wgsl__shaders-3.snap new file mode 100644 index 000000000..3505e6c62 --- /dev/null +++ b/render/naga-agal/tests/snapshots/wgsl__shaders-3.snap @@ -0,0 +1,21 @@ +--- +source: tests/wgsl.rs +assertion_line: 70 +expression: output +--- +struct FragmentOutput { + @location(0) member: vec4, +} + +@group(0) @binding(1) +var constant_registers: array,28u>; + +@fragment +fn main(@location(0) param: vec4) -> FragmentOutput { + var dest_temp: vec4; + + dest_temp = param; + let _e3: vec4 = dest_temp; + return FragmentOutput(_e3); +} + diff --git a/render/naga-agal/tests/snapshots/wgsl__shaders.snap b/render/naga-agal/tests/snapshots/wgsl__shaders.snap new file mode 100644 index 000000000..aed3cd053 --- /dev/null +++ b/render/naga-agal/tests/snapshots/wgsl__shaders.snap @@ -0,0 +1,29 @@ +--- +source: tests/wgsl.rs +assertion_line: 35 +expression: output +--- +struct VertexOutput { + @builtin(position) member: vec4, + @location(0) varying_0_: vec4, +} + +@group(0) @binding(0) +var constant_registers: array,128u>; + +@vertex +fn main(@location(0) param: vec3, @location(1) param_1: vec3) -> VertexOutput { + var dest_temp: vec4; + var varying_0_: vec4; + + let _e4: vec4 = constant_registers[0u]; + let _e7: vec4 = constant_registers[1u]; + let _e10: vec4 = constant_registers[2u]; + let _e13: vec4 = constant_registers[3u]; + dest_temp = (transpose(mat4x4(_e4, _e7, _e10, _e13)) * vec4(param.x, param.y, param.z, 1.0)); + varying_0_ = vec4(param_1.x, param_1.y, param_1.z, 1.0); + let _e30: vec4 = dest_temp; + let _e31: vec4 = varying_0_; + return VertexOutput(_e30, _e31); +} + diff --git a/render/naga-agal/tests/wgsl.rs b/render/naga-agal/tests/wgsl.rs new file mode 100644 index 000000000..79aa66e80 --- /dev/null +++ b/render/naga-agal/tests/wgsl.rs @@ -0,0 +1,83 @@ +use naga::{ + valid::{Capabilities, ValidationFlags, Validator}, + Module, +}; +use naga_agal::{agal_to_naga, VertexAttributeFormat}; + +pub fn to_wgsl(module: &Module) -> String { + let mut out = String::new(); + + let mut validator = Validator::new(ValidationFlags::all(), 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 +} + +// Making this a macro gives us a better span in 'inta' +macro_rules! test_shader { + ($shader:expr, $attrs:expr, $shader_type:expr $(,)?) => { + let module = agal_to_naga(&$shader, $attrs).unwrap(); + let output = to_wgsl(&module); + insta::assert_display_snapshot!(output); + }; +} + +#[test] +fn test_shaders() { + test_shader!( + // m44 op, va0, vc0 + // mov v0, va1 + [ + 160, 1, 0, 0, 0, 161, 0, 24, 0, 0, 0, 0, 0, 15, 3, 0, 0, 0, 228, 0, 0, 0, 0, 0, 0, 0, + 228, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 15, 4, 1, 0, 0, 228, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0 + ], + &[ + Some(VertexAttributeFormat::Float3), + Some(VertexAttributeFormat::Float3), + None, + None, + None, + None, + None, + None, + ], + ShaderType::Vertex, + ); + + test_shader!( + // mov op, va0 + // mov v0, va1 + [ + 160, 1, 0, 0, 0, 161, 0, 0, 0, 0, 0, 0, 0, 15, 3, 0, 0, 0, 228, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 15, 4, 1, 0, 0, 228, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 + ], + &[ + Some(VertexAttributeFormat::Float4), + Some(VertexAttributeFormat::Float4), + None, + None, + None, + None, + None, + None, + ], + ShaderType::Vertex, + ); + + test_shader!( + // mov oc, v0 + [ + 160, 1, 0, 0, 0, 161, 1, 0, 0, 0, 0, 0, 0, 15, 3, 0, 0, 0, 228, 4, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0 + ], + &[None, None, None, None, None, None, None, None], + ShaderType::Fragment, + ); +} diff --git a/tests/tests/regression_tests.rs b/tests/tests/regression_tests.rs index dbb4605a3..8b2b8b02f 100644 --- a/tests/tests/regression_tests.rs +++ b/tests/tests/regression_tests.rs @@ -150,6 +150,7 @@ swf_tests! { (as2_super_and_this_v8, "avm1/as2_super_and_this_v8", 1), (as2_super_via_manual_prototype, "avm1/as2_super_via_manual_prototype", 1), (as3_add, "avm2/add", 1), + (as3_agal_compiler, "avm2/agal_compiler", 1), (as3_application_domain, "avm2/application_domain", 1), (as3_array_access, "avm2/array_access", 1), (as3_array_concat, "avm2/array_concat", 1), diff --git a/tests/tests/swfs/avm2/agal_compiler/AGALMiniAssembler.as b/tests/tests/swfs/avm2/agal_compiler/AGALMiniAssembler.as new file mode 100644 index 000000000..4c1d68b75 --- /dev/null +++ b/tests/tests/swfs/avm2/agal_compiler/AGALMiniAssembler.as @@ -0,0 +1,819 @@ +/* +Copyright (c) 2015, Adobe Systems Incorporated +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are +met: + +* Redistributions of source code must retain the above copyright notice, +this list of conditions and the following disclaimer. + +* Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +* Neither the name of Adobe Systems Incorporated nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, +THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR +PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR +CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF +LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING +NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*/ +package +{ + // =========================================================================== + // Imports + // --------------------------------------------------------------------------- + import flash.display3D.*; + import flash.utils.*; + + // =========================================================================== + // Class + // --------------------------------------------------------------------------- + public class AGALMiniAssembler + { // ====================================================================== + // Constants + // ---------------------------------------------------------------------- + protected static const REGEXP_OUTER_SPACES:RegExp = /^\s+|\s+$/g; + + // ====================================================================== + // Properties + // ---------------------------------------------------------------------- + // AGAL bytes and error buffer + private var _agalcode:ByteArray = null; + private var _error:String = ""; + + private var debugEnabled:Boolean = false; + + private static var initialized:Boolean = false; + public var verbose:Boolean = false; + + // ====================================================================== + // Getters + // ---------------------------------------------------------------------- + public function get error():String { return _error; } + public function get agalcode():ByteArray { return _agalcode; } + + // ====================================================================== + // Constructor + // ---------------------------------------------------------------------- + public function AGALMiniAssembler( debugging:Boolean = false ):void + { + debugEnabled = debugging; + if ( !initialized ) + init(); + } + // ====================================================================== + // Methods + // ---------------------------------------------------------------------- + + public function assemble2( ctx3d : Context3D, version:uint, vertexsrc:String, fragmentsrc:String ) : Program3D + { + var agalvertex : ByteArray = assemble ( VERTEX, vertexsrc, version ); + var agalfragment : ByteArray = assemble ( FRAGMENT, fragmentsrc, version ); + var prog : Program3D = ctx3d.createProgram(); + prog.upload(agalvertex,agalfragment); + return prog; + } + + public function assemble( mode:String, source:String, version:uint=1, ignorelimits:Boolean=false ):ByteArray + { + var start:uint = getTimer(); + + _agalcode = new ByteArray(); + _error = ""; + + var isFrag:Boolean = false; + + if ( mode == FRAGMENT ) + isFrag = true; + else if ( mode != VERTEX ) + _error = 'ERROR: mode needs to be "' + FRAGMENT + '" or "' + VERTEX + '" but is "' + mode + '".'; + + agalcode.endian = Endian.LITTLE_ENDIAN; + agalcode.writeByte( 0xa0 ); // tag version + agalcode.writeUnsignedInt( version ); // AGAL version, big endian, bit pattern will be 0x01000000 + agalcode.writeByte( 0xa1 ); // tag program id + agalcode.writeByte( isFrag ? 1 : 0 ); // vertex or fragment + + initregmap(version, ignorelimits); + + var lines:Array = source.replace( /[\f\n\r\v]+/g, "\n" ).split( "\n" ); + var nest:int = 0; + var nops:int = 0; + var i:int; + var lng:int = lines.length; + + for ( i = 0; i < lng && _error == ""; i++ ) + { + var line:String = new String( lines[i] ); + line = line.replace( REGEXP_OUTER_SPACES, "" ); + + // remove comments + var startcomment:int = line.search( "//" ); + if ( startcomment != -1 ) + line = line.slice( 0, startcomment ); + + // grab options + var optsi:int = line.search( /<.*>/g ); + var opts:Array; + if ( optsi != -1 ) + { + opts = line.slice( optsi ).match( /([\w\.\-\+]+)/gi ); + line = line.slice( 0, optsi ); + } + + // find opcode + var opCode:Array = line.match( /^\w{3}/ig ); + if ( !opCode ) + { + if ( line.length >= 3 ) + trace( "warning: bad line "+i+": "+lines[i] ); + continue; + } + var opFound:OpCode = OPMAP[ opCode[0] ]; + + // if debug is enabled, output the opcodes + if ( debugEnabled ) + trace( opFound ); + + if ( opFound == null ) + { + if ( line.length >= 3 ) + trace( "warning: bad line "+i+": "+lines[i] ); + continue; + } + + line = line.slice( line.search( opFound.name ) + opFound.name.length ); + + if ( ( opFound.flags & OP_VERSION2 ) && version<2 ) + { + _error = "error: opcode requires version 2."; + break; + } + + if ( ( opFound.flags & OP_VERT_ONLY ) && isFrag ) + { + _error = "error: opcode is only allowed in vertex programs."; + break; + } + + if ( ( opFound.flags & OP_FRAG_ONLY ) && !isFrag ) + { + _error = "error: opcode is only allowed in fragment programs."; + break; + } + if ( verbose ) + trace( "emit opcode=" + opFound ); + + agalcode.writeUnsignedInt( opFound.emitCode ); + nops++; + + if ( nops > MAX_OPCODES ) + { + _error = "error: too many opcodes. maximum is "+MAX_OPCODES+"."; + break; + } + + // get operands, use regexp + var regs:Array; + + // will match both syntax + regs = line.match( /vc\[([vofi][acostdip]?[d]?)(\d*)?((\.[xyzw])?(\+\d{1,3})?)?\](\.[xyzw]{1,4})?|([vofi][acostdip]?[d]?)(\d*)?(\.[xyzw]{1,4})?/gi ); + + if ( !regs || regs.length != opFound.numRegister ) + { + _error = "error: wrong number of operands. found "+regs.length+" but expected "+opFound.numRegister+"."; + break; + } + + var badreg:Boolean = false; + var pad:uint = 64 + 64 + 32; + var regLength:uint = regs.length; + + for ( var j:int = 0; j < regLength; j++ ) + { + var isRelative:Boolean = false; + var relreg:Array = regs[ j ].match( /\[.*\]/ig ); + if ( relreg && relreg.length > 0 ) + { + regs[ j ] = regs[ j ].replace( relreg[ 0 ], "0" ); + + if ( verbose ) + trace( "IS REL" ); + isRelative = true; + } + + var res:Array = regs[j].match( /^\b[A-Za-z]{1,3}/ig ); + if ( !res ) + { + _error = "error: could not parse operand "+j+" ("+regs[j]+")."; + badreg = true; + break; + } + var regFound:Register = REGMAP[ res[ 0 ] ]; + + // if debug is enabled, output the registers + if ( debugEnabled ) + trace( regFound ); + + if ( regFound == null ) + { + _error = "error: could not find register name for operand "+j+" ("+regs[j]+")."; + badreg = true; + break; + } + + if ( isFrag ) + { + if ( !( regFound.flags & REG_FRAG ) ) + { + _error = "error: register operand "+j+" ("+regs[j]+") only allowed in vertex programs."; + badreg = true; + break; + } + if ( isRelative ) + { + _error = "error: register operand "+j+" ("+regs[j]+") relative adressing not allowed in fragment programs."; + badreg = true; + break; + } + } + else + { + if ( !( regFound.flags & REG_VERT ) ) + { + _error = "error: register operand "+j+" ("+regs[j]+") only allowed in fragment programs."; + badreg = true; + break; + } + } + + regs[j] = regs[j].slice( regs[j].search( regFound.name ) + regFound.name.length ); + //trace( "REGNUM: " +regs[j] ); + var idxmatch:Array = isRelative ? relreg[0].match( /\d+/ ) : regs[j].match( /\d+/ ); + var regidx:uint = 0; + + if ( idxmatch ) + regidx = uint( idxmatch[0] ); + + if ( regFound.range < regidx ) + { + _error = "error: register operand "+j+" ("+regs[j]+") index exceeds limit of "+(regFound.range+1)+"."; + badreg = true; + break; + } + + var regmask:uint = 0; + var maskmatch:Array = regs[j].match( /(\.[xyzw]{1,4})/ ); + var isDest:Boolean = ( j == 0 && !( opFound.flags & OP_NO_DEST ) ); + var isSampler:Boolean = ( j == 2 && ( opFound.flags & OP_SPECIAL_TEX ) ); + var reltype:uint = 0; + var relsel:uint = 0; + var reloffset:int = 0; + + if ( isDest && isRelative ) + { + _error = "error: relative can not be destination"; + badreg = true; + break; + } + + if ( maskmatch ) + { + regmask = 0; + var cv:uint; + var maskLength:uint = maskmatch[0].length; + for ( var k:int = 1; k < maskLength; k++ ) + { + cv = maskmatch[0].charCodeAt(k) - "x".charCodeAt(0); + if ( cv > 2 ) + cv = 3; + if ( isDest ) + regmask |= 1 << cv; + else + regmask |= cv << ( ( k - 1 ) << 1 ); + } + if ( !isDest ) + for ( ; k <= 4; k++ ) + regmask |= cv << ( ( k - 1 ) << 1 ); // repeat last + } + else + { + regmask = isDest ? 0xf : 0xe4; // id swizzle or mask + } + + if ( isRelative ) + { + var relname:Array = relreg[0].match( /[A-Za-z]{1,3}/ig ); + var regFoundRel:Register = REGMAP[ relname[0]]; + if ( regFoundRel == null ) + { + _error = "error: bad index register"; + badreg = true; + break; + } + reltype = regFoundRel.emitCode; + var selmatch:Array = relreg[0].match( /(\.[xyzw]{1,1})/ ); + if ( selmatch.length==0 ) + { + _error = "error: bad index register select"; + badreg = true; + break; + } + relsel = selmatch[0].charCodeAt(1) - "x".charCodeAt(0); + if ( relsel > 2 ) + relsel = 3; + var relofs:Array = relreg[0].match( /\+\d{1,3}/ig ); + if ( relofs.length > 0 ) + reloffset = relofs[0]; + if ( reloffset < 0 || reloffset > 255 ) + { + _error = "error: index offset "+reloffset+" out of bounds. [0..255]"; + badreg = true; + break; + } + if ( verbose ) + trace( "RELATIVE: type="+reltype+"=="+relname[0]+" sel="+relsel+"=="+selmatch[0]+" idx="+regidx+" offset="+reloffset ); + } + + if ( verbose ) + trace( " emit argcode="+regFound+"["+regidx+"]["+regmask+"]" ); + if ( isDest ) + { + agalcode.writeShort( regidx ); + agalcode.writeByte( regmask ); + agalcode.writeByte( regFound.emitCode ); + pad -= 32; + } else + { + if ( isSampler ) + { + if ( verbose ) + trace( " emit sampler" ); + var samplerbits:uint = 5; // type 5 + var optsLength:uint = opts == null ? 0 : opts.length; + var bias:Number = 0; + for ( k = 0; kFPIg-E!-x`^WE{=bi6%&gcEiJm>p7=fEt1%wPb39RL7ls5#cBOhgL+ z002hQZvpULcyz!SA{s>``gnPu0=x*mNYiYP`>-8BSjwkc|tX1XK={f2JJ`#f{W@*K4)BKDVfG^IY>X z4WpCzKmZ!`5z#gM1ZmGlxY|WvAEE|5Ia%M?mO+Z{7t_T&Grk{j&<7^vND*t^J0Iy* zDbOCu@+1c9nU&AdE5&bo;=tjK6ikvmtD>)$XG%EmfGk!~?a3X~yxNYONp~{OqI(5@ zq?Sn)%qv%m!;x(KW|DTbEmT_hZ0We;hFEAMe;KDKbMY&fvQXKB!++S7RGg_=;^V{i z)XIeqVdS};5&e(dDj*;zsnuD{8fT|$3A$YG;+21M;g#&$hjFRUG@{*&%ic*!dx0R# zxbe(tk#vr1p&@6Dt8`;r*x1rv6Hr`0^i(z~FdCC5LUMZSEp&}~16ec3A!qem_k9Pl zud42W_I`l}Gq0#_FNK0XN1)F+L3b^UQ-vH{5{hd?6Q!Tu>m~c?#(13>eoEevPwk7< zUr#D|LW(K@yYtE4WBKB#r+V8bEj^02<36r%XV*qnhUld1AZOE&cs)kL zi{x9e66UByOP`i*3o zLc=-MqxYGMbV9$pUxE@T-cp*=ALN9oUduC%;lRG}3sqD8ZE zYQcI++?N+Nn_*m7MH#p#BpZxgs4{3gdE3EP2VV5`(B;WT5*nM#Z07-6;$`~hg`z9# zd+nv4eHSnNy+V1S0#FRGeWPqCldm-EI4hlAFAM&VL#jiE76Wvf)6rTc=xCnq$Z+i`Ob=xxGTEXHbSDEl$FwR_j;&L zh;-j9BWkP{Im{U%>%M&NPB#7?Wa zTY+{(wuQBT3f@IX=vSXRwVXiVHog`!o<_N-@Nmd z&b#0vt`H-_0=o~sur90KKc<5mZ>91dLXGs5D8Sg{v(&$ReW3wsKwnD$1k z=+*7E{urb0vQL?<^U`9Bh8aqpHED=vN_|Rt2y>@Kyd!vxoO#jpZiC2Wj6o~)x{NiI z!w(bN_v#RH?t@lP)auG#(92c{NFC;e2@- zKne4sLgGxW=%pW8#wYa6JFVC1!qyXqgZ>sD7Y zj}#vJQx5xwewls6ne5Jr7xI_xj?SwJefVhHu!#&*$XT>1Tn#H)EX^+O56`L_&gNLX zQ5L~f0-qPOIJt0QzOI069)Azwsy{DW(t5r~WvoMZg(qeKQEMUrNg*%QL@TficE+=Y zQ`d^u2bW7S3U3MIqS(i{GY`by9aJ;w=C(HE*G7(TniU$S^Xgcb;X*Z1$))fmQzq;g*mIkv&wN>*4XQv8CCVv`z+`*oC`~@fGBFSnF zV^(mQ%3|X!h7imfJYeJGE@}LNBCVTT#rEX zM#Ma^QC)zgQRZo0_tl9#qa4=aX@(Yy?g|!mn}xMqKT}!y#-9=g!X)8b94Xp`@(GnbhC_#Cb_w+aK&z_^>!eh) z)ZBLK<7h}YqQk{(CsNuw1*BecSHQYxzD<{4>vxB{ub8}CiX<$?s zFzWTlA*_<`XPL$}Ve)cm2DW!D@F`BRPM>Q9o*((-?Fo)JF_^)r(`$#}6d+i^2-rCn zD&yie++i2bAB`E-NmVh1roZEz4aweXChNKG(%k9C*4bvI@k~X-+vYY&#fh+{;89nh z?&i@Lo4LWCfF4&hk`B^KWwm>hG1+ukJdH`C;GCNKGmdpW6HvmvQ|jbtq?YSZ@Rc)= zfrjGPDtBw0PMtu*On&!@BNCvu%5I2*r@hCL7CEPGfF3Iw!+WAhO0$PwLYWi$wSo1A z+PN-`zje)dSiJiAhGUGyruhdjBk^!rKnb{gyp%mO>NbCO{`QlKg6GA6lx$4N9u}Gw)kB2wFS( z@{J_K*_f=6^l%i{2yC$=e;2zvdQ@rlMz=w2`)8fxDoqX3LH&lg zYC*i2OUvx4MDMLa+^f~#^#@QREoVyXYiUcrJ3iU7jZ}~R4NH}c{@w-e{L)E{zg|ev z&a<*tvXxz6+vRjWVJO93Q8(;V7wPO=KdJIyp{gJ`2^Hse?~zh8I^f)=yvHk2V_fR~ zHX$Mf%`fKDq{)^-WWL)_8@|GG8Ho*#3Ed63V_AJGrW+%0ktX!B9rY`k6M81lsQ}p! zzch+KSgx;98?WIt2r;jDE(K7LaWG0J*{G6uJhVTs+t9naCcNoWcgAKPS6dyp61U@R zy-RNB9pD4q@ZsK)%#x;=bl?GT%P8RyC9P++%axz4*3hv75`?p{9Vf4hbsZS{NE z-6nL00bn%4X1II#D#6d-jnG6i-QatmJaA|%{q`HSod9M4gbv{bDATVQ)mTP%B|jMd z|06f-U;+cSgU8RwVQUcMTbJJ7bA;uOtZmd*-q^b4rXPIiy&Wfh&te3MtzxjQ()+tZ zgZ;-aM!?uA2Cbgnzla-KQ-2pY{^8RP^Ys2Hetet#I~%!8wgdia%RidRALsvUE4P=3 zq%ZLcllfnJUUaNIZa@C*KMQV~zd+~zLVn*i@Vjujy@2|j?WhW~WMO4I*+u^( M(F0yE{TTr8KlO0pqW}N^ literal 0 HcmV?d00001 diff --git a/tests/tests/swfs/avm2/agal_compiler/test.swf b/tests/tests/swfs/avm2/agal_compiler/test.swf new file mode 100644 index 0000000000000000000000000000000000000000..efb7ac6782ebcf74ada639ea50f10b5887bf6de5 GIT binary patch literal 10580 zcmV-aDXZ2)S5qsjL;wJIoQ+%ucvMvue&^m-W>P01384>R0t6l@fDJ_mDHu$_Bp^ug zGV^98c_T?C&P)PWBUr$Kpi*r?MM3Ok*R`*`tp>2{+IKB$-FdcEcim;rc~c0BF3ab6 z_ul`Md(SQBo}1)zr69ElLgs)Vnd4M6>dJR zueWC5DaR~ouI_HED-|BkO+{G{n!NsQzq3pb9?vPQBmn<8gkm8J#D1qO&5@HonmOn8 zLqF>m`_cR-T#%b0p+I*>YI~fKP%+- z_=pvTK&Hj;^>+0%86KZMWVhP%5&MDx-`Mi`<+XKz?m&4s?C;V${h_201S6vkXZu5; zV5p?SANd7UPlu<|8}#{!t;`DsJN=&S35SavhVZ8O^_elhe1ULJr)SlSO8df4up{K@ z!dgYJJK|p%nNcbB2D*P~OrCgcI$K;Z^4HnCslLkfza0+wJGwkQQx6j=_Jm!>l+9}B zsUY|KU4ckN+|S8JkVUsT!X!B92!h?7F3Vu?k>q;28vPxCaKvvdL*Ws`W5tm=j|5}8 zYuc~LCA%Ez`;nzPgWVnVamR{|D4m#sDMyq{&gzjxL!R!A|Fi!@<;aqW`TEt>Z}N2Y zbjAnzKetIju*(x({(lK3Ch1p~G$9x-k7zLAOCbh~NE;4#R-Gr%-5luhcLus~Yb4Ni zJdOUcsYj3+ii#9x&7R#0pwI>6BEm2E3c$$a*$^^)QH`SifswRd88d?Z%Z5W?c zfTuHXrr#%0U20fZSH3W%GT@B_g593bDjAg#u{V|1Ev&7=T54VSlD7J)rskUZ`5c42 z)oxf=(NI~{WUFp0pI=v1-)zI4JG%Vc5qWV{V{_FK8FvL<5ZN&+jcv_K7gn`3om5lZ zoRxqoYwBW1ZUSjqSXEI|UK>Mm5@=)9!m9G-{{ri3XxRS+R@c;84ijicH`-2ZtZA;Y zQLVA6ywa-BfUOFR)82q4D#~li8#AaD?;eMrHKbWr-rQKTBn=VLZEL8nT}lnow;Czf zsBdV)*=SCU;Y~FS_0tKisjr}kqiSW9RpoUJi+NYDkC#_ga*xl)n-<+)|D6%6ck`yIQ+U{a1`ng%=33P2^s&X|tk1*xylinZ>x;0( z4XiK77T2)80IRNG?Os;h%-SQYx{0-iS#<+z53=e?*6!mqHI+ONKm$#zFU%^Zv*|uo zIfKpci4}|HRf*nS-JgQHtiGzek(RJ4(9`Yrgz(ha5v{Eu+>Rg}J5NX<$%gtWg~Wp0 zeyP3z5ve$gn~T*@AQ>Jc{2~; z%{+v6>><2kS7r?1l^wk#V+gbC=#{n#Jg#@zlohh2n~kr$?%&MRCd7iDHiUfnC&Qfo+Q& zG3mD6HjCfZ=ZG;DC64xDM~u=Ap@jXu z%!6oo?ZT7FGY_&nojry}tX$GuCi+%J%BXe})sB)D*Hl(DNPPjHKUi|wtQLx(;rKVa zqCMQ=>Fb)_Libl2Sqq@M#~)9(B}Q-2=ZScDus0%EslBZy7!FXpbkGgDs5^|40L`tVc zx;aJK7?2`i#K-U)33k4UI-Hq^x$!ofo{J~b8SGJ=eV$VMcuSr5 zFC`V@jkAvcPnm=2aS4l=VgX(RN;4N*)}7&)xmlU9-jcUiR&~6;q(oL?1P5~kttCA5 zY$O+Sobdln@-N+&b}Pr|z#S=hBP(ZEX+{`l+#VZLIOI)DbdL2DvKZr}Q8mF5d+^1Q z?!{G#_~VZkxwoe?=<$hNls<%my&aLn3AOaKPF3)H zbmD$Uv&M)MNIpuzo=7+#(pL^z`H>RrsSfsb`$SB(!m{r1Vaa#Z8Cy0)u&!7McX+v-J~2ihTuIkhARNy} zQKU1ns>d=H_IKhg=?u2Fhy4+`Wm)l*mR2a0Ynq&mRkaSSMPF7v?KID{Gh3}`7#DZJ z8|>}$IlF_n8lF(tpK!s+oRe8map_@F?Ez$l5_vbJ=ICoEsg@sCNJ=Ca1w}{PIY$?v zT5NipwDH4qsac)=&&}gH-16ZXIZ(j7f(Ox zsMbu}D&@^Ji>vVX;kKI1X3v($Bz|#2&bGx3%+7Q=tt0CRby#PXwJuAsqF<8{3`J7o zDTj_RIfa5-Y=T5R;0b5MHRB>SWLsy&PFET>W;x}0>pZ0pKpL~0@oEICsaW?!VgU@C zicjyT#j16O#5$2IM~8LzyK&#qq0=!k(->zdRKthk_+PrSVl`(3FNax92s`3?(-Xmm zCZ1t`{6fWp(2W7SJCYEH_`7<7A$)tq?{}Kmp5BPkhY;@G}bvJV*XCY1(ke zFO$lohqfN7UiEhylG^abQ81P%tlVjd!I{+1i5maI2^!k%K>JW z^SA=`K~R7xOvI`J=?c$Mc(ww=6qu$!u_Bcya;XAkify`Lo1xfdDsYqnM=NlQ0>>(F zoMM}$*iTSWPgGLpD5-Om)N%#pDNvz6r2g&11r*!qify@K>r`xAimh9*1r=M5Vmm{zg%n#@u|*VHuVU*{Y%3JoO2xKH zv7M>d&QfgWD7Nzy+xei_)*xMobTQI8q)U-5N4f%OBhpn!n~=63Z9}>S={lqvkai&5 zgtQasR;1lXdywuxx(n$Zr2CK_K>97x?~oowdKBq#q~9Ywh4c*4b4Zzr?L{nxE4Dvi z@hVcTV%vwsYe;V({SoOcq<4_sL;3*eFGwFDeS-8U(&tEDBK-}?MEVBlJEXrO{fP7r zq@R)gjnvN++iIi>kS;=6i*yOnWk~ChHXvPzbT!guq^(HXnF2Q;4Iu4c_8S@Qpq)&~ zyA9iW*eI|c2livYzLSA{8XJXs>=|S{i+gMk>2<8Xf%GQQACdlqifh*WhUS*@26kZ*o4?`ZIE~h7O8^3(gfgQ##%^8$rL470m$jk#eX6LP{h2o z40uY)aQw%(BIfsJW?I_H2&9o%9fjp+K@>nylrdHx0*sCIpp3)9i8LPT1xOPF2}gv2 zu9%lJ=NHpNd0w+sF79XY@9IY@FNsZNIo82x7oTZ?DselgH?TTQxr;{&f87KmP zgA`Lz9M(6B2}(8-Y{NiE$pIlX7lgDtyFEq7whzZ{g?##qAj>1EHi`)uqnVI7h6#?b zOvoAs!Z4@Zo+gY;B|n4$lAJ&v7c~pxI=K^>kcT?MQ8#}Q6GouksL6JFhTv3+NM;Is zrjnk@gfY{YFczi9q15=|5?i)V2&I-qDOe`MbYX_TaOg9$xO|k%j~0%R#iNB|W$9?) zI9Z;>`0=tbi}Mp?8~9c{p-w<$UDU$s?3>UN4VSGI)bL zO3CC4<CSPQU3cEKe8 zOXWhi4B#|*BCH2kCQpJ504?%lxDudME`qB8Go1<>VU zxE{bOm%spkPcDTU0sL|q+zikzPlsCoI^-F!3&44M0#n9v%efkxzh!0M3wS!y^D8`9ydOAS};;CjcVyTzC?o zS1yOA0s7>5@GQUzxdNUCSSeS+3jnL+DtHOtOt~6f1~^Ne53c~6EuRE~0O!axuovK5 zc>(MPI8QzqUI&~uxfb38I3MIX_!Hp#$@TCyz#5Pn;9bCplo!JL02hLM3j7&xM&(BM z5a42vo8V)>sg;}IuK?>nUId>3&au20z5uus22Hyf) z0rE2V9&q~Q7We^RBgn1r6W}_?8vGOBDv;aYUw{iDyWs%9CXhXFKEsufby&k-3&>u$ zkl_-^KDd~{Hjw?Wj^SF#?QkiBYe4RR%NZ`3Y`_%^t^+v$8yT*id^%jk;0BPF!zPBy zD0jjZ20K9Rf^7`fQ|^Xq7^Dk1!e~5bTm+ja2nAvq2AMojU`{;ylSM(C2*n~ulkm86 zkq?&6$D_qWQCjr%wak!@8u46IQ=zX1DGJ^>(&@pkX)mkTqQr zv;KVGfK9TAW)89LV1xUrQ(!>7lf^`G(IZS;MCp1)y$+{dy%Lg)^%@}xcVVQ7!rcs< zSF?qR%Gy_c!?F>~aAdcRc}Gq9MgeqDWlnd27&dhVjWe#5LD3N7#F;UJ^%TP994 zC$^?7NE5IZbCNEp9eek;-v%z2)UXFcbMm0o)TQ>@p&^x-MQ8}E2xeh?vhIM~2XL~c zCh%$0(}S^|im@lJ=%}{3`W=%>%#wIlvRSIjs(+s@4(Y3GFxJ-PcxTfSMvHZEBU;kM zebnFG*xyW?Cc3Nm$iY&YC*FD|xb3BG`x0S{VIzd0ge2lF8AzjTu~aD0%G}e{hnV^> zo5zf@C@h0n+Kfq_8Bu6uzY(>gT900Ev|t<+g*N_MQ9CBqJSGZm@j+2LHr6~g3Lfbp zQ9I5h8plO3B0Vf>#~1cHjpL)>Qyvwy6I8!(LKNC` zAmvF>o8w9|=0xH2)TcykuFGM}jl%M@r$w#Ym1C4gp)>s%QJd$=H|9m5E8|&Ft8k4r zDx%Px`GTlby1=N6LeTM|s8zX`Q5A)rtd~Ts+T}E=qi{y{%c3^lHP@IQh0w59MeQV% zKPd{~oIz2m!71Qw@=hgM%|O6&^LU)sMVrwZ4_4Izb0yRh5caE zMPcQL*F~+qu%8+AQCKzd4N+?->=%uOD4aR!O;KA|*l#lyM&Yc{Z;9F|>ec$h+311wwj_C@W#C>YR&3uFq)$n?ws$5+9LG=Fcw8I*o}W*)E283 zfw4G>p|0QqQ9D&yEY40IFzENV;COTbtX#W3gki>NJCF9Ty~6oXvhhoW|xx*m+v zq8Q>Pek5wk)D2)Pi(-JA^s%V5s8@o~62f>Q*q?q8QkwekN*ebvqdDD26rlb5ZlC*Mi}}5ChY`5H(%B9t=H-A+7jJQS+(; zV0fb#&`Q1%HJ^GT7``an1f_oyHNSc@82%{igtDlpwX3&)(T*VmrkkSHq3!~sBMQ4= z#@C`|sJDS(L}3ri{6^FQ>g`|zqHqTs^{uF#uHFg8=@Bp8t> zJOy+AEo!~$(_r*Q;Tb4DAZmT;vtaZ^;W?PsFKH{(=fPMJh0KccC2ghp0vIc!@FG;M zmb6vsOJJ;u!tkm!l6I#0G8kt@;SW%KfuxQsQFs;Rua&g3)j=@MjzaE9>m=Hwv%8f=easJoR-j&WpkuaPno6)(`5NVD#fP_(!O{T++@5 z^-o}&kL|ahZoQ?NU&`0^?F_{|!#PM$#?=6@$lR*fwFwwUTx@ zs9%F|IkvxnrPoQ?dQiUwV?DONgVU~;v@1aU9*irn{dZV)gQRT$^#?FEVEadC8IZJ% zp#B8LMr{8BT6aj=m7x9;j4QGIGiWzT+Et+b3yiC<{cmWyNz$$c^#B-GW4oWZZ?X}FeOVYN1 zx{ev!uzd;h@0PUfpkB(1?byDIwcjRbc*rhi#x>Yp&pP%<+O?ow!HjFMy@46GOWJjy zZe+%F*uIhl?vS+WLA{C@*JJx?cKV%?b_1xJm~jKPH?!q;N!kFYTbMC`?X9fyZb{n# z>NaNV!1i|5b&sUo2l;{q>=P3d1h!py&5R$+d`wHXsVGt9qo_fhUsTJzke zxR^_rvLH|N+y&;Rgf=4D=;AJL332!{f}0R-LYO7t&k1fuxY@<0B+)MjU4-Z&6c&^4 zmjo|HcrjX#lJHjqpNjCQC@m-9zY)9y;Ux$wNjOUIQiPYH5nB>A2|f+s(-5{N;jal^ zhVU|kQI6DddgJ2)RK7@xQ;eQhBN7#?HbCU4S z1h*sHj&N=g{ujX=2zMZymxTXKuz|3FwuieH#Nh)32M`V*oS%gIMTAdB_;iFvB;oT3 zUXJi`v^~;Y?>>2VoU@uZoyh4#&ZwcBHN@#cP8V`UyX)NbyWI`D-6tjZ7ZAT2`Q6AL zb2$G(;s=o*ME=;r`4>BV`P>^|9Dli*xVoId3AA*U$04C@JAf$$214<4y22wsWs zN_1js65c@YDuh>|w3>uB5_~4YXCgc;313O@SqPtnaIw4g;9xe3?ql88LF;7vxKBid z$(-OV1Yb|^^#ns4-b(Nd1m8d~i^JOp9w2yNsPuM%cM!Z|sPr`i-$?L{1oJq2Ex~IE zUJLFy$&p+~@H&Fm5u7k`J;9d{d7+Cj{gpVw-LN;D10Zu+X>!IaAFefBKTT@uO&F)?cD@lL+~{OC%nB!#M?;$ z5T#VQoaPRe4z_3DyVRD1$_YwG$j+E3O(Ce79z2gSp&Dy7f_c*5UL(`2anHz71i@U; zIV4jS2Q$(96olY7HN#ijlU>5&41N=V07=C}xIDVgGB4+VVF$vT6r2Oq-TiI&_XGBMez7FpjZ#GgtZ zb+UJwvwyO;7zbKHKq-C7ur(c1TqfZ2<9P;4Y2FxqUtk=+FESCI9xpKopAIiG8Pn}Q zFa^`>SC|dcYuS!j^{Xr;Y%XiP7pm_CG$@&^gRN`<6O43I8&tu_Fx``IHub&k6LRdh zY~G1p*-?lU-(L6ZVWY9)A5`};_Xz{)J~rTyYGYc`@!=7 zmpg7aMY!8Qh0Y=eJ@{ zXZlIUgTUfH@SxdJ%pROdn~9f$*KUS&rTtb2dsEF`T42X#sr3Bc#ZvJV@D5AEto~h= zPFeYTEQ6Bq_gN<8#2+vRrK*2sS&veN&9a}z;zKs<1BxN<%oVMG`zodV zE?5A7yT@Fqr@6~ipg~ddEIm#Am{tD{a180@*?PL`z_3i5uXA+VXnK0}Lou;_K(*=W zr!l8u;u+>@&@)iJ809ZO0X+l70mVLJ>gOc(aLnUO^Fpx9WhBh4i<$e9sb39|bC{Qa z?#L=bbC;n2&Wnp@(&Aw#-3{vBn4U7AM%h5zq1car)`R1o6v4z8d6bS>5aReB9S8F& z7;MkXPDA~xNng*h(r*3+ zbi4cbjNvH#AlmtkxsM-EzqgLT@4(TI=AMYH&EF4Q!))_O(6cvB$x5Sk2rbAN?3gc98=Kh^MHZLd;wf@rwq890~nyp7s1lNFlD|B6sa6L zZM6A+VtKF7vx^OX6^nne{6Ep7i;egIi+{0^YZ*Sd*r;__9AKk2VbRaWY{uezK6VQh ztNFOCSghgBZCG5u$8X2tLSAqUGhe|$3^!lJe)T+DT=eoH9@~MB#JtQm2SLxr`L>gn z`LsKPyRuI}7x!68aYa4f+Gy&FO zUCkV;v9pWV&^3t{Bk^=T23BdE&c^|)gqz{OMhHTD>2M{k8^d&4g}a23{3#sve3qp{I)=1S&x zKoZPNRwrZ3&FtVvu0A*t?^tsS)A4b&C%!o;G}P_d+~*xJ#2+2J{8)Yk? z&)4vYkD;^I@=1?laUGxh92VE}qP@&~j2+xZHynDRCz+2kJfD+>oX^ST6HK4%%%O>X z5+{0qCi*FBqMv4t2f1KAi~ROcv{cUxvAko*@)w!UGrh3Q7t4uG5PxqiT-38{}A-&yns2@|zeN3NXb(leMZ9ldw(J=0y*O-o( z;T=;bM=sJQ6dr)Pvd5s%8#t<)xjtb)-Dw@yH?6__kvaA=!F-ERjGT>PZ!?OKw{U&@ zfO;z)i`ZlhsJpOU5QB>1&~B`c9#C(?VnXavd(SdcfM(ve%QV>*(5VRRFYZg>j853TwbD`)%2s*l7Ql3E`#J;&w5$;xr$;($M8x-)wXZG}7e zlrM4d@8VOx!s2d@d4l;lON9I{ta02E3;AC%mxS4{`4!VMoqMbeCPW;VE9#ka**%Su6?pX4@Pm_k0o7Y)>C;@c8>?r<^5qK& zizmq!S+Y~Dgk&wJ!41Q{*KzAtZrGtq+7WXVufiPHTx%`Xp#;TK-1>BIH^j}6iH$t5 zV&+vtMsc-OpJ4U4i630<&3x!`Z?RhAt@>80KAi4QJSE`D+_^iR+LOKQ$-!W%@?;)f zI=Q(`KO8UzxRtk9A!o-BF9&F}wt=vhM00X6ceHiTq^JebAd&um4GDT0(Q}38U zU=hAkRoIiozB^#v%B{w*nJBcI4^rc|oU+mfc`_@dvvpgnKhLfB2&d^oL2Kgf;*<+K zNm2WEy!aj-g@<^_y;wZVOYh_ABiy`)t%XrZa+@c6Xc9uV)^ovwL_egVp==oF&=*d>e7zFPnDYn z)W>;jyF5$%?v$V45Y<1B>ba`8@S0?&1fPxfNn8 z>A2{tB-#%s+Yjo~+$H|cp6sMsO+6j7Fh~~8Cky9;`phA2jvr!SRNNrTq1K$3)6b#P z)d$DDmrSiDQ>#IJ_Wy9Z`VhDGk@^}^UjyoM+$A0Ac;+E}?I+<2NcaL!pFd=lhV~`I zba||=xbADDdm-sw2xW*-R3p`GoDY z{s5&?*egY$3g|nJ(YJ!{mCSt-{VZ<3G`QCoVZJ5}?lVT3uS!w1yYsLZ!A0~|%pK?ie_gZ$PSHoM>mN59C!L>bo6BS*>ak` zwVb{c;?I)#FU$3Rlj{ZM0djqU*)OB(t!->yK0x#;3xSn2ON{r(#TU=-0_K^de-BO=8zu*c8OZC9xYU zY${?-`n#gz2)kL1xp=c2g;RjOjqGOG+$ob=x5(($RP$CD-6}G7ky}&D-Q>o2a$^!1 zxRc!Y6F2WhH>M7#Z_#jz2GqByoR%1N0sT2k(_kqnOl^HO{^u+(#<0fqh}8OA9AaaP zJc2%!TAzPfV?=?c}i}3Y($aM*W?WGMbLhNs8CgSkyWrL-Df$YBorS6xJ{($rf(yK@%_}>y~ iOzfXg{PoA&%tysJ!$!vcn)4U`(p;E}oc{m<*xjmshsk09 literal 0 HcmV?d00001