mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-07-08 14:15:58 +00:00
This change implements pointers and references as described by the WGSL specification change in https://github.com/gpuweb/gpuweb/pull/1569. reader/spirv: * Now emits address-of `&expr` and indirection `*expr` operators as needed. * As an identifier may now resolve to a pointer or reference type depending on whether the declaration is a `var`, `let` or parameter, `Function::identifier_values_` has been changed from an ID set to an ID -> Type* map. resolver: * Now correctly resolves all expressions to either a value type, reference type or pointer type. * Validates pointer / reference rules on assignment, `var` and `let` construction, and usage. * Handles the address-of and indirection operators. * No longer does any implicit loads of pointer types. * Storage class validation is still TODO (crbug.com/tint/809) writer/spirv: * Correctly handles variables and expressions of pointer and reference types, emitting OpLoads where necessary. test: * Lots of new test cases Fixed: tint:727 Change-Id: I77d3281590e35e5a3122f5b74cdeb71a6fe51f74 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/50740 Commit-Queue: Ben Clayton <bclayton@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: David Neto <dneto@google.com>
121 lines
3.5 KiB
Plaintext
121 lines
3.5 KiB
Plaintext
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
struct tint_symbol_1 {
|
|
float2 a_particlePos [[attribute(0)]];
|
|
float2 a_particleVel [[attribute(1)]];
|
|
float2 a_pos [[attribute(2)]];
|
|
};
|
|
struct tint_symbol_2 {
|
|
float4 value [[position]];
|
|
};
|
|
struct tint_symbol_3 {
|
|
float4 value [[color(0)]];
|
|
};
|
|
struct Particle {
|
|
/* 0x0000 */ packed_float2 pos;
|
|
/* 0x0008 */ packed_float2 vel;
|
|
};
|
|
struct SimParams {
|
|
/* 0x0000 */ float deltaT;
|
|
/* 0x0004 */ float rule1Distance;
|
|
/* 0x0008 */ float rule2Distance;
|
|
/* 0x000c */ float rule3Distance;
|
|
/* 0x0010 */ float rule1Scale;
|
|
/* 0x0014 */ float rule2Scale;
|
|
/* 0x0018 */ float rule3Scale;
|
|
};
|
|
struct Particles {
|
|
/* 0x0000 */ Particle particles[5];
|
|
};
|
|
struct tint_symbol_5 {
|
|
uint3 gl_GlobalInvocationID [[thread_position_in_grid]];
|
|
};
|
|
|
|
vertex tint_symbol_2 vert_main(tint_symbol_1 tint_symbol [[stage_in]]) {
|
|
float2 const a_particlePos = tint_symbol.a_particlePos;
|
|
float2 const a_particleVel = tint_symbol.a_particleVel;
|
|
float2 const a_pos = tint_symbol.a_pos;
|
|
float angle = -( atan2(a_particleVel.x, a_particleVel.y));
|
|
float2 pos = float2(((a_pos.x * cos(angle)) - (a_pos.y * sin(angle))), ((a_pos.x * sin(angle)) + (a_pos.y * cos(angle))));
|
|
return {float4((pos + a_particlePos), 0.0f, 1.0f)};
|
|
}
|
|
|
|
fragment tint_symbol_3 frag_main() {
|
|
return {float4(1.0f, 1.0f, 1.0f, 1.0f)};
|
|
}
|
|
|
|
kernel void comp_main(tint_symbol_5 tint_symbol_4 [[stage_in]], constant SimParams& params [[buffer(0)]], device Particles& particlesA [[buffer(1)]], device Particles& particlesB [[buffer(2)]]) {
|
|
uint3 const gl_GlobalInvocationID = tint_symbol_4.gl_GlobalInvocationID;
|
|
uint index = gl_GlobalInvocationID.x;
|
|
if ((index >= 5u)) {
|
|
return;
|
|
}
|
|
float2 vPos = particlesA.particles[index].pos;
|
|
float2 vVel = particlesA.particles[index].vel;
|
|
float2 cMass = float2(0.0f, 0.0f);
|
|
float2 cVel = float2(0.0f, 0.0f);
|
|
float2 colVel = float2(0.0f, 0.0f);
|
|
int cMassCount = 0;
|
|
int cVelCount = 0;
|
|
float2 pos = 0.0f;
|
|
float2 vel = 0.0f;
|
|
{
|
|
uint i = 0u;
|
|
{
|
|
bool tint_msl_is_first_1 = true;
|
|
for(;;) {
|
|
if (!tint_msl_is_first_1) {
|
|
i = (i + 1u);
|
|
}
|
|
tint_msl_is_first_1 = false;
|
|
|
|
if (!((i < 5u))) {
|
|
break;
|
|
}
|
|
if ((i == index)) {
|
|
continue;
|
|
}
|
|
pos = particlesA.particles[i].pos.xy;
|
|
vel = particlesA.particles[i].vel.xy;
|
|
if (( distance(pos, vPos) < params.rule1Distance)) {
|
|
cMass = (cMass + pos);
|
|
cMassCount = (cMassCount + 1);
|
|
}
|
|
if (( distance(pos, vPos) < params.rule2Distance)) {
|
|
colVel = (colVel - (pos - vPos));
|
|
}
|
|
if (( distance(pos, vPos) < params.rule3Distance)) {
|
|
cVel = (cVel + vel);
|
|
cVelCount = (cVelCount + 1);
|
|
}
|
|
}
|
|
}
|
|
}
|
|
if ((cMassCount > 0)) {
|
|
cMass = ((cMass / float2(float(cMassCount), float(cMassCount))) - vPos);
|
|
}
|
|
if ((cVelCount > 0)) {
|
|
cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
|
|
}
|
|
vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
|
|
vVel = ( normalize(vVel) * clamp( length(vVel), 0.0f, 0.100000001f));
|
|
vPos = (vPos + (vVel * params.deltaT));
|
|
if ((vPos.x < -1.0f)) {
|
|
vPos.x = 1.0f;
|
|
}
|
|
if ((vPos.x > 1.0f)) {
|
|
vPos.x = -1.0f;
|
|
}
|
|
if ((vPos.y < -1.0f)) {
|
|
vPos.y = 1.0f;
|
|
}
|
|
if ((vPos.y > 1.0f)) {
|
|
vPos.y = -1.0f;
|
|
}
|
|
particlesB.particles[index].pos = vPos;
|
|
particlesB.particles[index].vel = vVel;
|
|
return;
|
|
}
|
|
|