diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 419452e1cd..7c302b5670 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -385,7 +385,11 @@ impl<'a, W: Write> Writer<'a, W> { // This are always ordered because of the IR is structured in a way that you can't make a // struct without adding all of it's members first for (handle, ty) in self.module.types.iter() { - if let TypeInner::Struct { ref members } = ty.inner { + if let TypeInner::Struct { + block: _, + ref members, + } = ty.inner + { self.write_struct(handle, members)? } } @@ -513,7 +517,7 @@ impl<'a, W: Write> Writer<'a, W> { /// - If type is either a image or sampler /// - If it's an Array with a [`ArraySize::Constant`](crate::ArraySize::Constant) with a /// constant that isn't [`Uint`](crate::ConstantInner::Uint) - fn write_type(&mut self, ty: Handle, block: bool) -> BackendResult { + fn write_type(&mut self, ty: Handle) -> BackendResult { match self.module.types[ty].inner { // Scalars are simple we just get the full name from `glsl_scalar` TypeInner::Scalar { kind, width } => { @@ -543,10 +547,10 @@ impl<'a, W: Write> Writer<'a, W> { rows as u8 )?, // glsl has no pointer types so just write types as normal and loads are skipped - TypeInner::Pointer { base, .. } => self.write_type(base, false)?, + TypeInner::Pointer { base, .. } => self.write_type(base)?, // Arrays are written as `base[size]` TypeInner::Array { base, size, .. } => { - self.write_type(base, false)?; + self.write_type(base)?; write!(self.out, "[")?; @@ -570,7 +574,7 @@ impl<'a, W: Write> Writer<'a, W> { // If it's a block we need to write `block_name { members }` where `block_name` must be // unique between blocks and structs so we add `_block_ID` where `ID` is a `IdGenerator` // generated number so it's unique and `members` are the same as in a struct - TypeInner::Struct { ref members } => { + TypeInner::Struct { block, ref members } => { // Get the struct name let name = &self.names[&NameKey::Type(ty)]; @@ -583,7 +587,7 @@ impl<'a, W: Write> Writer<'a, W> { // Add a tab for identation (readability only) writeln!(self.out, "\t")?; // Write the member type - self.write_type(member.ty, false)?; + self.write_type(member.ty)?; // Finish the member with the name, a semicolon and a newline // The leading space is important @@ -694,20 +698,13 @@ impl<'a, W: Write> Writer<'a, W> { }; } - // glsl doesn't allow structures as types in `buffer` and `uniform` instead blocks must be - // used so we set block to true in `write_type` - let block = match global.class { - StorageClass::Storage | StorageClass::Uniform => true, - _ => false, - }; - // Write the storage class // Trailing space is important write!(self.out, "{} ", glsl_storage_class(global.class))?; // Write the type // `write_type` adds no leading or trailing spaces - self.write_type(global.ty, block)?; + self.write_type(global.ty)?; // Finally write the global name and end the global with a `;` and a newline // Leading space is important @@ -762,7 +759,7 @@ impl<'a, W: Write> Writer<'a, W> { // This is the only place where `void` is a valid type // (though it's more a keyword than a type) if let Some(ty) = func.return_type { - self.write_type(ty, false)?; + self.write_type(ty)?; } else { write!(self.out, "void")?; } @@ -777,7 +774,7 @@ impl<'a, W: Write> Writer<'a, W> { self.write_slice(&func.arguments, |this, i, arg| { // Write the argument type // `write_type` adds no trailing spaces - this.write_type(arg.ty, false)?; + this.write_type(arg.ty)?; // Write the argument name // The leading space is important @@ -797,7 +794,7 @@ impl<'a, W: Write> Writer<'a, W> { // Write identation (only for readability) and the type // `write_type` adds no trailing space write!(self.out, "\t")?; - self.write_type(local.ty, false)?; + self.write_type(local.ty)?; // Write the local name // The leading space is important @@ -883,7 +880,7 @@ impl<'a, W: Write> Writer<'a, W> { // Composite constant are created using the same syntax as compose // `type(components)` where `components` is a comma separated list of constants ConstantInner::Composite(ref components) => { - self.write_type(constant.ty, false)?; + self.write_type(constant.ty)?; write!(self.out, "(")?; // Write the comma separated constants @@ -919,7 +916,7 @@ impl<'a, W: Write> Writer<'a, W> { // Write the member type // Adds no trailing space - self.write_type(member.ty, false)?; + self.write_type(member.ty)?; // Write the member name and put a semicolon // The leading space is important @@ -1139,7 +1136,7 @@ impl<'a, W: Write> Writer<'a, W> { // `Compose` is pretty simple we just write `type(components)` where `components` is a // comma separated list of expressions Expression::Compose { ty, ref components } => { - self.write_type(ty, false)?; + self.write_type(ty)?; write!(self.out, "(")?; self.write_slice(components, |this, _, arg| this.write_expr(*arg, ctx))?; @@ -1492,7 +1489,7 @@ impl<'a, W: Write> Writer<'a, W> { convert, } => { if convert { - self.write_type(ctx.typifier.get_handle(expr).unwrap(), false)?; + self.write_type(ctx.typifier.get_handle(expr).unwrap())?; } else { let source_kind = match *ctx.typifier.get(expr, &self.module.types) { TypeInner::Scalar { diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 710948d625..e151e9f69c 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -694,7 +694,10 @@ impl Writer { crate::ArraySize::Dynamic => write!(self.out, "1]")?, } } - crate::TypeInner::Struct { ref members } => { + crate::TypeInner::Struct { + block: _, + ref members, + } => { writeln!(self.out, "struct {} {{", name)?; for (index, member) in members.iter().enumerate() { let member_name = &self.names[&NameKey::StructMember(handle, index as u32)]; @@ -893,7 +896,10 @@ impl Writer { continue; } // if it's a struct, lift all the built-in contents up to the root - if let crate::TypeInner::Struct { ref members } = module.types[var.ty].inner + if let crate::TypeInner::Struct { + block: _, + ref members, + } = module.types[var.ty].inner { for (index, member) in members.iter().enumerate() { if let crate::MemberOrigin::BuiltIn(built_in) = member.origin { @@ -932,7 +938,10 @@ impl Writer { continue; } // if it's a struct, lift all the built-in contents up to the root - if let crate::TypeInner::Struct { ref members } = module.types[var.ty].inner + if let crate::TypeInner::Struct { + block: _, + ref members, + } = module.types[var.ty].inner { for (index, member) in members.iter().enumerate() { let name = diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index db06ba320a..7650dac671 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -142,6 +142,7 @@ pub struct Writer { lookup_function_type: crate::FastHashMap, lookup_constant: crate::FastHashMap, Word>, lookup_global_variable: crate::FastHashMap, Word>, + storage_type_handles: crate::FastHashSet>, } // type alias, for success return of write_expression @@ -167,6 +168,7 @@ impl Writer { lookup_function_type: crate::FastHashMap::default(), lookup_constant: crate::FastHashMap::default(), lookup_global_variable: crate::FastHashMap::default(), + storage_type_handles: crate::FastHashSet::default(), } } @@ -588,7 +590,20 @@ impl Writer { } } } - crate::TypeInner::Struct { ref members } => { + crate::TypeInner::Struct { block, ref members } => { + if block { + let decoration = if self.storage_type_handles.contains(&handle) { + spirv::Decoration::BufferBlock + } else { + spirv::Decoration::Block + }; + self.annotations + .push(super::instructions::instruction_decorate( + id, + decoration, + &[], + )); + } let mut member_ids = Vec::with_capacity(members.len()); for member in members { let member_id = self.get_type_id(arena, LookupType::Handle(member.ty))?; @@ -956,7 +971,10 @@ impl Writer { LookupType::Local(LocalType::Scalar { kind, width }), ) } - crate::TypeInner::Struct { ref members } => { + crate::TypeInner::Struct { + block: _, + ref members, + } => { let member = &members[index as usize]; let type_id = self.get_type_id(&ir_module.types, LookupType::Handle(member.ty))?; @@ -1625,6 +1643,12 @@ impl Writer { )); } + for (_, var) in ir_module.global_variables.iter() { + if !var.storage_access.is_empty() { + self.storage_type_handles.insert(var.ty); + } + } + for (handle, ir_function) in ir_module.functions.iter() { let id = self.write_function(ir_function, ir_module)?; self.lookup_function.insert(handle, id); diff --git a/src/front/glsl/parser.rs b/src/front/glsl/parser.rs index 2d03b3a44b..4372721912 100644 --- a/src/front/glsl/parser.rs +++ b/src/front/glsl/parser.rs @@ -541,14 +541,16 @@ pomelo! { if i.1 == "gl_PerVertex" { None } else { + let block = !t.is_empty(); Some(VarDeclaration { type_qualifiers: t, ids_initializers: vec![(None, None)], ty: extra.module.types.fetch_or_append(Type{ name: Some(i.1), inner: TypeInner::Struct { - members: sdl - } + block, + members: sdl, + }, }), }) } @@ -556,14 +558,16 @@ pomelo! { declaration ::= type_qualifier(t) Identifier(i1) LeftBrace struct_declaration_list(sdl) RightBrace Identifier(i2) Semicolon { + let block = !t.is_empty(); Some(VarDeclaration { type_qualifiers: t, ids_initializers: vec![(Some(i2.1), None)], ty: extra.module.types.fetch_or_append(Type{ name: Some(i1.1), inner: TypeInner::Struct { - members: sdl - } + block, + members: sdl, + }, }), }) } @@ -720,7 +724,8 @@ pomelo! { Type{ name: Some(i.1), inner: TypeInner::Struct { - members: vec![] + block: false, + members: vec![], } } } @@ -1013,7 +1018,7 @@ pomelo! { } else { let ty = &extra.module.types[var.ty]; // anonymous structs - if let TypeInner::Struct { members } = &ty.inner { + if let TypeInner::Struct { block: _, ref members } = ty.inner { let base = extra.context.expressions.append( Expression::GlobalVariable(var_handle) ); diff --git a/src/front/glsl/variables.rs b/src/front/glsl/variables.rs index 50343ba922..7fb1a3518e 100644 --- a/src/front/glsl/variables.rs +++ b/src/front/glsl/variables.rs @@ -142,7 +142,10 @@ impl Program { meta: TokenMetadata, ) -> Result, ErrorKind> { match *self.resolve_type(expression)? { - TypeInner::Struct { ref members } => { + TypeInner::Struct { + block: _, + ref members, + } => { let index = members .iter() .position(|m| m.name == Some(name.into())) diff --git a/src/front/spv/mod.rs b/src/front/spv/mod.rs index 782c62a6fa..0f16da0304 100644 --- a/src/front/spv/mod.rs +++ b/src/front/spv/mod.rs @@ -142,7 +142,7 @@ fn check_sample_coordinates( type MemberIndex = u32; -#[derive(Debug, Default)] +#[derive(Clone, Debug, Default, PartialEq)] struct Block { buffer: bool, } @@ -2014,12 +2014,6 @@ impl> Parser { inst.expect_at_least(2)?; let id = self.next()?; let parent_decor = self.future_decor.remove(&id); - let is_buffer_block = parent_decor - .as_ref() - .map_or(false, |decor| match decor.block { - Some(Block { buffer }) => buffer, - _ => false, - }); let mut members = Vec::with_capacity(inst.wc as usize - 2); let mut member_type_ids = Vec::with_capacity(members.capacity()); @@ -2038,13 +2032,18 @@ impl> Parser { ty, }); } - let inner = crate::TypeInner::Struct { members }; + + let block_decor = parent_decor.as_ref().and_then(|decor| decor.block.clone()); + let inner = crate::TypeInner::Struct { + block: block_decor.is_some(), + members, + }; let ty_handle = module.types.append(crate::Type { name: parent_decor.and_then(|dec| dec.name), inner, }); - if is_buffer_block { + if block_decor == Some(Block { buffer: true }) { self.lookup_storage_buffer_types.insert(ty_handle); } for (i, type_id) in member_type_ids.into_iter().enumerate() { diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 365955d2f1..edc949c004 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -315,7 +315,7 @@ impl Parser { Ok(type_arena.fetch_or_append(crate::Type { name: None, inner })) } crate::TypeInner::Array { base, .. } => Ok(base), - crate::TypeInner::Struct { ref members } => Ok(members[index].ty), + crate::TypeInner::Struct { ref members, .. } => Ok(members[index].ty), _ => Err(Error::NotCompositeType(ty)), } } @@ -659,7 +659,7 @@ impl Parser { let _ = lexer.next(); let name = lexer.next_ident()?; let expression = match *ctx.resolve_type(handle)? { - crate::TypeInner::Struct { ref members } => { + crate::TypeInner::Struct { ref members, .. } => { let index = members .iter() .position(|m| m.name.as_deref() == Some(name)) @@ -1179,10 +1179,6 @@ impl Parser { crate::TypeInner::Array { base, size, stride } } - "struct" => { - let members = self.parse_struct_body(lexer, type_arena, const_arena)?; - crate::TypeInner::Struct { members } - } "sampler" => crate::TypeInner::Sampler { comparison: false }, "sampler_comparison" => crate::TypeInner::Sampler { comparison: true }, "texture_sampled_1d" => { @@ -1698,6 +1694,7 @@ impl Parser { // Perspective is the default qualifier. let mut interpolation = None; let mut stage = None; + let mut is_block = false; let mut workgroup_size = [0u32; 3]; if lexer.skip(Token::DoubleParen('[')) { @@ -1722,6 +1719,9 @@ impl Parser { bind_index = Some(lexer.next_uint_literal()?); lexer.expect(Token::Paren(')'))?; } + "block" => { + is_block = true; + } "group" => { lexer.expect(Token::Paren('('))?; bind_group = Some(lexer.next_uint_literal()?); @@ -1771,9 +1771,24 @@ impl Parser { } self.scopes.pop(); } + // read items match lexer.next() { Token::Separator(';') => {} + Token::Word("struct") => { + let name = lexer.next_ident()?; + let members = + self.parse_struct_body(lexer, &mut module.types, &mut module.constants)?; + let ty = module.types.fetch_or_append(crate::Type { + name: Some(name.to_string()), + inner: crate::TypeInner::Struct { + block: is_block, + members, + }, + }); + self.lookup_type.insert(name.to_owned(), ty); + lexer.expect(Token::Separator(';'))?; + } Token::Word("type") => { let name = lexer.next_ident()?; lexer.expect(Token::Operation('='))?; @@ -1857,6 +1872,7 @@ impl Parser { Token::End => return Ok(false), other => return other.unexpected("global item"), } + match binding { None => Ok(true), // we had the decoration but no var? diff --git a/src/lib.rs b/src/lib.rs index dff0233ace..6cc3bd4db1 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -386,7 +386,10 @@ pub enum TypeInner { stride: Option, }, /// User-defined structure. - Struct { members: Vec }, + Struct { + block: bool, + members: Vec, + }, /// Possibly multidimensional array of texels. Image { dim: ImageDimension, diff --git a/src/proc/namer.rs b/src/proc/namer.rs index 03b508904b..4e6ddf9977 100644 --- a/src/proc/namer.rs +++ b/src/proc/namer.rs @@ -80,7 +80,11 @@ impl Namer { let ty_name = this.call_or(&ty.name, "type"); output.insert(NameKey::Type(ty_handle), ty_name); - if let crate::TypeInner::Struct { ref members } = ty.inner { + if let crate::TypeInner::Struct { + block: _, + ref members, + } = ty.inner + { for (index, member) in members.iter().enumerate() { let name = this.call_or(&member.name, "member"); output.insert(NameKey::StructMember(ty_handle, index as u32), name); diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 8fe370163b..07eb1274fe 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -153,7 +153,10 @@ impl Typifier { }) } crate::TypeInner::Array { base, .. } => Resolution::Handle(base), - crate::TypeInner::Struct { ref members } => { + crate::TypeInner::Struct { + block: _, + ref members, + } => { let member = members .get(index as usize) .ok_or(ResolveError::InvalidAccessIndex)?; diff --git a/src/proc/validator.rs b/src/proc/validator.rs index f39a5bf85f..fdb7fd0d84 100644 --- a/src/proc/validator.rs +++ b/src/proc/validator.rs @@ -186,7 +186,10 @@ impl crate::GlobalVariable { None => { match types[self.ty].inner { //TODO: check the member types - crate::TypeInner::Struct { members: _ } => self.forbid_interpolation()?, + crate::TypeInner::Struct { + block: _, + members: _, + } => self.forbid_interpolation()?, _ => return Err(GlobalVariableError::InvalidType), } } @@ -492,8 +495,11 @@ impl Validator { } } } - Ti::Struct { ref members } => { - //TODO: check that offsets are not intersecting? + Ti::Struct { + block: _, + ref members, + } => { + //TODO: check the offsets for member in members { if member.ty >= handle { return Err(ValidationError::UnresolvedType(member.ty)); diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 630e48cfa4..32072f6800 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -120,6 +120,6 @@ fn converts_wgsl_simple() { #[cfg(feature = "wgsl-in")] #[test] -fn converts_wgsl_function() { - convert_wgsl("function"); +fn converts_wgsl_boids() { + convert_wgsl("boids"); } diff --git a/test-data/boids.param.ron b/tests/snapshots/in/boids.param.ron similarity index 100% rename from test-data/boids.param.ron rename to tests/snapshots/in/boids.param.ron diff --git a/test-data/boids.wgsl b/tests/snapshots/in/boids.wgsl similarity index 95% rename from test-data/boids.wgsl rename to tests/snapshots/in/boids.wgsl index bdcf36bed7..8c5ff2fd15 100644 --- a/test-data/boids.wgsl +++ b/tests/snapshots/in/boids.wgsl @@ -12,8 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -import "GLSL.std.450" as std; - # vertex shader [[location(0)]] var a_particlePos : vec2; @@ -22,7 +20,7 @@ import "GLSL.std.450" as std; [[builtin(position)]] var gl_Position : vec4; [[stage(vertex)]] -fn main() -> void { +fn main() { const angle : f32 = -atan2(a_particleVel.x, a_particleVel.y); const pos : vec2 = vec2( (a_pos.x * cos(angle)) - (a_pos.y * sin(angle)), @@ -34,17 +32,19 @@ fn main() -> void { [[location(0)]] var fragColor : vec4; [[stage(fragment)]] -fn main() -> void { +fn main() { fragColor = vec4(1.0, 1.0, 1.0, 1.0); } # compute shader -type Particle = struct { +[[block]] +struct Particle { [[offset(0)]] pos : vec2; [[offset(8)]] vel : vec2; }; -type SimParams = struct { +[[block]] +struct SimParams { [[offset(0)]] deltaT : f32; [[offset(4)]] rule1Distance : f32; [[offset(8)]] rule2Distance : f32; @@ -54,7 +54,8 @@ type SimParams = struct { [[offset(24)]] rule3Scale : f32; }; -type Particles = struct { +[[block]] +struct Particles { [[offset(0)]] particles : [[stride 16]] array; }; @@ -66,7 +67,7 @@ type Particles = struct { # https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp [[stage(compute), workgroup_size(1)]] -fn main() -> void { +fn main() { const index : u32 = gl_GlobalInvocationID.x; if (index >= u32(5)) { return; diff --git a/tests/snapshots/in/function.param.ron b/tests/snapshots/in/function.param.ron deleted file mode 100644 index 4adad36220..0000000000 --- a/tests/snapshots/in/function.param.ron +++ /dev/null @@ -1,5 +0,0 @@ -( - spv_capabilities: [ Shader ], - mtl_bindings: { - } -) diff --git a/tests/snapshots/in/function.wgsl b/tests/snapshots/in/function.wgsl deleted file mode 100644 index 12549c7f97..0000000000 --- a/tests/snapshots/in/function.wgsl +++ /dev/null @@ -1,8 +0,0 @@ -fn test_function(test: f32) -> f32 { - return test; -} - -[[stage(vertex)]] -fn main() -> void { - var test: f32 = test_function(1.0); -} diff --git a/tests/snapshots/in/quad.wgsl b/tests/snapshots/in/quad.wgsl index 16cbd2698a..50b367cf9a 100644 --- a/tests/snapshots/in/quad.wgsl +++ b/tests/snapshots/in/quad.wgsl @@ -6,10 +6,9 @@ const c_scale: f32 = 1.2; [[builtin(position)]] var o_position : vec4; [[stage(vertex)]] -fn main() -> void { +fn main() { v_uv = a_uv; o_position = vec4(c_scale * a_pos, 0.0, 1.0); - return; } # fragment @@ -19,7 +18,6 @@ fn main() -> void { [[location(0)]] var o_color : vec4; [[stage(fragment)]] -fn main() -> void { +fn main() { o_color = textureSample(u_texture, u_sampler, v_uv); - return; } diff --git a/tests/snapshots/in/simple.wgsl b/tests/snapshots/in/simple.wgsl index bbe591d8f3..ec43564123 100644 --- a/tests/snapshots/in/simple.wgsl +++ b/tests/snapshots/in/simple.wgsl @@ -2,7 +2,6 @@ [[builtin(position)]] var o_position : vec4; [[stage(vertex)]] -fn main() -> void { +fn main() { o_position = vec4(1); - return; } diff --git a/tests/snapshots/snapshots__boids.msl.snap b/tests/snapshots/snapshots__boids.msl.snap new file mode 100644 index 0000000000..2e3efb40c2 --- /dev/null +++ b/tests/snapshots/snapshots__boids.msl.snap @@ -0,0 +1,127 @@ +--- +source: tests/snapshots.rs +expression: msl +--- +#include +#include + +typedef float2 type; +typedef float4 type1; +typedef float type2; +struct Particle { + type pos; + type vel; +}; +struct SimParams { + type2 deltaT; + type2 rule1Distance; + type2 rule2Distance; + type2 rule3Distance; + type2 rule1Scale; + type2 rule2Scale; + type2 rule3Scale; +}; +typedef uint type3; +typedef Particle type4[5]; +struct Particles { + type4 particles; +}; +typedef uint3 type5; +typedef int type6; + +struct main1Input { + type a_particlePos [[attribute(0)]]; + type a_particleVel [[attribute(1)]]; + type a_pos [[attribute(2)]]; +}; +struct main1Output { + type1 gl_Position [[position]]; +}; +vertex main1Output main1( + main1Input input [[stage_in]] +) { + main1Output output; + output.gl_Position = float4(float2(input.a_pos.x * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) - input.a_pos.y * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)), input.a_pos.x * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) + input.a_pos.y * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y))) + input.a_particlePos, 0.0, 1.0); + return output; +} +struct main2Input { +}; +struct main2Output { + type1 fragColor [[color(0)]]; +}; +fragment main2Output main2( + main2Input input [[stage_in]] +) { + main2Output output; + output.fragColor = float4(1.0, 1.0, 1.0, 1.0); + return output; +} +kernel void main3( + constant SimParams& params [[buffer(0)]], + constant Particles& particlesA [[buffer(1)]], + device Particles& particlesB [[buffer(2)]], + type5 gl_GlobalInvocationID [[thread_position_in_grid]] +) { + type vPos; + type vVel; + type cMass; + type cVel; + type colVel; + type6 cMassCount = 0; + type6 cVelCount = 0; + type pos1; + type vel1; + type3 i = 0; + if (gl_GlobalInvocationID.x >= static_cast(5)) { + } + vPos = particlesA.particles[gl_GlobalInvocationID.x].pos; + vVel = particlesA.particles[gl_GlobalInvocationID.x].vel; + cMass = float2(0.0, 0.0); + cVel = float2(0.0, 0.0); + colVel = float2(0.0, 0.0); + while(true) { + if (i >= static_cast(5)) { + break; + } + if (i == gl_GlobalInvocationID.x) { + continue; + } + pos1 = float2(particlesA.particles[i].pos.x, particlesA.particles[i].pos.y); + vel1 = float2(particlesA.particles[i].vel.x, particlesA.particles[i].vel.y); + if (metal::distance(pos1, vPos) < params.rule1Distance) { + cMass = cMass + pos1; + cMassCount = cMassCount + 1; + } + if (metal::distance(pos1, vPos) < params.rule2Distance) { + colVel = colVel - pos1 - vPos; + } + if (metal::distance(pos1, vPos) < params.rule3Distance) { + cVel = cVel + vel1; + cVelCount = cVelCount + 1; + } + } + if (cMassCount == 0) { + cMass = cMass / float2(cMassCount, cMassCount) + vPos; + } + if (cVelCount == 0) { + cVel = cVel / float2(cVelCount, cVelCount); + } + vVel = vVel + cMass * params.rule1Scale + colVel * params.rule2Scale + cVel * params.rule3Scale; + vVel = metal::normalize(vVel) * metal::clamp(metal::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; + } + particlesB.particles[gl_GlobalInvocationID.x].pos = vPos; + particlesB.particles[gl_GlobalInvocationID.x].vel = vVel; +} + diff --git a/tests/snapshots/snapshots__function.msl.snap b/tests/snapshots/snapshots__function.msl.snap deleted file mode 100644 index cbaab56341..0000000000 --- a/tests/snapshots/snapshots__function.msl.snap +++ /dev/null @@ -1,27 +0,0 @@ ---- -source: tests/snapshots.rs -expression: msl ---- -#include -#include - -typedef float type; - -type test_function( - type test -) { - return test; -} -struct main1Input { -}; -struct main1Output { -}; -vertex main1Output main1( - main1Input input [[stage_in]] -) { - main1Output output; - type test1; - test1 = test_function(1.0); - return output; -} - diff --git a/tests/snapshots/snapshots__function.spvasm.snap b/tests/snapshots/snapshots__function.spvasm.snap deleted file mode 100644 index a28f333e15..0000000000 --- a/tests/snapshots/snapshots__function.spvasm.snap +++ /dev/null @@ -1,31 +0,0 @@ ---- -source: tests/snapshots.rs -expression: dis ---- -; SPIR-V -; Version: 1.0 -; Generator: rspirv -; Bound: 16 -OpCapability Shader -%1 = OpExtInstImport "GLSL.std.450" -OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %11 "main" -%2 = OpTypeFloat 32 -%4 = OpTypePointer Function %2 -%6 = OpTypeFunction %2 %4 -%10 = OpTypeVoid -%12 = OpTypeFunction %10 -%15 = OpConstant %2 1.0 -%5 = OpFunction %2 None %6 -%3 = OpFunctionParameter %4 -%7 = OpLabel -%8 = OpLoad %2 %3 -OpReturnValue %8 -OpFunctionEnd -%11 = OpFunction %10 None %12 -%13 = OpLabel -%9 = OpVariable %4 Function -%14 = OpFunctionCall %2 %5 %15 -OpStore %9 %14 -OpReturn -OpFunctionEnd