[wgsl-out] Implement all statements and expressions. Enable all wgsl-in tests

src/back/wgsl/writer.rs

@@ -1,13 +1,12 @@
-// TODO: temp
-#![allow(dead_code)]
 use super::Error;
 use crate::{
     back::{binary_operation_str, vector_size_str, wgsl::keywords::RESERVED},
     proc::{EntryPointIndex, NameKey, Namer, TypeResolution},
     valid::{FunctionInfo, ModuleInfo},
-    Arena, ArraySize, Binding, Constant, Expression, FastHashMap, Function, GlobalVariable, Handle,
-    ImageClass, ImageDimension, Interpolation, Module, SampleLevel, Sampling, ScalarKind,
-    ScalarValue, ShaderStage, Statement, StorageFormat, StructMember, Type, TypeInner,
+    Arena, ArraySize, Binding, Constant, ConstantInner, Expression, FastHashMap, Function,
+    GlobalVariable, Handle, ImageClass, ImageDimension, Interpolation, Module, SampleLevel,
+    Sampling, ScalarKind, ScalarValue, ShaderStage, Statement, StorageClass, StorageFormat,
+    StructMember, Type, TypeInner,
 };
 use bit_set::BitSet;
 use std::fmt::Write;
@@ -107,7 +106,7 @@ impl<W: Write> Writer<W> {
         // Write all constants
         for (handle, constant) in module.constants.iter() {
             if constant.name.is_some() {
-                self.write_global_constant(&constant, handle)?;
+                self.write_global_constant(&module, &constant.inner, handle)?;
             }
         }
 
@@ -173,7 +172,7 @@ impl<W: Write> Writer<W> {
     fn write_scalar_value(&mut self, value: ScalarValue) -> BackendResult {
         match value {
             ScalarValue::Sint(value) => write!(self.out, "{}", value)?,
-            ScalarValue::Uint(value) => write!(self.out, "{}", value)?,
+            ScalarValue::Uint(value) => write!(self.out, "{}u", value)?,
             // Floats are written using `Debug` instead of `Display` because it always appends the
             // decimal part even it's zero
             ScalarValue::Float(value) => write!(self.out, "{:?}", value)?,
@@ -501,7 +500,7 @@ impl<W: Write> Writer<W> {
                         if multi { "multisampled_" } else { "" },
                         format!("<{}>", scalar_kind_str(kind)),
                     ),
-                    ImageClass::Depth => ("depth", "", String::from("")),
+                    ImageClass::Depth => ("depth_", "", String::from("")),
                     ImageClass::Storage(storage_format) => (
                         "storage_",
                         "",
@@ -682,8 +681,105 @@ impl<W: Write> Writer<W> {
                 self.write_expr(module, value, func_ctx)?;
                 writeln!(self.out, ");")?;
             }
-            _ => {
-                return Err(Error::Unimplemented(format!("write_stmt {:?}", stmt)));
+            // TODO: copy-paste from glsl-out
+            Statement::Block(ref block) => {
+                write!(self.out, "{}", INDENT.repeat(indent))?;
+                writeln!(self.out, "{{")?;
+                for sta in block.iter() {
+                    // Increase the indentation to help with readability
+                    self.write_stmt(module, sta, func_ctx, indent + 1)?
+                }
+                writeln!(self.out, "{}}}", INDENT.repeat(indent))?
+            }
+            Statement::Switch {
+                selector,
+                ref cases,
+                ref default,
+            } => {
+                // Start the switch
+                write!(self.out, "{}", INDENT.repeat(indent))?;
+                write!(self.out, "switch(")?;
+                self.write_expr(module, selector, func_ctx)?;
+                writeln!(self.out, ") {{")?;
+
+                // Write all cases
+                let mut write_case = true;
+                let all_fall_through = cases
+                    .iter()
+                    .all(|case| case.fall_through && case.body.is_empty());
+                if !cases.is_empty() {
+                    for case in cases {
+                        if write_case {
+                            write!(self.out, "{}case ", INDENT.repeat(indent + 1))?;
+                        }
+                        if !all_fall_through && case.fall_through && case.body.is_empty() {
+                            write_case = false;
+                            write!(self.out, "{}, ", case.value)?;
+                            continue;
+                        } else {
+                            write_case = true;
+                            writeln!(self.out, "{}: {{", case.value)?;
+                        }
+
+                        for sta in case.body.iter() {
+                            self.write_stmt(module, sta, func_ctx, indent + 2)?;
+                        }
+
+                        if case.fall_through {
+                            writeln!(self.out, "{}fallthrough;", INDENT.repeat(indent + 2))?;
+                        }
+                    }
+
+                    writeln!(self.out, "{}}}", INDENT.repeat(indent + 1))?;
+                }
+
+                if !default.is_empty() {
+                    writeln!(self.out, "{}default: {{", INDENT.repeat(indent + 1))?;
+
+                    for sta in default {
+                        self.write_stmt(module, sta, func_ctx, indent + 2)?;
+                    }
+
+                    writeln!(self.out, "{}}}", INDENT.repeat(indent + 1))?
+                }
+
+                writeln!(self.out, "{}}}", INDENT.repeat(indent))?
+            }
+            Statement::Loop {
+                ref body,
+                ref continuing,
+            } => {
+                write!(self.out, "{}", INDENT.repeat(indent))?;
+                writeln!(self.out, "loop {{")?;
+
+                for sta in body.iter() {
+                    self.write_stmt(module, sta, func_ctx, indent + 1)?;
+                }
+
+                if !continuing.is_empty() {
+                    writeln!(self.out, "{}continuing {{", INDENT.repeat(indent + 1))?;
+                    for sta in continuing.iter() {
+                        self.write_stmt(module, sta, func_ctx, indent + 2)?;
+                    }
+                    writeln!(self.out, "{}}}", INDENT.repeat(indent + 1))?;
+                }
+
+                writeln!(self.out, "{}}}", INDENT.repeat(indent))?
+            }
+            Statement::Break => {
+                writeln!(self.out, "{}break;", INDENT.repeat(indent))?;
+            }
+            Statement::Continue => {
+                writeln!(self.out, "{}continue;", INDENT.repeat(indent))?;
+            }
+            Statement::Barrier(barrier) => {
+                if barrier.contains(crate::Barrier::STORAGE) {
+                    writeln!(self.out, "{}storageBarrier();", INDENT.repeat(indent))?;
+                }
+
+                if barrier.contains(crate::Barrier::WORK_GROUP) {
+                    writeln!(self.out, "{}workgroupBarrier();", INDENT.repeat(indent))?;
+                }
             }
         }
 
@@ -969,7 +1065,6 @@ impl<W: Write> Writer<W> {
                 }
                 write!(self.out, ")")?;
             }
-            // TODO: copy-paste from msl-out
             Expression::GlobalVariable(handle) => {
                 let name = &self.names[&NameKey::GlobalVariable(handle)];
                 write!(self.out, "{}", name)?;
@@ -1023,7 +1118,6 @@ impl<W: Write> Writer<W> {
                 self.write_expr(module, value, func_ctx)?;
                 write!(self.out, ")")?;
             }
-            //TODO: add pointer logic
             Expression::Load { pointer } => self.write_expr(module, pointer, func_ctx)?,
             Expression::LocalVariable(handle) => {
                 let name_key = match func_ctx.ty {
@@ -1036,6 +1130,9 @@ impl<W: Write> Writer<W> {
             }
             Expression::ArrayLength(expr) => {
                 write!(self.out, "arrayLength(")?;
+                if is_deref_required(expr, module, func_ctx.info) {
+                    write!(self.out, "&")?;
+                };
                 self.write_expr(module, expr, func_ctx)?;
                 write!(self.out, ")")?;
             }
@@ -1128,8 +1225,67 @@ impl<W: Write> Writer<W> {
                     self.out.write_char(COMPONENTS[sc as usize])?;
                 }
             }
-            _ => {
-                return Err(Error::Unimplemented(format!("write_expr {:?}", expression)));
+            Expression::Unary { op, expr } => {
+                let unary = match op {
+                    crate::UnaryOperator::Negate => "-",
+                    crate::UnaryOperator::Not => {
+                        match *func_ctx.info[expr].ty.inner_with(&module.types) {
+                            TypeInner::Scalar {
+                                kind: ScalarKind::Bool,
+                                ..
+                            }
+                            | TypeInner::Vector { .. } => "!",
+                            _ => "~",
+                        }
+                    }
+                };
+
+                write!(self.out, "{}(", unary)?;
+                self.write_expr(module, expr, func_ctx)?;
+
+                write!(self.out, ")")?
+            }
+            Expression::Select {
+                condition,
+                accept,
+                reject,
+            } => {
+                write!(self.out, "select(")?;
+                self.write_expr(module, accept, func_ctx)?;
+                write!(self.out, ", ")?;
+                self.write_expr(module, reject, func_ctx)?;
+                write!(self.out, ", ")?;
+                self.write_expr(module, condition, func_ctx)?;
+                write!(self.out, ")")?
+            }
+            Expression::Derivative { axis, expr } => {
+                let op = match axis {
+                    crate::DerivativeAxis::X => "dpdx",
+                    crate::DerivativeAxis::Y => "dpdy",
+                    crate::DerivativeAxis::Width => "fwidth",
+                };
+                write!(self.out, "{}(", op)?;
+                self.write_expr(module, expr, func_ctx)?;
+                write!(self.out, ")")?
+            }
+            Expression::Relational { fun, argument } => {
+                let fun_name = match fun {
+                    crate::RelationalFunction::IsFinite => "isFinite",
+                    crate::RelationalFunction::IsInf => "isInf",
+                    crate::RelationalFunction::IsNan => "isNan",
+                    crate::RelationalFunction::IsNormal => "isNormal",
+                    crate::RelationalFunction::All => "all",
+                    crate::RelationalFunction::Any => "any",
+                };
+                write!(self.out, "{}(", fun_name)?;
+
+                self.write_expr(module, argument, func_ctx)?;
+
+                write!(self.out, ")")?
+            }
+            Expression::Call(function) => {
+                let func_name = &self.names[&NameKey::Function(function)];
+                write!(self.out, "{}(", func_name)?
             }
         }
 
@@ -1158,8 +1314,12 @@ impl<W: Write> Writer<W> {
             writeln!(self.out)?;
         }
 
-        // First write only global name
-        write!(self.out, "var {}: ", name)?;
+        // First write global name and storage class if supported
+        write!(self.out, "var")?;
+        if let Some(storage_class) = storage_class_str(global.class) {
+            write!(self.out, "<{}>", storage_class)?;
+        }
+        write!(self.out, " {}: ", name)?;
         // Write access attribute if present
         if !global.storage_access.is_empty() {
             self.write_attributes(&[Attribute::Access(global.storage_access)], true)?;
@@ -1215,10 +1375,11 @@ impl<W: Write> Writer<W> {
     /// Ends in a newline
     fn write_global_constant(
         &mut self,
-        constant: &Constant,
+        module: &Module,
+        inner: &ConstantInner,
         handle: Handle<Constant>,
     ) -> BackendResult {
-        match constant.inner {
+        match *inner {
             crate::ConstantInner::Scalar {
                 width: _,
                 ref value,
@@ -1232,7 +1393,7 @@ impl<W: Write> Writer<W> {
                         write!(self.out, "i32 = {}", value)?;
                     }
                     crate::ScalarValue::Uint(value) => {
-                        write!(self.out, "u32 = {}", value)?;
+                        write!(self.out, "u32 = {}u", value)?;
                     }
                     crate::ScalarValue::Float(value) => {
                         // Floats are written using `Debug` instead of `Display` because it always appends the
@@ -1243,18 +1404,33 @@ impl<W: Write> Writer<W> {
                         write!(self.out, "bool = {}", value)?;
                     }
                 };
-                // End with semicolon and extra newline for readability
+                // End with semicolon
                 writeln!(self.out, ";")?;
-                writeln!(self.out)?;
             }
-            _ => {
-                return Err(Error::Unimplemented(format!(
-                    "write_global_constant {:?}",
-                    constant.inner
-                )));
+            ConstantInner::Composite { ty, ref components } => {
+                let name = self.names[&NameKey::Constant(handle)].clone();
+                // First write only constant name
+                write!(self.out, "let {}: ", name)?;
+                // Next write constant type
+                self.write_type(module, ty)?;
+
+                write!(self.out, " = ")?;
+                self.write_type(module, ty)?;
+
+                write!(self.out, "(")?;
+                for (index, constant) in components.iter().enumerate() {
+                    self.write_constant(module, *constant)?;
+                    // Only write a comma if isn't the last element
+                    if index != components.len().saturating_sub(1) {
+                        // The leading space is for readability only
+                        write!(self.out, ", ")?;
+                    }
+                }
+                write!(self.out, ");")?;
             }
         }
-
+        // End with extra newline for readability
+        writeln!(self.out)?;
         Ok(())
     }
 
@@ -1355,6 +1531,17 @@ fn sampling_str(sampling: Sampling) -> &'static str {
     }
 }
 
+fn storage_class_str(storage_class: StorageClass) -> Option<&'static str> {
+    match storage_class {
+        StorageClass::Private => Some("private"),
+        StorageClass::Uniform => Some("uniform"),
+        StorageClass::Storage => Some("storage"),
+        StorageClass::PushConstant => Some("push_constant"),
+        StorageClass::WorkGroup => Some("workgroup"),
+        StorageClass::Function | StorageClass::Handle => None,
+    }
+}
+
 fn map_binding_to_attribute(binding: &Binding) -> Vec<Attribute> {
     match *binding {
         Binding::BuiltIn(built_in) => vec![Attribute::BuiltIn(built_in)],
@@ -1368,3 +1555,16 @@ fn map_binding_to_attribute(binding: &Binding) -> Vec<Attribute> {
         ],
     }
 }
+
+fn is_deref_required(expr: Handle<Expression>, module: &Module, info: &FunctionInfo) -> bool {
+    let base_ty_res = &info[expr].ty;
+    let resolved = base_ty_res.inner_with(&module.types);
+    match *resolved {
+        TypeInner::Pointer { base, class: _ } => match module.types[base].inner {
+            TypeInner::Scalar { .. } | TypeInner::Vector { .. } | TypeInner::Array { .. } => true,
+            _ => false,
+        },
+        TypeInner::ValuePointer { .. } => true,
+        _ => false,
+    }
+}

tests/out/access.wgsl

@@ -0,0 +1,13 @@
+[[block]]
+struct Bar {
+    matrix: mat4x4<f32>;
+    data: [[stride(4)]] array<i32>;
+};
+
+[[group(0), binding(0)]]
+var<storage> bar: [[access(read_write)]] Bar;
+
+[[stage(vertex)]]
+fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
+    return vec4<f32>(vec4<i32>(array<i32,5>(bar.data[(arrayLength(&bar.data) - 1u)], i32(bar.matrix[3u].x), 3, 4, 5)[vi]));
+}

tests/out/boids.wgsl

@@ -0,0 +1,100 @@
+struct Particle {
+    pos: vec2<f32>;
+    vel: vec2<f32>;
+};
+
+[[block]]
+struct SimParams {
+    deltaT: f32;
+    rule1Distance: f32;
+    rule2Distance: f32;
+    rule3Distance: f32;
+    rule1Scale: f32;
+    rule2Scale: f32;
+    rule3Scale: f32;
+};
+
+[[block]]
+struct Particles {
+    particles: [[stride(16)]] array<Particle>;
+};
+
+let NUM_PARTICLES: u32 = 1500u;
+
+[[group(0), binding(0)]]
+var<uniform> params: SimParams;
+[[group(0), binding(1)]]
+var<storage> particlesSrc: [[access(read)]] Particles;
+[[group(0), binding(2)]]
+var<storage> particlesDst: [[access(read_write)]] Particles;
+
+[[stage(compute), workgroup_size(64, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3<u32>) {
+    var vPos: vec2<f32>;
+    var vVel: vec2<f32>;
+    var cMass: vec2<f32>;
+    var cVel: vec2<f32>;
+    var colVel: vec2<f32>;
+    var cMassCount: i32 = 0;
+    var cVelCount: i32 = 0;
+    var pos1: vec2<f32>;
+    var vel1: vec2<f32>;
+    var i: u32 = 0u;
+
+    if ((global_invocation_id.x >= NUM_PARTICLES)) {
+        return;
+    }
+    vPos = particlesSrc.particles[global_invocation_id.x].pos;
+    vVel = particlesSrc.particles[global_invocation_id.x].vel;
+    cMass = vec2<f32>(0.0, 0.0);
+    cVel = vec2<f32>(0.0, 0.0);
+    colVel = vec2<f32>(0.0, 0.0);
+    loop {
+        if ((i >= NUM_PARTICLES)) {
+            break;
+        }
+        if ((i == global_invocation_id.x)) {
+            continue;
+        }
+        pos1 = particlesSrc.particles[i].pos;
+        vel1 = particlesSrc.particles[i].vel;
+        if ((distance(pos1, vPos) < params.rule1Distance)) {
+            cMass = (cMass + pos1);
+            cMassCount = (cMassCount + 1);
+        }
+        if ((distance(pos1, vPos) < params.rule2Distance)) {
+            colVel = (colVel - (pos1 - vPos));
+        }
+        if ((distance(pos1, vPos) < params.rule3Distance)) {
+            cVel = (cVel + vel1);
+            cVelCount = (cVelCount + 1);
+        }
+        continuing {
+            i = (i + 1u);
+        }
+    }
+    if ((cMassCount > 0)) {
+        cMass = ((cMass / vec2<f32>(f32(cMassCount))) - vPos);
+    }
+    if ((cVelCount > 0)) {
+        cVel = (cVel / vec2<f32>(f32(cVelCount)));
+    }
+    vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
+    vVel = (normalize(vVel) * clamp(length(vVel), 0.0, 0.1));
+    vPos = (vPos + (vVel * params.deltaT));
+    if ((vPos.x < -1.0)) {
+        vPos.x = 1.0;
+    }
+    if ((vPos.x > 1.0)) {
+        vPos.x = -1.0;
+    }
+    if ((vPos.y < -1.0)) {
+        vPos.y = 1.0;
+    }
+    if ((vPos.y > 1.0)) {
+        vPos.y = -1.0;
+    }
+    particlesDst.particles[global_invocation_id.x].pos = vPos;
+    particlesDst.particles[global_invocation_id.x].vel = vVel;
+    return;
+}

tests/out/collatz.wgsl

@@ -0,0 +1,33 @@
+[[block]]
+struct PrimeIndices {
+    data: [[stride(4)]] array<u32>;
+};
+
+[[group(0), binding(0)]]
+var<storage> v_indices: [[access(read_write)]] PrimeIndices;
+
+fn collatz_iterations(n_base: u32) -> u32 {
+    var n: u32;
+    var i: u32 = 0u;
+
+    n = n_base;
+    loop {
+        if ((n <= 1u)) {
+            break;
+        }
+        if (((n % 2u) == 0u)) {
+            n = (n / 2u);
+        } else {
+            n = ((3u * n) + 1u);
+        }
+        i = (i + 1u);
+    }
+    return i;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+    let _e9: u32 = collatz_iterations(v_indices.data[global_id.x]);
+    v_indices.data[global_id.x] = _e9;
+    return;
+}

tests/out/control-flow.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+    storageBarrier();
+    workgroupBarrier();
+    return;
+}

tests/out/extra.wgsl

@@ -0,0 +1,12 @@
+[[block]]
+struct PushConstants {
+    index: u32;
+    double: vec2<f32>;
+};
+
+var<push_constant> pc: PushConstants;
+
+[[stage(fragment)]]
+fn main([[location(0), interpolate(perspective)]] color: vec4<f32>) -> [[location(0)]] vec4<f32> {
+    return color;
+}

tests/out/interface.wgsl

@@ -0,0 +1,25 @@
+struct VertexOutput {
+    [[builtin(position)]] position: vec4<f32>;
+    [[location(1), interpolate(perspective)]] varying: f32;
+};
+
+struct FragmentOutput {
+    [[builtin(frag_depth)]] depth: f32;
+    [[builtin(sample_mask)]] sample_mask: u32;
+    [[location(0)]] color: f32;
+};
+
+[[stage(vertex)]]
+fn vertex([[builtin(vertex_index)]] vertex_index: u32, [[builtin(instance_index)]] instance_index: u32, [[location(10)]] color1: u32) -> VertexOutput {
+    return VertexOutput(vec4<f32>(1.0), f32(((vertex_index + instance_index) + color1)));
+}
+
+[[stage(fragment)]]
+fn fragment(in: VertexOutput, [[builtin(front_facing)]] front_facing: bool, [[builtin(sample_index)]] sample_index: u32, [[builtin(sample_mask)]] sample_mask1: u32) -> FragmentOutput {
+    return FragmentOutput(in.varying, (sample_mask1 & (1u << sample_index)), select(0.0, 1.0, front_facing));
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn compute([[builtin(global_invocation_id)]] global_id: vec3<u32>, [[builtin(local_invocation_id)]] local_id: vec3<u32>, [[builtin(local_invocation_index)]] local_index: u32, [[builtin(workgroup_id)]] wg_id: vec3<u32>) {
+    return;
+}

tests/out/interpolate.wgsl

@@ -0,0 +1,30 @@
+struct FragmentInput {
+    [[builtin(position)]] position: vec4<f32>;
+    [[location(0), interpolate(flat)]] flat: u32;
+    [[location(1), interpolate(linear)]] linear: f32;
+    [[location(2), interpolate(linear,centroid)]] linear_centroid: vec2<f32>;
+    [[location(3), interpolate(linear,sample)]] linear_sample: vec3<f32>;
+    [[location(4), interpolate(perspective)]] perspective: vec4<f32>;
+    [[location(5), interpolate(perspective,centroid)]] perspective_centroid: f32;
+    [[location(6), interpolate(perspective,sample)]] perspective_sample: f32;
+};
+
+[[stage(vertex)]]
+fn main() -> FragmentInput {
+    var out: FragmentInput;
+
+    out.position = vec4<f32>(2.0, 4.0, 5.0, 6.0);
+    out.flat = 8u;
+    out.linear = 27.0;
+    out.linear_centroid = vec2<f32>(64.0, 125.0);
+    out.linear_sample = vec3<f32>(216.0, 343.0, 512.0);
+    out.perspective = vec4<f32>(729.0, 1000.0, 1331.0, 1728.0);
+    out.perspective_centroid = 2197.0;
+    out.perspective_sample = 2744.0;
+    return out;
+}
+
+[[stage(fragment)]]
+fn main1(val: FragmentInput) {
+    return;
+}

tests/out/operators.wgsl

@@ -0,0 +1,12 @@
+fn splat() -> vec4<f32> {
+    return ((((vec2<f32>(1.0) + vec2<f32>(2.0)) - vec2<f32>(3.0)) / vec2<f32>(4.0)).xyxy + vec4<f32>((vec4<i32>(5) % vec4<i32>(2))));
+}
+
+fn unary() -> i32 {
+    if (!(true)) {
+        return 1;
+    } else {
+        return ~(1);
+    }
+}
+

tests/out/quad-vert.wgsl

@@ -8,10 +8,10 @@ struct VertexOutput {
     [[builtin(position)]] gl_Position1: vec4<f32>;
 };
 
-var v_uv: vec2<f32>;
-var a_uv: vec2<f32>;
-var perVertexStruct: gl_PerVertex;
-var a_pos: vec2<f32>;
+var<private> v_uv: vec2<f32>;
+var<private> a_uv: vec2<f32>;
+var<private> perVertexStruct: gl_PerVertex;
+var<private> a_pos: vec2<f32>;
 
 fn main() {
     v_uv = a_uv;

tests/out/shadow.wgsl

@@ -0,0 +1,54 @@
+[[block]]
+struct Globals {
+    num_lights: vec4<u32>;
+};
+
+struct Light {
+    proj: mat4x4<f32>;
+    pos: vec4<f32>;
+    color: vec4<f32>;
+};
+
+[[block]]
+struct Lights {
+    data: [[stride(96)]] array<Light>;
+};
+
+let c_ambient: vec3<f32> = vec3<f32>(0.05, 0.05, 0.05);
+let c_max_lights: u32 = 10u;
+
+[[group(0), binding(0)]]
+var<uniform> u_globals: Globals;
+[[group(0), binding(1)]]
+var<storage> s_lights: [[access(read)]] Lights;
+[[group(0), binding(2)]]
+var t_shadow: texture_depth_2d_array;
+[[group(0), binding(3)]]
+var sampler_shadow: sampler_comparison;
+
+fn fetch_shadow(light_id: u32, homogeneous_coords: vec4<f32>) -> f32 {
+    if ((homogeneous_coords.w <= 0.0)) {
+        return 1.0;
+    }
+    let _e26: f32 = textureSampleCompare(t_shadow, sampler_shadow, (((homogeneous_coords.xy * vec2<f32>(0.5, -0.5)) / vec2<f32>(homogeneous_coords.w)) + vec2<f32>(0.5, 0.5)), i32(light_id), (homogeneous_coords.z / homogeneous_coords.w));
+    return _e26;
+}
+
+[[stage(fragment)]]
+fn fs_main([[location(0), interpolate(perspective)]] raw_normal: vec3<f32>, [[location(1), interpolate(perspective)]] position: vec4<f32>) -> [[location(0)]] vec4<f32> {
+    var color1: vec3<f32> = vec3<f32>(0.05, 0.05, 0.05);
+    var i: u32 = 0u;
+
+    loop {
+        if ((i >= min(u_globals.num_lights.x, c_max_lights))) {
+            break;
+        }
+        let _e21: Light = s_lights.data[i];
+        let _e25: f32 = fetch_shadow(i, (_e21.proj * position));
+        color1 = (color1 + ((_e25 * max(0.0, dot(normalize(raw_normal), normalize((_e21.pos.xyz - position.xyz))))) * _e21.color.xyz));
+        continuing {
+            i = (i + 1u);
+        }
+    }
+    return vec4<f32>(color1, 1.0);
+}

tests/out/skybox.wgsl

@@ -0,0 +1,34 @@
+struct VertexOutput {
+    [[builtin(position)]] position: vec4<f32>;
+    [[location(0), interpolate(perspective)]] uv: vec3<f32>;
+};
+
+[[block]]
+struct Data {
+    proj_inv: mat4x4<f32>;
+    view: mat4x4<f32>;
+};
+
+[[group(0), binding(0)]]
+var<uniform> r_data: Data;
+[[group(0), binding(1)]]
+var r_texture: texture_cube<f32>;
+[[group(0), binding(2)]]
+var r_sampler: sampler;
+
+[[stage(vertex)]]
+fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
+    var tmp1_: i32;
+    var tmp2_: i32;
+
+    tmp1_ = (i32(vertex_index) / 2);
+    tmp2_ = (i32(vertex_index) & 1);
+    let _e24: vec4<f32> = vec4<f32>(((f32(tmp1_) * 4.0) - 1.0), ((f32(tmp2_) * 4.0) - 1.0), 0.0, 1.0);
+    return VertexOutput(_e24, (transpose(mat3x3<f32>(r_data.view[0].xyz, r_data.view[1].xyz, r_data.view[2].xyz)) * (r_data.proj_inv * _e24).xyz));
+}
+
+[[stage(fragment)]]
+fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+    let _e5: vec4<f32> = textureSample(r_texture, r_sampler, in.uv);
+    return _e5;
+}

tests/out/standard.wgsl

@@ -0,0 +1,7 @@
+[[stage(fragment)]]
+fn derivatives([[builtin(position)]] foo: vec4<f32>) -> [[location(0)]] vec4<f32> {
+    let _e1: vec4<f32> = dpdx(foo);
+    let _e2: vec4<f32> = dpdy(foo);
+    let _e3: vec4<f32> = fwidth(foo);
+    return ((_e1 + _e2) * _e3);
+}

tests/snapshots.rs

@@ -252,28 +252,47 @@ fn convert_wgsl() {
             "quad",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::DOT | Targets::WGSL,
         ),
-        ("boids", Targets::SPIRV | Targets::METAL | Targets::GLSL),
-        ("skybox", Targets::SPIRV | Targets::METAL | Targets::GLSL),
+        (
+            "boids",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
+        (
+            "skybox",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
         (
             "collatz",
-            Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS,
+            Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS | Targets::WGSL,
+        ),
+        (
+            "shadow",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
         ),
-        ("shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL),
         ("image", Targets::SPIRV | Targets::METAL | Targets::WGSL),
-        ("extra", Targets::SPIRV | Targets::METAL),
-        ("operators", Targets::SPIRV | Targets::METAL | Targets::GLSL),
+        ("extra", Targets::SPIRV | Targets::METAL | Targets::WGSL),
+        (
+            "operators",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
         (
             "interpolate",
-            Targets::SPIRV | Targets::METAL | Targets::GLSL,
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
         ),
-        ("access", Targets::SPIRV | Targets::METAL),
+        ("access", Targets::SPIRV | Targets::METAL | Targets::WGSL),
         (
             "control-flow",
-            Targets::SPIRV | Targets::METAL | Targets::GLSL,
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
+        (
+            "standard",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+        ),
+        (
+            "standard",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
         ),
-        ("standard", Targets::SPIRV | Targets::METAL | Targets::GLSL),
         //TODO: GLSL https://github.com/gfx-rs/naga/issues/874
-        ("interface", Targets::SPIRV | Targets::METAL),
+        ("interface", Targets::SPIRV | Targets::METAL | Targets::WGSL),
     ];
 
     for &(name, targets) in inputs.iter() {