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; k