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

Add block boolean to the structure types in IR

Browse Source 
Also update WGSL syntax for structures.
Move the boids code into snapshots.
This commit is contained in:
Dzmitry Malyshau
2020-12-11 00:50:36 -05:00
committed by Dzmitry Malyshau
parent d18b73c3a7
commit dafca34877
21 changed files with 264 additions and 141 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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<Type>, block: bool) -> BackendResult {
fn write_type(&mut self, ty: Handle<Type>) -> 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 {

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

@@ -694,7 +694,10 @@ impl<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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 =

28
src/back/spv/writer.rs
View File

@@ -142,6 +142,7 @@ pub struct Writer {
lookup_function_type: crate::FastHashMap<LookupFunctionType, Word>,
lookup_constant: crate::FastHashMap<crate::Handle<crate::Constant>, Word>,
lookup_global_variable: crate::FastHashMap<crate::Handle<crate::GlobalVariable>, Word>,
storage_type_handles: crate::FastHashSet<crate::Handle<crate::Type>>,
}
// 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);

17
src/front/glsl/parser.rs
View File

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

5
src/front/glsl/variables.rs
View File

@@ -142,7 +142,10 @@ impl Program {
meta: TokenMetadata,
) -> Result<Handle<Expression>, 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()))

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

@@ -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<I: Iterator<Item = u32>> Parser<I> {
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<I: Iterator<Item = u32>> Parser<I> {
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() {

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

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

5
src/lib.rs
View File

@@ -386,7 +386,10 @@ pub enum TypeInner {
stride: Option<NonZeroU32>,
},
/// User-defined structure.
Struct { members: Vec<StructMember> },
Struct {
block: bool,
members: Vec<StructMember>,
},
/// Possibly multidimensional array of texels.
Image {
dim: ImageDimension,

6
src/proc/namer.rs
View File

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

5
src/proc/typifier.rs
View File

@@ -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)?;

12
src/proc/validator.rs
View File

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

4
tests/snapshots.rs
View File

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

0
test-data/boids.param.ron → tests/snapshots/in/boids.param.ron
View File

17
test-data/boids.wgsl → tests/snapshots/in/boids.wgsl
View File

@@ -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<in> a_particlePos : vec2<f32>;
@@ -22,7 +20,7 @@ import "GLSL.std.450" as std;
[[builtin(position)]] var gl_Position : vec4<f32>;
[[stage(vertex)]]
fn main() -> void {
fn main() {
const angle : f32 = -atan2(a_particleVel.x, a_particleVel.y);
const pos : vec2<f32> = vec2<f32>(
(a_pos.x * cos(angle)) - (a_pos.y * sin(angle)),
@@ -34,17 +32,19 @@ fn main() -> void {
[[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]]
fn main() -> void {
fn main() {
fragColor = vec4<f32>(1.0, 1.0, 1.0, 1.0);
}
# compute shader
type Particle = struct {
[[block]]
struct Particle {
[[offset(0)]] pos : vec2<f32>;
[[offset(8)]] vel : vec2<f32>;
};
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<Particle, 5>;
};
@@ -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;

5
tests/snapshots/in/function.param.ron
View File

@@ -1,5 +0,0 @@
(
spv_capabilities: [ Shader ],
mtl_bindings: {
}
)

8
tests/snapshots/in/function.wgsl
View File

@@ -1,8 +0,0 @@
fn test_function(test: f32) -> f32 {
return test;
}
[[stage(vertex)]]
fn main() -> void {
var test: f32 = test_function(1.0);
}

6
tests/snapshots/in/quad.wgsl
View File

@@ -6,10 +6,9 @@ const c_scale: f32 = 1.2;
[[builtin(position)]] var<out> o_position : vec4<f32>;
[[stage(vertex)]]
fn main() -> void {
fn main() {
v_uv = a_uv;
o_position = vec4<f32>(c_scale * a_pos, 0.0, 1.0);
return;
}
# fragment
@@ -19,7 +18,6 @@ fn main() -> void {
[[location(0)]] var<out> o_color : vec4<f32>;
[[stage(fragment)]]
fn main() -> void {
fn main() {
o_color = textureSample(u_texture, u_sampler, v_uv);
return;
}

3
tests/snapshots/in/simple.wgsl
View File

@@ -2,7 +2,6 @@
[[builtin(position)]] var<out> o_position : vec4<f32>;
[[stage(vertex)]]
fn main() -> void {
fn main() {
o_position = vec4<f32>(1);
return;
}

127
tests/snapshots/snapshots__boids.msl.snap Normal file
View File

@@ -0,0 +1,127 @@
---
source: tests/snapshots.rs
expression: msl
---
#include <metal_stdlib>
#include <simd/simd.h>
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<uint>(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<uint>(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;
}

27
tests/snapshots/snapshots__function.msl.snap
View File

@@ -1,27 +0,0 @@
---
source: tests/snapshots.rs
expression: msl
---
#include <metal_stdlib>
#include <simd/simd.h>
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;
}

31
tests/snapshots/snapshots__function.spvasm.snap
View File

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