[wgsl] Add support for simplified vector construction, add a snapshot test

This commit is contained in:
Dzmitry Malyshau
2021-04-14 10:05:48 -04:00
committed by Dzmitry Malyshau
parent c8d26a1f47
commit 96f9cb4ce1
9 changed files with 131 additions and 66 deletions

View File

@@ -1477,9 +1477,13 @@ impl<'a, W: Write> Writer<'a, W> {
Expression::Constant(constant) => {
self.write_constant(&self.module.constants[constant])?
}
// `Splat` is just writing `value`
// `Splat` needs to actually write down a vector, it's not always inferred in GLSL.
Expression::Splat { size: _, value } => {
let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
self.write_value_type(resolved)?;
write!(self.out, "(")?;
self.write_expr(value, ctx)?;
write!(self.out, ")")?
}
// `Compose` is pretty simple we just write `type(components)` where `components` is a
// comma separated list of expressions

View File

@@ -1132,6 +1132,13 @@ impl Parser {
let expr = if components.is_empty() {
let last_component_inner = ctx.resolve_type(last_component)?;
match (&inner, last_component_inner) {
(
&crate::TypeInner::Vector { size, .. },
&crate::TypeInner::Scalar { .. },
) => crate::Expression::Splat {
size,
value: last_component,
},
(
&crate::TypeInner::Scalar { .. },
&crate::TypeInner::Scalar { .. },
@@ -1215,6 +1222,7 @@ impl Parser {
crate::TypeInner::Vector { size, kind, width } => {
match Composition::make(handle, size, name, name_span, ctx.expressions)?
{
//TODO: Swizzling in IR
Composition::Multi(size, components) => {
let inner = crate::TypeInner::Vector { size, kind, width };
crate::Expression::Compose {
@@ -1238,6 +1246,7 @@ impl Parser {
name_span,
ctx.expressions,
)? {
//TODO: is this really supported?
Composition::Multi(columns, components) => {
let inner = crate::TypeInner::Matrix {
columns,

View File

@@ -0,0 +1,7 @@
(
spv_version: (1, 0),
spv_capabilities: [ Shader ],
spv_debug: false,
spv_adjust_coordinate_space: false,
msl_custom: false,
)

6
tests/in/operators.wgsl Normal file
View File

@@ -0,0 +1,6 @@
[[stage(vertex)]]
fn splat() -> [[builtin(position)]] vec4<f32> {
let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
let b = vec4<i32>(5) % 2;
return a.xyxy + vec4<f32>(b);
}

View File

@@ -0,0 +1,11 @@
#version 310 es
precision highp float;
void main() {
vec2 _expr10 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0));
gl_Position = (vec4(_expr10[0], _expr10[1], _expr10[0], _expr10[1]) + vec4((ivec4(5) % ivec4(2))));
return;
}

12
tests/out/operators.msl Normal file
View File

@@ -0,0 +1,12 @@
#include <metal_stdlib>
#include <simd/simd.h>
struct splatOutput {
metal::float4 member [[position]];
};
vertex splatOutput splat(
) {
metal::float2 _e10 = ((1.0 + 2.0) - 3.0) / 4.0;
return splatOutput { metal::float4(_e10.x, _e10.y, _e10.x, _e10.y) + static_cast<float4>(5 % 2) };
}

View File

@@ -0,0 +1,48 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 37
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %15 "splat" %13
OpDecorate %13 BuiltIn Position
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 1.0
%5 = OpConstant %4 2.0
%6 = OpConstant %4 3.0
%7 = OpConstant %4 4.0
%9 = OpTypeInt 32 1
%8 = OpConstant %9 5
%10 = OpConstant %9 2
%11 = OpTypeVector %4 4
%14 = OpTypePointer Output %11
%13 = OpVariable %14 Output
%16 = OpTypeFunction %2
%18 = OpTypeVector %4 2
%26 = OpTypeVector %9 4
%15 = OpFunction %2 None %16
%12 = OpLabel
OpBranch %17
%17 = OpLabel
%19 = OpCompositeConstruct %18 %5 %5
%20 = OpCompositeConstruct %18 %3 %3
%21 = OpFAdd %18 %20 %19
%22 = OpCompositeConstruct %18 %6 %6
%23 = OpFSub %18 %21 %22
%24 = OpCompositeConstruct %18 %7 %7
%25 = OpFDiv %18 %23 %24
%27 = OpCompositeConstruct %26 %8 %8 %8 %8
%28 = OpCompositeConstruct %26 %10 %10 %10 %10
%29 = OpSMod %26 %27 %28
%30 = OpCompositeExtract %4 %25 0
%31 = OpCompositeExtract %4 %25 1
%32 = OpCompositeExtract %4 %25 0
%33 = OpCompositeExtract %4 %25 1
%34 = OpCompositeConstruct %11 %30 %31 %32 %33
%35 = OpConvertSToF %11 %29
%36 = OpFAdd %11 %34 %35
OpStore %13 %36
OpReturn
OpFunctionEnd

View File

@@ -26,7 +26,7 @@ float fetch_shadow(uint light_id, vec4 homogeneous_coords) {
if((homogeneous_coords[3] <= 0.0)) {
return 1.0;
}
float _expr28 = textureGrad(_group_0_binding_2, vec4((((vec2(homogeneous_coords[0], homogeneous_coords[1]) * vec2(0.5, -0.5)) / homogeneous_coords[3]) + vec2(0.5, 0.5)), int(light_id), (homogeneous_coords[2] / homogeneous_coords[3])), vec2(0, 0), vec2(0,0));
float _expr28 = textureGrad(_group_0_binding_2, vec4((((vec2(homogeneous_coords[0], homogeneous_coords[1]) * vec2(0.5, -0.5)) / vec2(homogeneous_coords[3])) + vec2(0.5, 0.5)), int(light_id), (homogeneous_coords[2] / homogeneous_coords[3])), vec2(0, 0), vec2(0,0));
return _expr28;
}

View File

@@ -195,72 +195,40 @@ fn check_output_hlsl(module: &naga::Module, destination: &PathBuf) {
}
#[cfg(feature = "wgsl-in")]
fn convert_wgsl(name: &str, targets: Targets) {
#[test]
fn convert_wgsl() {
let root = env!("CARGO_MANIFEST_DIR");
let module = naga::front::wgsl::parse_str(
&fs::read_to_string(format!("{}/{}/{}.wgsl", root, DIR_IN, name))
.expect("Couldn't find wgsl file"),
)
.unwrap();
check_targets(&module, name, targets);
}
let inputs = [
(
"empty",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL,
),
(
"quad",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::DOT,
),
("boids", Targets::SPIRV | Targets::METAL),
("skybox", Targets::SPIRV | Targets::METAL | Targets::GLSL),
(
"collatz",
Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS,
),
("shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL),
//SPIR-V is blocked by https://github.com/gfx-rs/naga/issues/646
("image-copy", Targets::METAL),
("texture-array", Targets::SPIRV | Targets::METAL),
("operators", Targets::SPIRV | Targets::METAL | Targets::GLSL),
];
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_quad() {
convert_wgsl(
"quad",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::DOT,
);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_empty() {
convert_wgsl(
"empty",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL,
);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_boids() {
convert_wgsl("boids", Targets::SPIRV | Targets::METAL);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_skybox() {
convert_wgsl("skybox", Targets::SPIRV | Targets::METAL | Targets::GLSL);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_collatz() {
convert_wgsl(
"collatz",
Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS,
);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_shadow() {
convert_wgsl("shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_image_copy() {
//SPIR-V is blocked by https://github.com/gfx-rs/naga/issues/646
convert_wgsl("image-copy", Targets::METAL);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn convert_wgsl_texture_array() {
convert_wgsl("texture-array", Targets::SPIRV | Targets::METAL);
for &(name, targets) in inputs.iter() {
println!("Processing '{}'", name);
let file = fs::read_to_string(format!("{}/{}/{}.wgsl", root, DIR_IN, name))
.expect("Couldn't find wgsl file");
match naga::front::wgsl::parse_str(&file) {
Ok(module) => check_targets(&module, name, targets),
Err(e) => panic!("{}", e),
}
}
}
#[cfg(feature = "spv-in")]