From 350ceb383d4c85e5ae05e5f9bd2b2881917dac2d Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 6 May 2021 00:04:43 -0400 Subject: [PATCH] WGSL support for f64, snapshot test for it --- src/back/spv/mod.rs | 2 + src/front/wgsl/conv.rs | 10 +++- src/front/wgsl/mod.rs | 19 ++------ src/valid/interface.rs | 22 ++++++--- src/valid/mod.rs | 7 ++- tests/in/boids.param.ron | 1 - tests/in/collatz.param.ron | 3 -- tests/in/control-flow.param.ron | 3 -- tests/in/empty.param.ron | 3 -- tests/in/extra.param.ron | 5 ++ tests/in/extra.wgsl | 11 +++++ tests/in/image.param.ron | 2 - tests/in/interpolate.param.ron | 1 - tests/in/operators.param.ron | 3 -- tests/in/quad-glsl.param.ron | 1 - tests/in/quad.param.ron | 1 - tests/in/shadow.param.ron | 2 - tests/in/texture-array.param.ron | 8 ---- tests/in/texture-array.wgsl | 18 -------- tests/out/extra.msl | 21 +++++++++ tests/out/extra.spvasm | 37 +++++++++++++++ tests/out/texture-array.msl | 29 ------------ tests/out/texture-array.spvasm | 78 -------------------------------- tests/snapshots.rs | 21 ++++++--- 24 files changed, 123 insertions(+), 185 deletions(-) create mode 100644 tests/in/extra.param.ron create mode 100644 tests/in/extra.wgsl delete mode 100644 tests/in/texture-array.param.ron delete mode 100644 tests/in/texture-array.wgsl create mode 100644 tests/out/extra.msl create mode 100644 tests/out/extra.spvasm delete mode 100644 tests/out/texture-array.msl delete mode 100644 tests/out/texture-array.spvasm diff --git a/src/back/spv/mod.rs b/src/back/spv/mod.rs index 9679251420..e571c3ba13 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -58,6 +58,8 @@ pub struct Options { /// Configuration flags for the writer. pub flags: WriterFlags, /// Set of SPIR-V allowed capabilities, if provided. + // Note: there is a major bug currently associated with deriving the capabilities. + // We are calling `required_capabilities`, but the semantics of this is broken. pub capabilities: Option>, } diff --git a/src/front/wgsl/conv.rs b/src/front/wgsl/conv.rs index 227327ad5e..f9b79290d1 100644 --- a/src/front/wgsl/conv.rs +++ b/src/front/wgsl/conv.rs @@ -98,10 +98,18 @@ pub fn map_storage_format(word: &str) -> Result> pub fn get_scalar_type(word: &str) -> Option<(crate::ScalarKind, crate::Bytes)> { match word { + "f16" => Some((crate::ScalarKind::Float, 2)), "f32" => Some((crate::ScalarKind::Float, 4)), + "f64" => Some((crate::ScalarKind::Float, 8)), + "i8" => Some((crate::ScalarKind::Sint, 1)), + "i16" => Some((crate::ScalarKind::Sint, 2)), "i32" => Some((crate::ScalarKind::Sint, 4)), + "i64" => Some((crate::ScalarKind::Sint, 8)), + "u8" => Some((crate::ScalarKind::Uint, 1)), + "u16" => Some((crate::ScalarKind::Uint, 2)), "u32" => Some((crate::ScalarKind::Uint, 4)), - "bool" => Some((crate::ScalarKind::Bool, 1)), + "u64" => Some((crate::ScalarKind::Uint, 8)), + "bool" => Some((crate::ScalarKind::Bool, crate::BOOL_WIDTH)), _ => None, } } diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index a9f4f999ba..51ebde4174 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -1637,23 +1637,10 @@ impl Parser { type_arena: &mut Arena, const_arena: &mut Arena, ) -> Result> { + if let Some((kind, width)) = conv::get_scalar_type(word) { + return Ok(crate::TypeInner::Scalar { kind, width }); + } Ok(match word { - "f32" => crate::TypeInner::Scalar { - kind: crate::ScalarKind::Float, - width: 4, - }, - "i32" => crate::TypeInner::Scalar { - kind: crate::ScalarKind::Sint, - width: 4, - }, - "u32" => crate::TypeInner::Scalar { - kind: crate::ScalarKind::Uint, - width: 4, - }, - "bool" => crate::TypeInner::Scalar { - kind: crate::ScalarKind::Bool, - width: crate::BOOL_WIDTH, - }, "vec2" => { let (kind, width) = lexer.next_scalar_generic()?; crate::TypeInner::Vector { diff --git a/src/valid/interface.rs b/src/valid/interface.rs index 711124508e..6a9949a652 100644 --- a/src/valid/interface.rs +++ b/src/valid/interface.rs @@ -1,6 +1,7 @@ use super::{ analyzer::{FunctionInfo, GlobalUse}, - Disalignment, FunctionError, ModuleInfo, ShaderStages, TypeFlags, ValidationFlags, + Capabilities, Disalignment, FunctionError, ModuleInfo, ShaderStages, TypeFlags, + ValidationFlags, }; use crate::arena::{Arena, Handle}; @@ -24,6 +25,8 @@ pub enum GlobalVariableError { required: TypeFlags, seen: TypeFlags, }, + #[error("Capability {0:?} is not supported")] + UnsupportedCapability(Capabilities), #[error("Binding decoration is missing or not applicable")] InvalidBinding, #[error("Alignment requirements for this storage class are not met by {0:?}")] @@ -332,11 +335,18 @@ impl super::Validator { crate::StorageClass::Private | crate::StorageClass::WorkGroup => { (crate::StorageAccess::empty(), TypeFlags::DATA, false) } - crate::StorageClass::PushConstant => ( - crate::StorageAccess::LOAD, - TypeFlags::DATA | TypeFlags::HOST_SHARED, - false, - ), + crate::StorageClass::PushConstant => { + if !self.capabilities.contains(Capabilities::PUSH_CONSTANT) { + return Err(GlobalVariableError::UnsupportedCapability( + Capabilities::PUSH_CONSTANT, + )); + } + ( + crate::StorageAccess::LOAD, + TypeFlags::DATA | TypeFlags::HOST_SHARED, + false, + ) + } }; if !allowed_storage_access.contains(var.storage_access) { diff --git a/src/valid/mod.rs b/src/valid/mod.rs index 558b75bd44..c34eeebd2c 100644 --- a/src/valid/mod.rs +++ b/src/valid/mod.rs @@ -49,11 +49,14 @@ impl Default for ValidationFlags { #[must_use] bitflags::bitflags! { /// Allowed IR capabilities. + #[derive(Default)] #[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] pub struct Capabilities: u8 { - /// Float values with width = 8 - const FLOAT64 = 0x1; + /// Support for `StorageClass:PushConstant`. + const PUSH_CONSTANT = 0x1; + /// Float values with width = 8. + const FLOAT64 = 0x2; } } diff --git a/tests/in/boids.param.ron b/tests/in/boids.param.ron index 450770cac6..16021ec3b1 100644 --- a/tests/in/boids.param.ron +++ b/tests/in/boids.param.ron @@ -1,5 +1,4 @@ ( - spv_flow_dump_prefix: "", spv_version: (1, 0), spv_capabilities: [ Shader ], spv_debug: true, diff --git a/tests/in/collatz.param.ron b/tests/in/collatz.param.ron index 984ca763e0..42c8480e4c 100644 --- a/tests/in/collatz.param.ron +++ b/tests/in/collatz.param.ron @@ -1,8 +1,5 @@ ( - spv_flow_dump_prefix: "", spv_version: (1, 0), spv_capabilities: [ Shader ], spv_debug: true, - spv_adjust_coordinate_space: false, - msl_custom: false, ) diff --git a/tests/in/control-flow.param.ron b/tests/in/control-flow.param.ron index 6de22960da..d6269ab629 100644 --- a/tests/in/control-flow.param.ron +++ b/tests/in/control-flow.param.ron @@ -1,7 +1,4 @@ ( spv_version: (1, 1), spv_capabilities: [ Shader ], - spv_debug: false, - spv_adjust_coordinate_space: false, - msl_custom: false, ) diff --git a/tests/in/empty.param.ron b/tests/in/empty.param.ron index 6de22960da..d6269ab629 100644 --- a/tests/in/empty.param.ron +++ b/tests/in/empty.param.ron @@ -1,7 +1,4 @@ ( spv_version: (1, 1), spv_capabilities: [ Shader ], - spv_debug: false, - spv_adjust_coordinate_space: false, - msl_custom: false, ) diff --git a/tests/in/extra.param.ron b/tests/in/extra.param.ron new file mode 100644 index 0000000000..1767cbe3f5 --- /dev/null +++ b/tests/in/extra.param.ron @@ -0,0 +1,5 @@ +( + god_mode: true, + spv_version: (1, 0), + spv_capabilities: [ Shader ], +) diff --git a/tests/in/extra.wgsl b/tests/in/extra.wgsl new file mode 100644 index 0000000000..11b8a2eb64 --- /dev/null +++ b/tests/in/extra.wgsl @@ -0,0 +1,11 @@ +[[block]] +struct PushConstants { + index: u32; + double: vec2; +}; +var pc: PushConstants; + +[[stage(fragment)]] +fn main([[location(0)]] color: vec4) -> [[location(0)]] vec4 { + return color; +} diff --git a/tests/in/image.param.ron b/tests/in/image.param.ron index 54df971171..7031137559 100644 --- a/tests/in/image.param.ron +++ b/tests/in/image.param.ron @@ -2,6 +2,4 @@ spv_version: (1, 1), spv_capabilities: [ Shader, ImageQuery, Image1D, Sampled1D ], spv_debug: true, - spv_adjust_coordinate_space: false, - msl_custom: false, ) diff --git a/tests/in/interpolate.param.ron b/tests/in/interpolate.param.ron index 804665025c..f0c6a13d56 100644 --- a/tests/in/interpolate.param.ron +++ b/tests/in/interpolate.param.ron @@ -3,6 +3,5 @@ spv_capabilities: [ Shader, SampleRateShading ], spv_debug: true, spv_adjust_coordinate_space: true, - msl_custom: false, glsl_desktop_version: Some(400) ) diff --git a/tests/in/operators.param.ron b/tests/in/operators.param.ron index bd5e5acc29..553f5e75fb 100644 --- a/tests/in/operators.param.ron +++ b/tests/in/operators.param.ron @@ -1,7 +1,4 @@ ( spv_version: (1, 0), spv_capabilities: [ Shader ], - spv_debug: false, - spv_adjust_coordinate_space: false, - msl_custom: false, ) diff --git a/tests/in/quad-glsl.param.ron b/tests/in/quad-glsl.param.ron index efded54c28..c073cf1a50 100644 --- a/tests/in/quad-glsl.param.ron +++ b/tests/in/quad-glsl.param.ron @@ -3,5 +3,4 @@ spv_capabilities: [ Shader ], spv_debug: true, spv_adjust_coordinate_space: true, - msl_custom: false, ) diff --git a/tests/in/quad.param.ron b/tests/in/quad.param.ron index efded54c28..c073cf1a50 100644 --- a/tests/in/quad.param.ron +++ b/tests/in/quad.param.ron @@ -3,5 +3,4 @@ spv_capabilities: [ Shader ], spv_debug: true, spv_adjust_coordinate_space: true, - msl_custom: false, ) diff --git a/tests/in/shadow.param.ron b/tests/in/shadow.param.ron index dfa9978c3d..c6b43089b9 100644 --- a/tests/in/shadow.param.ron +++ b/tests/in/shadow.param.ron @@ -1,8 +1,6 @@ ( - spv_flow_dump_prefix: "", spv_version: (1, 2), spv_capabilities: [ Shader ], spv_debug: true, spv_adjust_coordinate_space: true, - msl_custom: false, ) diff --git a/tests/in/texture-array.param.ron b/tests/in/texture-array.param.ron deleted file mode 100644 index b46594e319..0000000000 --- a/tests/in/texture-array.param.ron +++ /dev/null @@ -1,8 +0,0 @@ -( - spv_flow_dump_prefix: "", - spv_version: (1, 5), - spv_capabilities: [ Shader ], - spv_debug: true, - spv_adjust_coordinate_space: false, - msl_custom: false, -) diff --git a/tests/in/texture-array.wgsl b/tests/in/texture-array.wgsl deleted file mode 100644 index 5f2e09b35f..0000000000 --- a/tests/in/texture-array.wgsl +++ /dev/null @@ -1,18 +0,0 @@ -[[group(0), binding(0)]] var texture0: texture_2d; -[[group(0), binding(1)]] var texture1: texture_2d; -[[group(0), binding(2)]] var sampler: sampler; - -[[block]] -struct PushConstants { - index: u32; -}; -var pc: PushConstants; - -[[stage(fragment)]] -fn main([[location(0)]] tex_coord: vec2) -> [[location(1)]] vec4 { - if (pc.index == 0u) { - return textureSample(texture0, sampler, tex_coord); - } else { - return textureSample(texture1, sampler, tex_coord); - } -} diff --git a/tests/out/extra.msl b/tests/out/extra.msl new file mode 100644 index 0000000000..b2a2363d90 --- /dev/null +++ b/tests/out/extra.msl @@ -0,0 +1,21 @@ +#include +#include + +struct PushConstants { + metal::uint index; + char _pad1[12]; + metal::float2 double1; +}; + +struct main1Input { + metal::float4 color [[user(loc0), center_perspective]]; +}; +struct main1Output { + metal::float4 member [[color(0)]]; +}; +fragment main1Output main1( + main1Input varyings [[stage_in]] +) { + const auto color = varyings.color; + return main1Output { color }; +} diff --git a/tests/out/extra.spvasm b/tests/out/extra.spvasm new file mode 100644 index 0000000000..ec05bdbdf1 --- /dev/null +++ b/tests/out/extra.spvasm @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 20 +OpCapability Shader +OpCapability Float64 +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint Fragment %17 "main" %12 %15 +OpExecutionMode %17 OriginUpperLeft +OpDecorate %6 Block +OpMemberDecorate %6 0 Offset 0 +OpMemberDecorate %6 1 Offset 16 +OpDecorate %12 Location 0 +OpDecorate %15 Location 0 +%2 = OpTypeVoid +%3 = OpTypeInt 32 0 +%5 = OpTypeFloat 64 +%4 = OpTypeVector %5 2 +%6 = OpTypeStruct %3 %4 +%8 = OpTypeFloat 32 +%7 = OpTypeVector %8 4 +%10 = OpTypePointer PushConstant %6 +%9 = OpVariable %10 PushConstant +%13 = OpTypePointer Input %7 +%12 = OpVariable %13 Input +%16 = OpTypePointer Output %7 +%15 = OpVariable %16 Output +%18 = OpTypeFunction %2 +%17 = OpFunction %2 None %18 +%11 = OpLabel +%14 = OpLoad %7 %12 +OpBranch %19 +%19 = OpLabel +OpStore %15 %14 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/texture-array.msl b/tests/out/texture-array.msl deleted file mode 100644 index 229d533579..0000000000 --- a/tests/out/texture-array.msl +++ /dev/null @@ -1,29 +0,0 @@ -#include -#include - -struct PushConstants { - metal::uint index; -}; - -struct main1Input { - metal::float2 tex_coord [[user(loc0), center_perspective]]; -}; -struct main1Output { - metal::float4 member [[color(1)]]; -}; -fragment main1Output main1( - main1Input varyings [[stage_in]] -, metal::texture2d texture0_ [[user(fake0)]] -, metal::texture2d texture1_ [[user(fake0)]] -, metal::sampler sampler [[user(fake0)]] -, constant PushConstants& pc [[user(fake0)]] -) { - const auto tex_coord = varyings.tex_coord; - if (pc.index == 0u) { - metal::float4 _e9 = texture0_.sample(sampler, tex_coord); - return main1Output { _e9 }; - } else { - metal::float4 _e10 = texture1_.sample(sampler, tex_coord); - return main1Output { _e10 }; - } -} diff --git a/tests/out/texture-array.spvasm b/tests/out/texture-array.spvasm deleted file mode 100644 index faefb5f6ed..0000000000 --- a/tests/out/texture-array.spvasm +++ /dev/null @@ -1,78 +0,0 @@ -; SPIR-V -; Version: 1.5 -; Generator: rspirv -; Bound: 43 -OpCapability Shader -%1 = OpExtInstImport "GLSL.std.450" -OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %24 "main" %19 %22 -OpExecutionMode %24 OriginUpperLeft -OpSource GLSL 450 -OpName %8 "PushConstants" -OpMemberName %8 0 "index" -OpName %11 "texture0" -OpName %13 "texture1" -OpName %14 "sampler" -OpName %16 "pc" -OpName %19 "tex_coord" -OpName %24 "main" -OpDecorate %8 Block -OpMemberDecorate %8 0 Offset 0 -OpDecorate %11 DescriptorSet 0 -OpDecorate %11 Binding 0 -OpDecorate %13 DescriptorSet 0 -OpDecorate %13 Binding 1 -OpDecorate %14 DescriptorSet 0 -OpDecorate %14 Binding 2 -OpDecorate %19 Location 0 -OpDecorate %22 Location 1 -%2 = OpTypeVoid -%4 = OpTypeInt 32 0 -%3 = OpConstant %4 0 -%6 = OpTypeFloat 32 -%5 = OpTypeImage %6 2D 0 0 0 1 Unknown -%7 = OpTypeSampler -%8 = OpTypeStruct %4 -%9 = OpTypeVector %6 2 -%10 = OpTypeVector %6 4 -%12 = OpTypePointer UniformConstant %5 -%11 = OpVariable %12 UniformConstant -%13 = OpVariable %12 UniformConstant -%15 = OpTypePointer UniformConstant %7 -%14 = OpVariable %15 UniformConstant -%17 = OpTypePointer PushConstant %8 -%16 = OpVariable %17 PushConstant -%20 = OpTypePointer Input %9 -%19 = OpVariable %20 Input -%23 = OpTypePointer Output %10 -%22 = OpVariable %23 Output -%25 = OpTypeFunction %2 -%30 = OpTypePointer PushConstant %4 -%33 = OpTypeBool -%38 = OpTypeSampledImage %5 -%24 = OpFunction %2 None %25 -%18 = OpLabel -%21 = OpLoad %9 %19 -%26 = OpLoad %5 %11 -%27 = OpLoad %5 %13 -%28 = OpLoad %7 %14 -OpBranch %29 -%29 = OpLabel -%31 = OpAccessChain %30 %16 %3 -%32 = OpLoad %4 %31 -%34 = OpIEqual %33 %32 %3 -OpSelectionMerge %35 None -OpBranchConditional %34 %36 %37 -%36 = OpLabel -%39 = OpSampledImage %38 %26 %28 -%40 = OpImageSampleImplicitLod %10 %39 %21 -OpStore %22 %40 -OpReturn -%37 = OpLabel -%41 = OpSampledImage %38 %27 %28 -%42 = OpImageSampleImplicitLod %10 %41 %21 -OpStore %22 %42 -OpReturn -%35 = OpLabel -OpReturn -OpFunctionEnd \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index e0aabd8c26..b628075ce5 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -21,18 +21,23 @@ bitflags::bitflags! { #[derive(Default, serde::Deserialize)] struct Parameters { + #[serde(default)] + god_mode: bool, #[cfg_attr(not(feature = "spv-out"), allow(dead_code))] spv_version: (u8, u8), #[cfg_attr(not(feature = "spv-out"), allow(dead_code))] spv_capabilities: naga::FastHashSet, #[cfg_attr(not(feature = "spv-out"), allow(dead_code))] + #[serde(default)] spv_debug: bool, #[cfg_attr(not(feature = "spv-out"), allow(dead_code))] + #[serde(default)] spv_adjust_coordinate_space: bool, #[cfg(all(feature = "deserialize", feature = "msl-out"))] #[serde(default)] msl: naga::back::msl::Options, #[cfg(all(not(feature = "deserialize"), feature = "msl-out"))] + #[serde(default)] msl_custom: bool, #[cfg_attr(not(feature = "glsl-out"), allow(dead_code))] #[serde(default)] @@ -46,12 +51,14 @@ fn check_targets(module: &naga::Module, name: &str, targets: Targets) { Ok(string) => ron::de::from_str(&string).expect("Couldn't find param file"), Err(_) => Parameters::default(), }; - let info = naga::valid::Validator::new( - naga::valid::ValidationFlags::all(), - naga::valid::Capabilities::empty(), - ) - .validate(module) - .unwrap(); + let capabilities = if params.god_mode { + naga::valid::Capabilities::all() + } else { + naga::valid::Capabilities::empty() + }; + let info = naga::valid::Validator::new(naga::valid::ValidationFlags::all(), capabilities) + .validate(module) + .unwrap(); let dest = PathBuf::from(root).join(DIR_OUT).join(name); @@ -248,7 +255,7 @@ fn convert_wgsl() { ), ("shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL), ("image", Targets::SPIRV | Targets::METAL), - ("texture-array", Targets::SPIRV | Targets::METAL), + ("extra", Targets::SPIRV | Targets::METAL), ("operators", Targets::SPIRV | Targets::METAL | Targets::GLSL), ( "interpolate",