liangsheng8708
/
ruffle
mirror of https://mirror.ghproxy.com/https://github.com/ruffle-rs/ruffle.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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.
This commit is contained in:
Aaron Hill 2022-09-22 19:44:49 -05:00
parent 45515be0a3
commit d0230a2bea
16 changed files with 2157 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

45
Cargo.lock generated
View File

@ -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"

1
Cargo.toml
View File

@ -13,6 +13,7 @@ members = [
"render",
"render/canvas",
"render/naga-agal",
"render/wgpu",
"render/webgl",

16
render/naga-agal/Cargo.toml Normal file
View File

@ -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"] }

761
render/naga-agal/src/builder.rs Normal file
View File

@ -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<T> = std::result::Result<T, Error>;
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<Expression>,
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<Option<Handle<Expression>>>,
varying_pointers: Vec<Option<Handle<Expression>>>,
// An `Expression::GlobalVariables` for the uniform buffer
// that stores all of the program constants.
constant_registers: Handle<Expression>,
// 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<Type>,
// The Naga representation of 'mat4x4f'
matrix4x4f: Handle<Type>,
}
impl VertexAttributeFormat {
fn to_naga_type(self, module: &mut Module) -> Handle<Type> {
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<Expression>,
builder: &mut NagaBuilder,
) -> Result<Handle<Expression>> {
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<VertexAttributeFormat>; 8],
}
impl<'a> NagaBuilder<'a> {
pub fn process_agal(
mut agal: &[u8],
vertex_attributes: &[Option<VertexAttributeFormat>; MAX_VERTEX_ATTRIBUTES],
) -> Result<Module> {
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::<f32>() 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<Handle<Expression>> {
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<Handle<Expression>> {
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<Handle<Expression>> {
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<Handle<Expression>> {
// 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<Handle<Expression>> {
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<f32>
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::<Vec<_>>()
.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<Expression>) -> 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<Expression> {
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<Module> {
// 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)
}
}

93
render/naga-agal/src/lib.rs Normal file
View File

@ -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<std::io::Error> 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<VertexAttributeFormat>; MAX_VERTEX_ATTRIBUTES],
) -> Result<Module, Error> {
NagaBuilder::process_agal(agal, vertex_attributes)
}

195
render/naga-agal/src/types.rs Normal file
View File

@ -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<DestField, Error> {
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<SourceField, Error> {
// 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<SamplerField, Error> {
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),
}

25
render/naga-agal/tests/snapshots/wgsl__shaders-2.snap Normal file
View File

@ -0,0 +1,25 @@
---
source: tests/wgsl.rs
assertion_line: 53
expression: output
---
struct VertexOutput {
@builtin(position) member: vec4<f32>,
@location(0) varying_0_: vec4<f32>,
}
@group(0) @binding(0)
var<uniform> constant_registers: array<vec4<f32>,128u>;
@vertex
fn main(@location(0) param: vec4<f32>, @location(1) param_1: vec4<f32>) -> VertexOutput {
var dest_temp: vec4<f32>;
var varying_0_: vec4<f32>;
dest_temp = param;
varying_0_ = param_1;
let _e5: vec4<f32> = dest_temp;
let _e6: vec4<f32> = varying_0_;
return VertexOutput(_e5, _e6);
}

21
render/naga-agal/tests/snapshots/wgsl__shaders-3.snap Normal file
View File

@ -0,0 +1,21 @@
---
source: tests/wgsl.rs
assertion_line: 70
expression: output
---
struct FragmentOutput {
@location(0) member: vec4<f32>,
}
@group(0) @binding(1)
var<uniform> constant_registers: array<vec4<f32>,28u>;
@fragment
fn main(@location(0) param: vec4<f32>) -> FragmentOutput {
var dest_temp: vec4<f32>;
dest_temp = param;
let _e3: vec4<f32> = dest_temp;
return FragmentOutput(_e3);
}

29
render/naga-agal/tests/snapshots/wgsl__shaders.snap Normal file
View File

@ -0,0 +1,29 @@
---
source: tests/wgsl.rs
assertion_line: 35
expression: output
---
struct VertexOutput {
@builtin(position) member: vec4<f32>,
@location(0) varying_0_: vec4<f32>,
}
@group(0) @binding(0)
var<uniform> constant_registers: array<vec4<f32>,128u>;
@vertex
fn main(@location(0) param: vec3<f32>, @location(1) param_1: vec3<f32>) -> VertexOutput {
var dest_temp: vec4<f32>;
var varying_0_: vec4<f32>;
let _e4: vec4<f32> = constant_registers[0u];
let _e7: vec4<f32> = constant_registers[1u];
let _e10: vec4<f32> = constant_registers[2u];
let _e13: vec4<f32> = constant_registers[3u];
dest_temp = (transpose(mat4x4<f32>(_e4, _e7, _e10, _e13)) * vec4<f32>(param.x, param.y, param.z, 1.0));
varying_0_ = vec4<f32>(param_1.x, param_1.y, param_1.z, 1.0);
let _e30: vec4<f32> = dest_temp;
let _e31: vec4<f32> = varying_0_;
return VertexOutput(_e30, _e31);
}

83
render/naga-agal/tests/wgsl.rs Normal file
View File

@ -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,
);
}

1
tests/tests/regression_tests.rs
View File

@ -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),

819
tests/tests/swfs/avm2/agal_compiler/AGALMiniAssembler.as Normal file
View File

@ -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<optsLength; k++ )
{
if ( verbose )
trace( " opt: "+opts[k] );
var optfound:Sampler = SAMPLEMAP [opts[k]];
if ( optfound == null )
{
// todo check that it's a number...
//trace( "Warning, unknown sampler option: "+opts[k] );
bias = Number(opts[k]);
if ( verbose )
trace( " bias: " + bias );
}
else
{
if ( optfound.flag != SAMPLER_SPECIAL_SHIFT )
samplerbits &= ~( 0xf << optfound.flag );
samplerbits |= uint( optfound.mask ) << uint( optfound.flag );
}
}
agalcode.writeShort( regidx );
agalcode.writeByte(int(bias*8.0));
agalcode.writeByte(0);
agalcode.writeUnsignedInt( samplerbits );
if ( verbose )
trace( " bits: " + ( samplerbits - 5 ) );
pad -= 64;
}
else
{
if ( j == 0 )
{
agalcode.writeUnsignedInt( 0 );
pad -= 32;
}
agalcode.writeShort( regidx );
agalcode.writeByte( reloffset );
agalcode.writeByte( regmask );
agalcode.writeByte( regFound.emitCode );
agalcode.writeByte( reltype );
agalcode.writeShort( isRelative ? ( relsel | ( 1 << 15 ) ) : 0 );
pad -= 64;
}
}
}
// pad unused regs
for ( j = 0; j < pad; j += 8 )
agalcode.writeByte( 0 );
if ( badreg )
break;
}
if ( _error != "" )
{
_error += "\n at line " + i + " " + lines[i];
agalcode.length = 0;
trace( _error );
}
// trace the bytecode bytes if debugging is enabled
if ( debugEnabled )
{
var dbgLine:String = "generated bytecode:";
var agalLength:uint = agalcode.length;
for ( var index:uint = 0; index < agalLength; index++ )
{
if ( !( index % 16 ) )
dbgLine += "\n";
if ( !( index % 4 ) )
dbgLine += " ";
var byteStr:String = agalcode[ index ].toString( 16 );
if ( byteStr.length < 2 )
byteStr = "0" + byteStr;
dbgLine += byteStr;
}
trace( dbgLine );
}
if ( verbose )
trace( "AGALMiniAssembler.assemble time: " + ( ( getTimer() - start ) / 1000 ) + "s" );
return agalcode;
}
private function initregmap ( version:uint, ignorelimits:Boolean ) : void {
// version changes limits
REGMAP[ VA ] = new Register( VA, "vertex attribute", 0x0, ignorelimits?1024:((version==1||version==2)?7:15), REG_VERT | REG_READ );
REGMAP[ VC ] = new Register( VC, "vertex constant", 0x1, ignorelimits?1024:(version==1?127:249), REG_VERT | REG_READ );
REGMAP[ VT ] = new Register( VT, "vertex temporary", 0x2, ignorelimits?1024:(version==1?7:25), REG_VERT | REG_WRITE | REG_READ );
REGMAP[ VO ] = new Register( VO, "vertex output", 0x3, ignorelimits?1024:0, REG_VERT | REG_WRITE );
REGMAP[ VI ] = new Register( VI, "varying", 0x4, ignorelimits?1024:(version==1?7:9), REG_VERT | REG_FRAG | REG_READ | REG_WRITE );
REGMAP[ FC ] = new Register( FC, "fragment constant", 0x1, ignorelimits?1024:(version==1?27:((version==2)?63:199)), REG_FRAG | REG_READ );
REGMAP[ FT ] = new Register( FT, "fragment temporary", 0x2, ignorelimits?1024:(version==1?7:25), REG_FRAG | REG_WRITE | REG_READ );
REGMAP[ FS ] = new Register( FS, "texture sampler", 0x5, ignorelimits?1024:15, REG_FRAG | REG_READ );
REGMAP[ FO ] = new Register( FO, "fragment output", 0x3, ignorelimits?1024:(version==1?0:3), REG_FRAG | REG_WRITE );
REGMAP[ FD ] = new Register( FD, "fragment depth output",0x6, ignorelimits?1024:(version==1?-1:0), REG_FRAG | REG_WRITE );
REGMAP[ IID ] = new Register( IID,"instance id", 0x7, ignorelimits?1024:0, REG_VERT | REG_READ );
REGMAP[ VS ] = new Register( VS, "vertex texture sampler", 0x5, ignorelimits?1024:3, REG_VERT | REG_READ );
// aliases
REGMAP[ "op" ] = REGMAP[ VO ];
REGMAP[ "i" ] = REGMAP[ VI ];
REGMAP[ "v" ] = REGMAP[ VI ];
REGMAP[ "oc" ] = REGMAP[ FO ];
REGMAP[ "od" ] = REGMAP[ FD ];
REGMAP[ "fi" ] = REGMAP[ VI ];
}
static private function init():void
{
initialized = true;
// Fill the dictionaries with opcodes and registers
OPMAP[ MOV ] = new OpCode( MOV, 2, 0x00, 0 );
OPMAP[ ADD ] = new OpCode( ADD, 3, 0x01, 0 );
OPMAP[ SUB ] = new OpCode( SUB, 3, 0x02, 0 );
OPMAP[ MUL ] = new OpCode( MUL, 3, 0x03, 0 );
OPMAP[ DIV ] = new OpCode( DIV, 3, 0x04, 0 );
OPMAP[ RCP ] = new OpCode( RCP, 2, 0x05, 0 );
OPMAP[ MIN ] = new OpCode( MIN, 3, 0x06, 0 );
OPMAP[ MAX ] = new OpCode( MAX, 3, 0x07, 0 );
OPMAP[ FRC ] = new OpCode( FRC, 2, 0x08, 0 );
OPMAP[ SQT ] = new OpCode( SQT, 2, 0x09, 0 );
OPMAP[ RSQ ] = new OpCode( RSQ, 2, 0x0a, 0 );
OPMAP[ POW ] = new OpCode( POW, 3, 0x0b, 0 );
OPMAP[ LOG ] = new OpCode( LOG, 2, 0x0c, 0 );
OPMAP[ EXP ] = new OpCode( EXP, 2, 0x0d, 0 );
OPMAP[ NRM ] = new OpCode( NRM, 2, 0x0e, 0 );
OPMAP[ SIN ] = new OpCode( SIN, 2, 0x0f, 0 );
OPMAP[ COS ] = new OpCode( COS, 2, 0x10, 0 );
OPMAP[ CRS ] = new OpCode( CRS, 3, 0x11, 0 );
OPMAP[ DP3 ] = new OpCode( DP3, 3, 0x12, 0 );
OPMAP[ DP4 ] = new OpCode( DP4, 3, 0x13, 0 );
OPMAP[ ABS ] = new OpCode( ABS, 2, 0x14, 0 );
OPMAP[ NEG ] = new OpCode( NEG, 2, 0x15, 0 );
OPMAP[ SAT ] = new OpCode( SAT, 2, 0x16, 0 );
OPMAP[ M33 ] = new OpCode( M33, 3, 0x17, OP_SPECIAL_MATRIX );
OPMAP[ M44 ] = new OpCode( M44, 3, 0x18, OP_SPECIAL_MATRIX );
OPMAP[ M34 ] = new OpCode( M34, 3, 0x19, OP_SPECIAL_MATRIX );
OPMAP[ DDX ] = new OpCode( DDX, 2, 0x1a, OP_VERSION2 | OP_FRAG_ONLY );
OPMAP[ DDY ] = new OpCode( DDY, 2, 0x1b, OP_VERSION2 | OP_FRAG_ONLY );
OPMAP[ IFE ] = new OpCode( IFE, 2, 0x1c, OP_NO_DEST | OP_VERSION2 | OP_INCNEST | OP_SCALAR );
OPMAP[ INE ] = new OpCode( INE, 2, 0x1d, OP_NO_DEST | OP_VERSION2 | OP_INCNEST | OP_SCALAR );
OPMAP[ IFG ] = new OpCode( IFG, 2, 0x1e, OP_NO_DEST | OP_VERSION2 | OP_INCNEST | OP_SCALAR );
OPMAP[ IFL ] = new OpCode( IFL, 2, 0x1f, OP_NO_DEST | OP_VERSION2 | OP_INCNEST | OP_SCALAR );
OPMAP[ ELS ] = new OpCode( ELS, 0, 0x20, OP_NO_DEST | OP_VERSION2 | OP_INCNEST | OP_DECNEST | OP_SCALAR );
OPMAP[ EIF ] = new OpCode( EIF, 0, 0x21, OP_NO_DEST | OP_VERSION2 | OP_DECNEST | OP_SCALAR );
// space
//OPMAP[ TED ] = new OpCode( TED, 3, 0x26, OP_FRAG_ONLY | OP_SPECIAL_TEX | OP_VERSION2); //ted is not available in AGAL2
OPMAP[ KIL ] = new OpCode( KIL, 1, 0x27, OP_NO_DEST | OP_FRAG_ONLY );
OPMAP[ TEX ] = new OpCode( TEX, 3, 0x28, OP_FRAG_ONLY | OP_SPECIAL_TEX );
OPMAP[ SGE ] = new OpCode( SGE, 3, 0x29, 0 );
OPMAP[ SLT ] = new OpCode( SLT, 3, 0x2a, 0 );
OPMAP[ SGN ] = new OpCode( SGN, 2, 0x2b, 0 );
OPMAP[ SEQ ] = new OpCode( SEQ, 3, 0x2c, 0 );
OPMAP[ SNE ] = new OpCode( SNE, 3, 0x2d, 0 );
OPMAP[ TLD ] = new OpCode( TLD, 3, 0x2e, OP_VERT_ONLY | OP_SPECIAL_TEX );
SAMPLEMAP[ RGBA ] = new Sampler( RGBA, SAMPLER_TYPE_SHIFT, 0 );
SAMPLEMAP[ COMPRESSED ] = new Sampler( COMPRESSED, SAMPLER_TYPE_SHIFT, 1 );
SAMPLEMAP[ COMPRESSEDALPHA ] = new Sampler( COMPRESSEDALPHA, SAMPLER_TYPE_SHIFT, 2 );
SAMPLEMAP[ DXT1 ] = new Sampler( DXT1, SAMPLER_TYPE_SHIFT, 1 );
SAMPLEMAP[ DXT5 ] = new Sampler( DXT5, SAMPLER_TYPE_SHIFT, 2 );
SAMPLEMAP[ VIDEO ] = new Sampler( VIDEO, SAMPLER_TYPE_SHIFT, 3 );
SAMPLEMAP[ D2 ] = new Sampler( D2, SAMPLER_DIM_SHIFT, 0 );
SAMPLEMAP[ D3 ] = new Sampler( D3, SAMPLER_DIM_SHIFT, 2 );
SAMPLEMAP[ CUBE ] = new Sampler( CUBE, SAMPLER_DIM_SHIFT, 1 );
SAMPLEMAP[ MIPNEAREST ] = new Sampler( MIPNEAREST, SAMPLER_MIPMAP_SHIFT, 1 );
SAMPLEMAP[ MIPLINEAR ] = new Sampler( MIPLINEAR, SAMPLER_MIPMAP_SHIFT, 2 );
SAMPLEMAP[ MIPNONE ] = new Sampler( MIPNONE, SAMPLER_MIPMAP_SHIFT, 0 );
SAMPLEMAP[ NOMIP ] = new Sampler( NOMIP, SAMPLER_MIPMAP_SHIFT, 0 );
SAMPLEMAP[ NEAREST ] = new Sampler( NEAREST, SAMPLER_FILTER_SHIFT, 0 );
SAMPLEMAP[ LINEAR ] = new Sampler( LINEAR, SAMPLER_FILTER_SHIFT, 1 );
SAMPLEMAP[ ANISOTROPIC2X ] = new Sampler( ANISOTROPIC2X, SAMPLER_FILTER_SHIFT, 2 );
SAMPLEMAP[ ANISOTROPIC4X ] = new Sampler( ANISOTROPIC4X, SAMPLER_FILTER_SHIFT, 3 );
SAMPLEMAP[ ANISOTROPIC8X ] = new Sampler( ANISOTROPIC8X, SAMPLER_FILTER_SHIFT, 4 );
SAMPLEMAP[ ANISOTROPIC16X ] = new Sampler( ANISOTROPIC16X, SAMPLER_FILTER_SHIFT,5 );
SAMPLEMAP[ CENTROID ] = new Sampler( CENTROID, SAMPLER_SPECIAL_SHIFT, 1 << 0 );
SAMPLEMAP[ SINGLE ] = new Sampler( SINGLE, SAMPLER_SPECIAL_SHIFT, 1 << 1 );
SAMPLEMAP[ IGNORESAMPLER ] = new Sampler( IGNORESAMPLER, SAMPLER_SPECIAL_SHIFT, 1 << 2 );
SAMPLEMAP[ REPEAT ] = new Sampler( REPEAT, SAMPLER_REPEAT_SHIFT, 1 );
SAMPLEMAP[ WRAP ] = new Sampler( WRAP, SAMPLER_REPEAT_SHIFT, 1 );
SAMPLEMAP[ CLAMP ] = new Sampler( CLAMP, SAMPLER_REPEAT_SHIFT, 0 );
SAMPLEMAP[ CLAMP_U_REPEAT_V ] = new Sampler( CLAMP_U_REPEAT_V, SAMPLER_REPEAT_SHIFT, 2 );
SAMPLEMAP[ REPEAT_U_CLAMP_V ] = new Sampler( REPEAT_U_CLAMP_V, SAMPLER_REPEAT_SHIFT, 3 );
}
// ======================================================================
// Constants
// ----------------------------------------------------------------------
private static const OPMAP:Dictionary = new Dictionary();
private static const REGMAP:Dictionary = new Dictionary();
private static const SAMPLEMAP:Dictionary = new Dictionary();
private static const MAX_NESTING:int = 4;
private static const MAX_OPCODES:int = 4096;
private static const FRAGMENT:String = "fragment";
private static const VERTEX:String = "vertex";
// masks and shifts
private static const SAMPLER_TYPE_SHIFT:uint = 8;
private static const SAMPLER_DIM_SHIFT:uint = 12;
private static const SAMPLER_SPECIAL_SHIFT:uint = 16;
private static const SAMPLER_REPEAT_SHIFT:uint = 20;
private static const SAMPLER_MIPMAP_SHIFT:uint = 24;
private static const SAMPLER_FILTER_SHIFT:uint = 28;
// regmap flags
private static const REG_WRITE:uint = 0x1;
private static const REG_READ:uint = 0x2;
private static const REG_FRAG:uint = 0x20;
private static const REG_VERT:uint = 0x40;
// opmap flags
private static const OP_SCALAR:uint = 0x1;
private static const OP_SPECIAL_TEX:uint = 0x8;
private static const OP_SPECIAL_MATRIX:uint = 0x10;
private static const OP_FRAG_ONLY:uint = 0x20;
private static const OP_VERT_ONLY:uint = 0x40;
private static const OP_NO_DEST:uint = 0x80;
private static const OP_VERSION2:uint = 0x100;
private static const OP_INCNEST:uint = 0x200;
private static const OP_DECNEST:uint = 0x400;
// opcodes
private static const MOV:String = "mov";
private static const ADD:String = "add";
private static const SUB:String = "sub";
private static const MUL:String = "mul";
private static const DIV:String = "div";
private static const RCP:String = "rcp";
private static const MIN:String = "min";
private static const MAX:String = "max";
private static const FRC:String = "frc";
private static const SQT:String = "sqt";
private static const RSQ:String = "rsq";
private static const POW:String = "pow";
private static const LOG:String = "log";
private static const EXP:String = "exp";
private static const NRM:String = "nrm";
private static const SIN:String = "sin";
private static const COS:String = "cos";
private static const CRS:String = "crs";
private static const DP3:String = "dp3";
private static const DP4:String = "dp4";
private static const ABS:String = "abs";
private static const NEG:String = "neg";
private static const SAT:String = "sat";
private static const M33:String = "m33";
private static const M44:String = "m44";
private static const M34:String = "m34";
private static const DDX:String = "ddx";
private static const DDY:String = "ddy";
private static const IFE:String = "ife";
private static const INE:String = "ine";
private static const IFG:String = "ifg";
private static const IFL:String = "ifl";
private static const ELS:String = "els";
private static const EIF:String = "eif";
private static const TED:String = "ted";
private static const KIL:String = "kil";
private static const TEX:String = "tex";
private static const SGE:String = "sge";
private static const SLT:String = "slt";
private static const SGN:String = "sgn";
private static const SEQ:String = "seq";
private static const SNE:String = "sne";
private static const TLD:String = "tld";
// registers
private static const VA:String = "va";
private static const VC:String = "vc";
private static const VT:String = "vt";
private static const VO:String = "vo";
private static const VI:String = "vi";
private static const FC:String = "fc";
private static const FT:String = "ft";
private static const FS:String = "fs";
private static const FO:String = "fo";
private static const FD:String = "fd";
private static const IID:String = "iid";
private static const VS:String = "vs";
// samplers
private static const D2:String = "2d";
private static const D3:String = "3d";
private static const CUBE:String = "cube";
private static const MIPNEAREST:String = "mipnearest";
private static const MIPLINEAR:String = "miplinear";
private static const MIPNONE:String = "mipnone";
private static const NOMIP:String = "nomip";
private static const NEAREST:String = "nearest";
private static const LINEAR:String = "linear";
private static const ANISOTROPIC2X:String = "anisotropic2x"; //Introduced by Flash 14
private static const ANISOTROPIC4X:String = "anisotropic4x"; //Introduced by Flash 14
private static const ANISOTROPIC8X:String = "anisotropic8x"; //Introduced by Flash 14
private static const ANISOTROPIC16X:String = "anisotropic16x"; //Introduced by Flash 14
private static const CENTROID:String = "centroid";
private static const SINGLE:String = "single";
private static const IGNORESAMPLER:String = "ignoresampler";
private static const REPEAT:String = "repeat";
private static const WRAP:String = "wrap";
private static const CLAMP:String = "clamp";
private static const REPEAT_U_CLAMP_V:String = "repeat_u_clamp_v"; //Introduced by Flash 13
private static const CLAMP_U_REPEAT_V:String = "clamp_u_repeat_v"; //Introduced by Flash 13
private static const RGBA:String = "rgba";
private static const COMPRESSED:String = "compressed";
private static const COMPRESSEDALPHA:String = "compressedalpha";
private static const DXT1:String = "dxt1";
private static const DXT5:String = "dxt5";
private static const VIDEO:String = "video";
}
}
// ================================================================================
// Helper Classes
// --------------------------------------------------------------------------------
{
// ===========================================================================
// Class
// ---------------------------------------------------------------------------
class OpCode
{
// ======================================================================
// Properties
// ----------------------------------------------------------------------
private var _emitCode:uint;
private var _flags:uint;
private var _name:String;
private var _numRegister:uint;
// ======================================================================
// Getters
// ----------------------------------------------------------------------
public function get emitCode():uint { return _emitCode; }
public function get flags():uint { return _flags; }
public function get name():String { return _name; }
public function get numRegister():uint { return _numRegister; }
// ======================================================================
// Constructor
// ----------------------------------------------------------------------
public function OpCode( name:String, numRegister:uint, emitCode:uint, flags:uint)
{
_name = name;
_numRegister = numRegister;
_emitCode = emitCode;
_flags = flags;
}
// ======================================================================
// Methods
// ----------------------------------------------------------------------
public function toString():String
{
return "[OpCode name=\""+_name+"\", numRegister="+_numRegister+", emitCode="+_emitCode+", flags="+_flags+"]";
}
}
// ===========================================================================
// Class
// ---------------------------------------------------------------------------
class Register
{
// ======================================================================
// Properties
// ----------------------------------------------------------------------
private var _emitCode:uint;
private var _name:String;
private var _longName:String;
private var _flags:uint;
private var _range:uint;
// ======================================================================
// Getters
// ----------------------------------------------------------------------
public function get emitCode():uint { return _emitCode; }
public function get longName():String { return _longName; }
public function get name():String { return _name; }
public function get flags():uint { return _flags; }
public function get range():uint { return _range; }
// ======================================================================
// Constructor
// ----------------------------------------------------------------------
public function Register( name:String, longName:String, emitCode:uint, range:uint, flags:uint)
{
_name = name;
_longName = longName;
_emitCode = emitCode;
_range = range;
_flags = flags;
}
// ======================================================================
// Methods
// ----------------------------------------------------------------------
public function toString():String
{
return "[Register name=\""+_name+"\", longName=\""+_longName+"\", emitCode="+_emitCode+", range="+_range+", flags="+ _flags+"]";
}
}
// ===========================================================================
// Class
// ---------------------------------------------------------------------------
class Sampler
{
// ======================================================================
// Properties
// ----------------------------------------------------------------------
private var _flag:uint;
private var _mask:uint;
private var _name:String;
// ======================================================================
// Getters
// ----------------------------------------------------------------------
public function get flag():uint { return _flag; }
public function get mask():uint { return _mask; }
public function get name():String { return _name; }
// ======================================================================
// Constructor
// ----------------------------------------------------------------------
public function Sampler( name:String, flag:uint, mask:uint )
{
_name = name;
_flag = flag;
_mask = mask;
}
// ======================================================================
// Methods
// ----------------------------------------------------------------------
public function toString():String
{
return "[Sampler name=\""+_name+"\", flag=\""+_flag+"\", mask="+mask+"]";
}
}
}

54
tests/tests/swfs/avm2/agal_compiler/Test.as Normal file
View File

@ -0,0 +1,54 @@
package {
import flash.utils.ByteArray;
public class Test {
private function readBytes(data: ByteArray): String {
var out = new Array();
data.position = 0;
for (var i = 0; i < data.length; i++) {
out.push(data.readUnsignedByte())
}
return out;
}
function dumpShader(text:String, code:ByteArray) {
// Output in a format that we can easily copy-paste into a Rust file
for each (var line in text.split("\n")) {
trace("// " + line);
}
trace("[" + readBytes(code) + "]")
}
public function Test() {
var vertexShaders = [
"m44 op, va0, vc0 \n" + // 4x4 matrix transform
"mov v0, va1", //copy color to varying variable v0
"mov op, va0 \n" + //copy position to output
"mov v0, va1" //copy color to varying variable v0
];
var fragmentShaders = [
"mov oc, v0", //Set the output color to the value interpolated from the three triangle vertices
]
trace("Vertex shaders:");
for (var i = 0; i < vertexShaders.length; i++) {
var vertexAssembler = new AGALMiniAssembler();
vertexAssembler.assemble( "vertex", vertexShaders[i], 1, false );
dumpShader(vertexShaders[i], vertexAssembler.agalcode);
trace();
}
trace();
trace("Fragment shaders:");
for (var j = 0; j < fragmentShaders.length; j++) {
var fragmentAssembler = new AGALMiniAssembler();
fragmentAssembler.assemble( "fragment", fragmentShaders[j], 1, false );
dumpShader(fragmentShaders[j], fragmentAssembler.agalcode);
trace();
}
}
}
}

14
tests/tests/swfs/avm2/agal_compiler/output.txt Normal file
View File

@ -0,0 +1,14 @@
Vertex shaders:
// 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]
// 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]
Fragment shaders:
// 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]

BIN
tests/tests/swfs/avm2/agal_compiler/test.fla Normal file
View File

Binary file not shown.

BIN
tests/tests/swfs/avm2/agal_compiler/test.swf Normal file
View File

Binary file not shown.