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

Update WGSL grammar for pointer access. (#1312)

Browse Source 
* Update WGSL grammar for pointer access.

Comes with a small test, which revealed a number of issues in the backends.

* Validate pointer arguments to functions to only have function/private/workgroup classes.

Comes with a small test. Also, "pointer-access.spv" test is temporarily disabled.
This commit is contained in:
Dzmitry Malyshau
2021-09-27 18:49:28 -04:00
committed by GitHub
parent 38d74a7f0a
commit 21324b8fea
19 changed files with 312 additions and 197 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
src/back/glsl/mod.rs
View File

@@ -715,7 +715,9 @@ impl<'a, W: Write> Writer<'a, W> {
TypeInner::Pointer { .. }
| TypeInner::Struct { .. }
| TypeInner::Image { .. }
| TypeInner::Sampler { .. } => unreachable!(),
| TypeInner::Sampler { .. } => {
return Err(Error::Custom(format!("Unable to write type {:?}", inner)))
}
}
Ok(())
@@ -1332,7 +1334,14 @@ impl<'a, W: Write> Writer<'a, W> {
// This is where we can generate intermediate constants for some expression types.
Statement::Emit(ref range) => {
for handle in range.clone() {
let expr_name = if let Some(name) = ctx.named_expressions.get(&handle) {
let info = &ctx.info[handle];
let ptr_class = info.ty.inner_with(&self.module.types).pointer_class();
let expr_name = if ptr_class.is_some() {
// GLSL can't save a pointer-valued expression in a variable,
// but we shouldn't ever need to: they should never be named expressions,
// and none of the expression types flagged by bake_ref_count can be pointer-valued.
None
} else if let Some(name) = ctx.named_expressions.get(&handle) {
// Front end provides names for all variables at the start of writing.
// But we write them to step by step. We need to recache them
// Otherwise, we could accidentally write variable name instead of full expression.
@@ -1340,7 +1349,7 @@ impl<'a, W: Write> Writer<'a, W> {
Some(self.namer.call_unique(name))
} else {
let min_ref_count = ctx.expressions[handle].bake_ref_count();
if min_ref_count <= ctx.info[handle].ref_count {
if min_ref_count <= info.ref_count {
Some(format!("{}{}", super::BAKE_PREFIX, handle.index()))
} else {
None

11
src/back/hlsl/writer.rs
View File

@@ -1057,7 +1057,14 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
match *stmt {
Statement::Emit(ref range) => {
for handle in range.clone() {
let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) {
let info = &func_ctx.info[handle];
let ptr_class = info.ty.inner_with(&module.types).pointer_class();
let expr_name = if ptr_class.is_some() {
// HLSL can't save a pointer-valued expression in a variable,
// but we shouldn't ever need to: they should never be named expressions,
// and none of the expression types flagged by bake_ref_count can be pointer-valued.
None
} else if let Some(name) = func_ctx.named_expressions.get(&handle) {
// Front end provides names for all variables at the start of writing.
// But we write them to step by step. We need to recache them
// Otherwise, we could accidentally write variable name instead of full expression.
@@ -1065,7 +1072,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
Some(self.namer.call_unique(name))
} else {
let min_ref_count = func_ctx.expressions[handle].bake_ref_count();
if min_ref_count <= func_ctx.info[handle].ref_count {
if min_ref_count <= info.ref_count {
Some(format!("_expr{}", handle.index()))
} else {
None

13
src/back/msl/writer.rs
View File

@@ -1355,7 +1355,7 @@ impl<W: Write> Writer<W> {
)?;
}
TypeResolution::Value(ref other) => {
log::error!("Type {:?} isn't a known local", other);
log::warn!("Type {:?} isn't a known local", other); //TEMP!
return Err(Error::FeatureNotImplemented("weird local type".to_string()));
}
}
@@ -1383,7 +1383,14 @@ impl<W: Write> Writer<W> {
match *statement {
crate::Statement::Emit(ref range) => {
for handle in range.clone() {
let expr_name = if let Some(name) =
let info = &context.expression.info[handle];
let ptr_class = info
.ty
.inner_with(&context.expression.module.types)
.pointer_class();
let expr_name = if ptr_class.is_some() {
None // don't bake pointer expressions (just yet)
} else if let Some(name) =
context.expression.function.named_expressions.get(&handle)
{
// Front end provides names for all variables at the start of writing.
@@ -1394,7 +1401,7 @@ impl<W: Write> Writer<W> {
} else {
let min_ref_count =
context.expression.function.expressions[handle].bake_ref_count();
if min_ref_count <= context.expression.info[handle].ref_count {
if min_ref_count <= info.ref_count {
Some(format!("{}{}", back::BAKE_PREFIX, handle.index()))
} else {
None

20
src/back/wgsl/writer.rs
View File

@@ -525,10 +525,13 @@ impl<W: Write> Writer<W> {
"storage_",
"",
storage_format_str(format),
if access.contains(crate::StorageAccess::STORE) {
",write"
if access.contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE)
{
",read_write"
} else if access.contains(crate::StorageAccess::LOAD) {
",read"
} else {
""
",write"
},
),
};
@@ -639,6 +642,7 @@ impl<W: Write> Writer<W> {
inner
)));
}
write!(self.out, ">")?;
}
_ => {
return Err(Error::Unimplemented(format!(
@@ -666,6 +670,7 @@ impl<W: Write> Writer<W> {
match *stmt {
Statement::Emit(ref range) => {
for handle in range.clone() {
let info = &func_ctx.info[handle];
let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) {
// Front end provides names for all variables at the start of writing.
// But we write them to step by step. We need to recache them
@@ -682,8 +687,7 @@ impl<W: Write> Writer<W> {
| Expression::ImageSample { .. } => true,
_ => false,
};
if min_ref_count <= func_ctx.info[handle].ref_count || required_baking_expr
{
if min_ref_count <= info.ref_count || required_baking_expr {
// If expression contains unsupported builtin we should skip it
if let Expression::Load { pointer } = func_ctx.expressions[handle] {
if let Expression::AccessIndex { base, index } =
@@ -809,8 +813,8 @@ impl<W: Write> Writer<W> {
}
let func_name = &self.names[&NameKey::Function(function)];
write!(self.out, "{}(", func_name)?;
for (index, argument) in arguments.iter().enumerate() {
self.write_expr(module, *argument, func_ctx)?;
for (index, &argument) in arguments.iter().enumerate() {
self.write_expr(module, argument, func_ctx)?;
// Only write a comma if isn't the last element
if index != arguments.len().saturating_sub(1) {
// The leading space is for readability only
@@ -1199,14 +1203,12 @@ impl<W: Write> Writer<W> {
self.write_expr(module, right, func_ctx)?;
write!(self.out, ")")?;
}
// TODO: copy-paste from glsl-out
Expression::Access { base, index } => {
self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
write!(self.out, "[")?;
self.write_expr(module, index, func_ctx)?;
write!(self.out, "]")?
}
// TODO: copy-paste from glsl-out
Expression::AccessIndex { base, index } => {
let base_ty_res = &func_ctx.info[base].ty;
let mut resolved = base_ty_res.inner_with(&module.types);

24
src/front/wgsl/lexer.rs
View File

@@ -558,24 +558,24 @@ impl<'a> Lexer<'a> {
Ok(pair)
}
// TODO relocate storage texture specifics
pub(super) fn next_storage_access(&mut self) -> Result<crate::StorageAccess, Error<'a>> {
let (ident, span) = self.next_ident_with_span()?;
match ident {
"read" => Ok(crate::StorageAccess::LOAD),
"write" => Ok(crate::StorageAccess::STORE),
"read_write" => Ok(crate::StorageAccess::LOAD | crate::StorageAccess::STORE),
_ => Err(Error::UnknownAccess(span)),
}
}
pub(super) fn next_format_generic(
&mut self,
) -> Result<(crate::StorageFormat, crate::StorageAccess), Error<'a>> {
self.expect(Token::Paren('<'))?;
let (ident, ident_span) = self.next_ident_with_span()?;
let format = conv::map_storage_format(ident, ident_span)?;
let access = if self.skip(Token::Separator(',')) {
let (raw, span) = self.next_ident_with_span()?;
match raw {
"read" => crate::StorageAccess::LOAD,
"write" => crate::StorageAccess::STORE,
"read_write" => crate::StorageAccess::all(),
_ => return Err(Error::UnknownAccess(span)),
}
} else {
crate::StorageAccess::LOAD
};
self.expect(Token::Separator(','))?;
let access = self.next_storage_access()?;
self.expect(Token::Paren('>'))?;
Ok((format, access))
}

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

@@ -2587,13 +2587,7 @@ impl Parser {
class = Some(match class_str {
"storage" => {
let access = if lexer.skip(Token::Separator(',')) {
let (ident, span) = lexer.next_ident_with_span()?;
match ident {
"read" => crate::StorageAccess::LOAD,
"write" => crate::StorageAccess::STORE,
"read_write" => crate::StorageAccess::all(),
_ => return Err(Error::UnknownAccess(span)),
}
lexer.next_storage_access()?
} else {
// defaulting to `read`
crate::StorageAccess::LOAD
@@ -2836,9 +2830,16 @@ impl Parser {
"ptr" => {
lexer.expect_generic_paren('<')?;
let (ident, span) = lexer.next_ident_with_span()?;
let class = conv::map_storage_class(ident, span)?;
let mut class = conv::map_storage_class(ident, span)?;
lexer.expect(Token::Separator(','))?;
let (base, _access) = self.parse_type_decl(lexer, None, type_arena, const_arena)?;
if let crate::StorageClass::Storage { ref mut access } = class {
*access = if lexer.skip(Token::Separator(',')) {
lexer.next_storage_access()?
} else {
crate::StorageAccess::LOAD
};
}
lexer.expect_generic_paren('>')?;
crate::TypeInner::Pointer { base, class }
}

4
src/front/wgsl/tests.rs
View File

@@ -92,7 +92,7 @@ fn parse_types() {
parse_str("var t: texture_cube_array<i32>;").unwrap();
parse_str("var t: texture_multisampled_2d<u32>;").unwrap();
parse_str("var t: texture_storage_1d<rgba8uint,write>;").unwrap();
parse_str("var t: texture_storage_3d<r32float>;").unwrap();
parse_str("var t: texture_storage_3d<r32float,read>;").unwrap();
}
#[test]
@@ -305,7 +305,7 @@ fn parse_texture_load() {
.unwrap();
parse_str(
"
var t: texture_storage_1d_array<r32float>;
var t: texture_storage_1d_array<r32float,read>;
fn foo() {
let r: vec4<f32> = textureLoad(t, 10, 2);
}

19
src/valid/function.rs
View File

@@ -76,6 +76,12 @@ pub enum FunctionError {
},
#[error("Argument '{name}' at index {index} has a type that can't be passed into functions.")]
InvalidArgumentType { index: usize, name: String },
#[error("Argument '{name}' at index {index} is a pointer of class {class:?}, which can't be passed into functions.")]
InvalidArgumentPointerClass {
index: usize,
name: String,
class: crate::StorageClass,
},
#[error("There are instructions after `return`/`break`/`continue`")]
InstructionsAfterReturn,
#[error("The `break` is used outside of a `loop` or `switch` context")]
@@ -696,6 +702,19 @@ impl super::Validator {
name: argument.name.clone().unwrap_or_default(),
});
}
match module.types[argument.ty].inner.pointer_class() {
Some(crate::StorageClass::Private)
| Some(crate::StorageClass::Function)
| Some(crate::StorageClass::WorkGroup)
| None => {}
Some(other) => {
return Err(FunctionError::InvalidArgumentPointerClass {
index,
name: argument.name.clone().unwrap_or_default(),
class: other,
})
}
}
}
self.valid_expression_set.clear();

8
tests/in/access.wgsl
View File

@@ -11,6 +11,10 @@ struct Bar {
[[group(0), binding(0)]]
var<storage,read_write> bar: Bar;
fn read_from_private(foo: ptr<function, f32>) -> f32 {
return *foo;
}
[[stage(vertex)]]
fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
var foo: f32 = 0.0;
@@ -25,6 +29,10 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let b = bar.matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 2u];
// test pointer types
let pointer: ptr<storage, i32, read_write> = &bar.data[0];
let foo_value = read_from_private(&foo);
// test storage stores
bar.matrix[1].z = 1.0;
bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));

2
tests/in/image.wgsl
View File

@@ -5,7 +5,7 @@ var image_multisampled_src: texture_multisampled_2d<u32>;
[[group(0), binding(4)]]
var image_depth_multisampled_src: texture_depth_multisampled_2d;
[[group(0), binding(1)]]
var image_storage_src: texture_storage_2d<rgba8uint>;
var image_storage_src: texture_storage_2d<rgba8uint, read>;
[[group(0), binding(5)]]
var image_array_src: texture_2d_array<u32>;
[[group(0), binding(6)]]

5
tests/out/glsl/access.atomics.Compute.glsl
View File

@@ -13,6 +13,11 @@ buffer Bar_block_0Cs {
} _group_0_binding_0;
float read_from_private(inout float foo2) {
float _e2 = foo2;
return _e2;
}
void main() {
int tmp = 0;
int value = _group_0_binding_0.atom;

6
tests/out/glsl/access.foo.Vertex.glsl
View File

@@ -11,6 +11,11 @@ buffer Bar_block_0Vs {
} _group_0_binding_0;
float read_from_private(inout float foo2) {
float _e2 = foo2;
return _e2;
}
void main() {
uint vi = uint(gl_VertexID);
float foo1 = 0.0;
@@ -21,6 +26,7 @@ void main() {
uvec2 arr[2] = _group_0_binding_0.arr;
float b = _group_0_binding_0.matrix[3][0];
int a = _group_0_binding_0.data[(uint(_group_0_binding_0.data.length()) - 2u)];
float _e25 = read_from_private(foo1);
_group_0_binding_0.matrix[1][2] = 1.0;
_group_0_binding_0.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0));
_group_0_binding_0.arr = uvec2[2](uvec2(0u), uvec2(1u));

7
tests/out/hlsl/access.hlsl
View File

@@ -1,6 +1,12 @@
RWByteAddressBuffer bar : register(u0);
float read_from_private(inout float foo2)
{
float _expr2 = foo2;
return _expr2;
}
uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
{
uint ret;
@@ -19,6 +25,7 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
uint2 arr[2] = {asuint(bar.Load2(72+0)), asuint(bar.Load2(72+8))};
float b = asfloat(bar.Load(0+48+0));
int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 88) / 4) - 2u)*4+88));
const float _e25 = read_from_private(foo1);
bar.Store(8+16+0, asuint(1.0));
{
float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx));

14
tests/out/msl/access.msl
View File

@@ -17,10 +17,17 @@ struct Bar {
type3 arr;
type5 data;
};
struct type9 {
struct type11 {
int inner[5];
};
float read_from_private(
thread float& foo2
) {
float _e2 = foo2;
return _e2;
}
struct fooInput {
};
struct fooOutput {
@@ -32,17 +39,18 @@ vertex fooOutput foo(
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
float foo1 = 0.0;
type9 c;
type11 c;
float baz = foo1;
foo1 = 1.0;
metal::float4x4 matrix = bar.matrix;
type3 arr = bar.arr;
float b = bar.matrix[3].x;
int a = bar.data[(1 + (_buffer_sizes.size0 - 88 - 4) / 4) - 2u];
float _e25 = read_from_private(foo1);
bar.matrix[1].z = 1.0;
bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(3.0));
for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type3 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type9 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type11 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
c.inner[vi + 1u] = 42;
int value = c.inner[vi];
return fooOutput { matrix * static_cast<metal::float4>(metal::int4(value)) };

296
tests/out/spv/access.spvasm
View File

@@ -1,41 +1,42 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 106
; Bound: 114
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %39 "foo" %34 %37
OpEntryPoint GLCompute %83 "atomics"
OpExecutionMode %83 LocalSize 1 1 1
OpEntryPoint Vertex %47 "foo" %42 %45
OpEntryPoint GLCompute %91 "atomics"
OpExecutionMode %91 LocalSize 1 1 1
OpSource GLSL 450
OpMemberName %25 0 "matrix"
OpMemberName %25 1 "atom"
OpMemberName %25 2 "arr"
OpMemberName %25 3 "data"
OpName %25 "Bar"
OpName %27 "bar"
OpName %29 "foo"
OpName %31 "c"
OpName %34 "vi"
OpName %39 "foo"
OpName %81 "tmp"
OpName %83 "atomics"
OpDecorate %23 ArrayStride 8
OpDecorate %24 ArrayStride 4
OpDecorate %25 Block
OpMemberDecorate %25 0 Offset 0
OpMemberDecorate %25 0 ColMajor
OpMemberDecorate %25 0 MatrixStride 16
OpMemberDecorate %25 1 Offset 64
OpMemberDecorate %25 2 Offset 72
OpMemberDecorate %25 3 Offset 88
OpDecorate %26 ArrayStride 4
OpDecorate %27 DescriptorSet 0
OpDecorate %27 Binding 0
OpDecorate %34 BuiltIn VertexIndex
OpDecorate %37 BuiltIn Position
OpMemberName %26 0 "matrix"
OpMemberName %26 1 "atom"
OpMemberName %26 2 "arr"
OpMemberName %26 3 "data"
OpName %26 "Bar"
OpName %30 "bar"
OpName %34 "read_from_private"
OpName %38 "foo"
OpName %39 "c"
OpName %42 "vi"
OpName %47 "foo"
OpName %89 "tmp"
OpName %91 "atomics"
OpDecorate %24 ArrayStride 8
OpDecorate %25 ArrayStride 4
OpDecorate %26 Block
OpMemberDecorate %26 0 Offset 0
OpMemberDecorate %26 0 ColMajor
OpMemberDecorate %26 0 MatrixStride 16
OpMemberDecorate %26 1 Offset 64
OpMemberDecorate %26 2 Offset 72
OpMemberDecorate %26 3 Offset 88
OpDecorate %29 ArrayStride 4
OpDecorate %30 DescriptorSet 0
OpDecorate %30 Binding 0
OpDecorate %42 BuiltIn VertexIndex
OpDecorate %45 BuiltIn Position
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 2
@@ -45,120 +46,131 @@ OpDecorate %37 BuiltIn Position
%9 = OpTypeInt 32 0
%8 = OpConstant %9 3
%10 = OpConstant %9 2
%11 = OpConstant %4 1
%12 = OpConstant %6 2.0
%13 = OpConstant %6 3.0
%14 = OpConstant %9 0
%15 = OpConstant %9 1
%16 = OpConstant %4 5
%17 = OpConstant %4 3
%18 = OpConstant %4 4
%19 = OpConstant %4 42
%21 = OpTypeVector %6 4
%20 = OpTypeMatrix %21 4
%22 = OpTypeVector %9 2
%23 = OpTypeArray %22 %3
%24 = OpTypeRuntimeArray %4
%25 = OpTypeStruct %20 %4 %23 %24
%26 = OpTypeArray %4 %16
%28 = OpTypePointer StorageBuffer %25
%27 = OpVariable %28 StorageBuffer
%30 = OpTypePointer Function %6
%32 = OpTypePointer Function %26
%35 = OpTypePointer Input %9
%34 = OpVariable %35 Input
%38 = OpTypePointer Output %21
%37 = OpVariable %38 Output
%40 = OpTypeFunction %2
%43 = OpTypePointer StorageBuffer %20
%46 = OpTypePointer StorageBuffer %23
%49 = OpTypePointer StorageBuffer %21
%50 = OpTypePointer StorageBuffer %6
%53 = OpTypePointer StorageBuffer %24
%56 = OpTypePointer StorageBuffer %4
%73 = OpTypePointer Function %4
%77 = OpTypeVector %4 4
%85 = OpTypePointer StorageBuffer %4
%88 = OpConstant %9 64
%39 = OpFunction %2 None %40
%33 = OpLabel
%29 = OpVariable %30 Function %5
%31 = OpVariable %32 Function
%36 = OpLoad %9 %34
OpBranch %41
%11 = OpConstant %4 0
%12 = OpConstant %4 1
%13 = OpConstant %6 2.0
%14 = OpConstant %6 3.0
%15 = OpConstant %9 0
%16 = OpConstant %9 1
%17 = OpConstant %4 5
%18 = OpConstant %4 3
%19 = OpConstant %4 4
%20 = OpConstant %4 42
%22 = OpTypeVector %6 4
%21 = OpTypeMatrix %22 4
%23 = OpTypeVector %9 2
%24 = OpTypeArray %23 %3
%25 = OpTypeRuntimeArray %4
%26 = OpTypeStruct %21 %4 %24 %25
%27 = OpTypePointer Function %6
%28 = OpTypePointer StorageBuffer %4
%29 = OpTypeArray %4 %17
%31 = OpTypePointer StorageBuffer %26
%30 = OpVariable %31 StorageBuffer
%35 = OpTypeFunction %6 %27
%40 = OpTypePointer Function %29
%43 = OpTypePointer Input %9
%42 = OpVariable %43 Input
%46 = OpTypePointer Output %22
%45 = OpVariable %46 Output
%48 = OpTypeFunction %2
%51 = OpTypePointer StorageBuffer %21
%54 = OpTypePointer StorageBuffer %24
%57 = OpTypePointer StorageBuffer %22
%58 = OpTypePointer StorageBuffer %6
%61 = OpTypePointer StorageBuffer %25
%81 = OpTypePointer Function %4
%85 = OpTypeVector %4 4
%93 = OpTypePointer StorageBuffer %4
%96 = OpConstant %9 64
%34 = OpFunction %6 None %35
%33 = OpFunctionParameter %27
%32 = OpLabel
OpBranch %36
%36 = OpLabel
%37 = OpLoad %6 %33
OpReturnValue %37
OpFunctionEnd
%47 = OpFunction %2 None %48
%41 = OpLabel
%42 = OpLoad %6 %29
OpStore %29 %7
%44 = OpAccessChain %43 %27 %14
%45 = OpLoad %20 %44
%47 = OpAccessChain %46 %27 %10
%48 = OpLoad %23 %47
%51 = OpAccessChain %50 %27 %14 %8 %14
%52 = OpLoad %6 %51
%54 = OpArrayLength %9 %27 3
%55 = OpISub %9 %54 %10
%57 = OpAccessChain %56 %27 %8 %55
%58 = OpLoad %4 %57
%59 = OpAccessChain %50 %27 %14 %15 %10
OpStore %59 %7
%60 = OpCompositeConstruct %21 %5 %5 %5 %5
%61 = OpCompositeConstruct %21 %7 %7 %7 %7
%62 = OpCompositeConstruct %21 %12 %12 %12 %12
%63 = OpCompositeConstruct %21 %13 %13 %13 %13
%64 = OpCompositeConstruct %20 %60 %61 %62 %63
%65 = OpAccessChain %43 %27 %14
OpStore %65 %64
%66 = OpCompositeConstruct %22 %14 %14
%67 = OpCompositeConstruct %22 %15 %15
%68 = OpCompositeConstruct %23 %66 %67
%69 = OpAccessChain %46 %27 %10
OpStore %69 %68
%70 = OpConvertFToS %4 %52
%71 = OpCompositeConstruct %26 %58 %70 %17 %18 %16
OpStore %31 %71
%72 = OpIAdd %9 %36 %15
%74 = OpAccessChain %73 %31 %72
OpStore %74 %19
%75 = OpAccessChain %73 %31 %36
%76 = OpLoad %4 %75
%78 = OpCompositeConstruct %77 %76 %76 %76 %76
%79 = OpConvertSToF %21 %78
%80 = OpMatrixTimesVector %21 %45 %79
OpStore %37 %80
%38 = OpVariable %27 Function %5
%39 = OpVariable %40 Function
%44 = OpLoad %9 %42
OpBranch %49
%49 = OpLabel
%50 = OpLoad %6 %38
OpStore %38 %7
%52 = OpAccessChain %51 %30 %15
%53 = OpLoad %21 %52
%55 = OpAccessChain %54 %30 %10
%56 = OpLoad %24 %55
%59 = OpAccessChain %58 %30 %15 %8 %15
%60 = OpLoad %6 %59
%62 = OpArrayLength %9 %30 3
%63 = OpISub %9 %62 %10
%64 = OpAccessChain %28 %30 %8 %63
%65 = OpLoad %4 %64
%66 = OpFunctionCall %6 %34 %38
%67 = OpAccessChain %58 %30 %15 %16 %10
OpStore %67 %7
%68 = OpCompositeConstruct %22 %5 %5 %5 %5
%69 = OpCompositeConstruct %22 %7 %7 %7 %7
%70 = OpCompositeConstruct %22 %13 %13 %13 %13
%71 = OpCompositeConstruct %22 %14 %14 %14 %14
%72 = OpCompositeConstruct %21 %68 %69 %70 %71
%73 = OpAccessChain %51 %30 %15
OpStore %73 %72
%74 = OpCompositeConstruct %23 %15 %15
%75 = OpCompositeConstruct %23 %16 %16
%76 = OpCompositeConstruct %24 %74 %75
%77 = OpAccessChain %54 %30 %10
OpStore %77 %76
%78 = OpConvertFToS %4 %60
%79 = OpCompositeConstruct %29 %65 %78 %18 %19 %17
OpStore %39 %79
%80 = OpIAdd %9 %44 %16
%82 = OpAccessChain %81 %39 %80
OpStore %82 %20
%83 = OpAccessChain %81 %39 %44
%84 = OpLoad %4 %83
%86 = OpCompositeConstruct %85 %84 %84 %84 %84
%87 = OpConvertSToF %22 %86
%88 = OpMatrixTimesVector %22 %53 %87
OpStore %45 %88
OpReturn
OpFunctionEnd
%83 = OpFunction %2 None %40
%82 = OpLabel
%81 = OpVariable %73 Function
OpBranch %84
%84 = OpLabel
%86 = OpAccessChain %85 %27 %15
%87 = OpAtomicLoad %4 %86 %11 %88
%90 = OpAccessChain %85 %27 %15
%89 = OpAtomicIAdd %4 %90 %11 %88 %16
OpStore %81 %89
%92 = OpAccessChain %85 %27 %15
%91 = OpAtomicISub %4 %92 %11 %88 %16
OpStore %81 %91
%94 = OpAccessChain %85 %27 %15
%93 = OpAtomicAnd %4 %94 %11 %88 %16
OpStore %81 %93
%96 = OpAccessChain %85 %27 %15
%95 = OpAtomicOr %4 %96 %11 %88 %16
OpStore %81 %95
%98 = OpAccessChain %85 %27 %15
%97 = OpAtomicXor %4 %98 %11 %88 %16
OpStore %81 %97
%100 = OpAccessChain %85 %27 %15
%99 = OpAtomicSMin %4 %100 %11 %88 %16
OpStore %81 %99
%102 = OpAccessChain %85 %27 %15
%101 = OpAtomicSMax %4 %102 %11 %88 %16
OpStore %81 %101
%104 = OpAccessChain %85 %27 %15
%103 = OpAtomicExchange %4 %104 %11 %88 %16
OpStore %81 %103
%105 = OpAccessChain %85 %27 %15
OpAtomicStore %105 %11 %88 %87
%91 = OpFunction %2 None %48
%90 = OpLabel
%89 = OpVariable %81 Function
OpBranch %92
%92 = OpLabel
%94 = OpAccessChain %93 %30 %16
%95 = OpAtomicLoad %4 %94 %12 %96
%98 = OpAccessChain %93 %30 %16
%97 = OpAtomicIAdd %4 %98 %12 %96 %17
OpStore %89 %97
%100 = OpAccessChain %93 %30 %16
%99 = OpAtomicISub %4 %100 %12 %96 %17
OpStore %89 %99
%102 = OpAccessChain %93 %30 %16
%101 = OpAtomicAnd %4 %102 %12 %96 %17
OpStore %89 %101
%104 = OpAccessChain %93 %30 %16
%103 = OpAtomicOr %4 %104 %12 %96 %17
OpStore %89 %103
%106 = OpAccessChain %93 %30 %16
%105 = OpAtomicXor %4 %106 %12 %96 %17
OpStore %89 %105
%108 = OpAccessChain %93 %30 %16
%107 = OpAtomicSMin %4 %108 %12 %96 %17
OpStore %89 %107
%110 = OpAccessChain %93 %30 %16
%109 = OpAtomicSMax %4 %110 %12 %96 %17
OpStore %89 %109
%112 = OpAccessChain %93 %30 %16
%111 = OpAtomicExchange %4 %112 %12 %96 %17
OpStore %89 %111
%113 = OpAccessChain %93 %30 %16
OpAtomicStore %113 %12 %96 %95
OpReturn
OpFunctionEnd

7
tests/out/wgsl/access.wgsl
View File

@@ -9,6 +9,11 @@ struct Bar {
[[group(0), binding(0)]]
var<storage, read_write> bar: Bar;
fn read_from_private(foo2: ptr<function, f32>) -> f32 {
let e2: f32 = (*foo2);
return e2;
}
[[stage(vertex)]]
fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
var foo1: f32 = 0.0;
@@ -20,6 +25,8 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let arr: array<vec2<u32>,2> = bar.arr;
let b: f32 = bar.matrix[3][0];
let a: i32 = bar.data[(arrayLength((&bar.data)) - 2u)];
let pointer1: ptr<storage, i32, read_write> = (&bar.data[0]);
let e25: f32 = read_from_private((&foo1));
bar.matrix[1][2] = 1.0;
bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));
bar.arr = array<vec2<u32>,2>(vec2<u32>(0u), vec2<u32>(1u));

4
tests/out/wgsl/image.wgsl
View File

@@ -5,11 +5,11 @@ var image_multisampled_src: texture_multisampled_2d<u32>;
[[group(0), binding(4)]]
var image_depth_multisampled_src: texture_depth_multisampled_2d;
[[group(0), binding(1)]]
var image_storage_src: texture_storage_2d<rgba8uint>;
var image_storage_src: texture_storage_2d<rgba8uint,read>;
[[group(0), binding(5)]]
var image_array_src: texture_2d_array<u32>;
[[group(0), binding(6)]]
var image_dup_src: texture_storage_1d<r32uint>;
var image_dup_src: texture_storage_1d<r32uint,read>;
[[group(0), binding(2)]]
var image_dst: texture_storage_1d<r32uint,write>;
[[group(0), binding(0)]]

17
tests/snapshots.rs
View File

@@ -267,7 +267,8 @@ fn write_output_msl(
allow_point_size: true,
};
let (string, tr_info) = msl::write_string(module, info, options, &pipeline_options).unwrap();
let (string, tr_info) =
msl::write_string(module, info, options, &pipeline_options).expect("Metal write failed");
for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_names) {
if let Err(error) = result {
@@ -308,9 +309,9 @@ fn write_output_glsl(
};
let mut buffer = String::new();
let mut writer =
glsl::Writer::new(&mut buffer, module, info, options, &pipeline_options).unwrap();
writer.write().unwrap();
let mut writer = glsl::Writer::new(&mut buffer, module, info, options, &pipeline_options)
.expect("GLSL init failed");
writer.write().expect("GLSL write failed");
fs::write(
destination.join(format!("glsl/{}.{}.{:?}.glsl", file_name, ep_name, stage)),
@@ -344,7 +345,7 @@ fn write_output_hlsl(
let mut buffer = String::new();
let mut writer = hlsl::Writer::new(&mut buffer, options);
let reflection_info = writer.write(module, info).unwrap();
let reflection_info = writer.write(module, info).expect("HLSL write failed");
fs::write(destination.join(format!("hlsl/{}.hlsl", file_name)), buffer).unwrap();
@@ -417,7 +418,7 @@ fn write_output_wgsl(
) {
use naga::back::wgsl;
let string = wgsl::write_string(module, info).unwrap();
let string = wgsl::write_string(module, info).expect("WGSL write failed");
fs::write(destination.join(format!("wgsl/{}.wgsl", file_name)), string).unwrap();
}
@@ -572,8 +573,8 @@ fn convert_spv_inverse_hyperbolic_trig_functions() {
}
#[cfg(all(feature = "spv-in", feature = "spv-out"))]
#[test]
fn convert_spv_pointer_access() {
//#[test] //TODO: https://github.com/gfx-rs/naga/issues/1432
fn _convert_spv_pointer_access() {
convert_spv("pointer-access", true, Targets::SPIRV);
}

20
tests/wgsl-errors.rs
View File

@@ -677,14 +677,30 @@ fn invalid_functions() {
if function_name == "unacceptable_unsized" && argument_name == "arg"
}
// A *valid* way to pass an unsized value.
check_validation_error! {
"
struct Unsized { data: array<f32>; };
fn acceptable_ptr_to_unsized(okay: ptr<storage, Unsized>) { }
fn acceptable_pointer_to_unsized(arg: ptr<workgroup, Unsized>) { }
":
Ok(_)
}
check_validation_error! {
"
struct Unsized { data: array<f32>; };
fn unacceptable_uniform_class(arg: ptr<uniform, f32>) { }
":
Err(naga::valid::ValidationError::Function {
name: function_name,
error: naga::valid::FunctionError::InvalidArgumentPointerClass {
index: 0,
name: argument_name,
class: naga::StorageClass::Uniform,
},
..
})
if function_name == "unacceptable_uniform_class" && argument_name == "arg"
}
}
#[test]