github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

WGSL support for f64, snapshot test for it

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-05-06 00:04:43 -04:00
committed by Dzmitry Malyshau
parent 59d1dcab5b
commit 350ceb383d
24 changed files with 123 additions and 185 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/back/spv/mod.rs
View File

@@ -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<crate::FastHashSet<Capability>>,
}

10
src/front/wgsl/conv.rs
View File

@@ -98,10 +98,18 @@ pub fn map_storage_format(word: &str) -> Result<crate::StorageFormat, Error<'_>>
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,
}
}

19
src/front/wgsl/mod.rs
View File

@@ -1637,23 +1637,10 @@ impl Parser {
type_arena: &mut Arena<crate::Type>,
const_arena: &mut Arena<crate::Constant>,
) -> Result<crate::TypeInner, Error<'a>> {
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 {

22
src/valid/interface.rs
View File

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

7
src/valid/mod.rs
View File

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

1
tests/in/boids.param.ron
View File

@@ -1,5 +1,4 @@
(
spv_flow_dump_prefix: "",
spv_version: (1, 0),
spv_capabilities: [ Shader ],
spv_debug: true,

3
tests/in/collatz.param.ron
View File

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

3
tests/in/control-flow.param.ron
View File

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

3
tests/in/empty.param.ron
View File

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

5
tests/in/extra.param.ron Normal file
View File

@@ -0,0 +1,5 @@
(
god_mode: true,
spv_version: (1, 0),
spv_capabilities: [ Shader ],
)

11
tests/in/extra.wgsl Normal file
View File

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

2
tests/in/image.param.ron
View File

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

1
tests/in/interpolate.param.ron
View File

@@ -3,6 +3,5 @@
spv_capabilities: [ Shader, SampleRateShading ],
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
glsl_desktop_version: Some(400)
)

3
tests/in/operators.param.ron
View File

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

1
tests/in/quad-glsl.param.ron
View File

@@ -3,5 +3,4 @@
spv_capabilities: [ Shader ],
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
)

1
tests/in/quad.param.ron
View File

@@ -3,5 +3,4 @@
spv_capabilities: [ Shader ],
spv_debug: true,
spv_adjust_coordinate_space: true,
msl_custom: false,
)

2
tests/in/shadow.param.ron
View File

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

8
tests/in/texture-array.param.ron
View File

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

18
tests/in/texture-array.wgsl
View File

@@ -1,18 +0,0 @@
[[group(0), binding(0)]] var texture0: texture_2d<f32>;
[[group(0), binding(1)]] var texture1: texture_2d<f32>;
[[group(0), binding(2)]] var sampler: sampler;
[[block]]
struct PushConstants {
index: u32;
};
var<push_constant> pc: PushConstants;
[[stage(fragment)]]
fn main([[location(0)]] tex_coord: vec2<f32>) -> [[location(1)]] vec4<f32> {
if (pc.index == 0u) {
return textureSample(texture0, sampler, tex_coord);
} else {
return textureSample(texture1, sampler, tex_coord);
}
}

21
tests/out/extra.msl Normal file
View File

@@ -0,0 +1,21 @@
#include <metal_stdlib>
#include <simd/simd.h>
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 };
}

37
tests/out/extra.spvasm Normal file
View File

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

29
tests/out/texture-array.msl
View File

@@ -1,29 +0,0 @@
#include <metal_stdlib>
#include <simd/simd.h>
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<float, metal::access::sample> texture0_ [[user(fake0)]]
, metal::texture2d<float, metal::access::sample> 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 };
}
}

78
tests/out/texture-array.spvasm
View File

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

21
tests/snapshots.rs
View File

@@ -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<spirv::Capability>,
#[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",