mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-26 00:33:51 +00:00
3fda684eb9
See the comments in the code for details. This patch emits the definition of the macro only when the first loop is encountered. This does make that first loop's code look a bit odd: it would be more natural to define the macro at the top of the file. (See the modified files in `naga/tests/out/msl`.) Rejected alternatives: - We could emit the macro definition unconditionally at the top of the file. But this changes every MSL snapshot output file, whereas only eight of them actually contain loops. - We could have the validator flag modules that contain loops. But the changes end up being not small, and spread across the validator, so this seems disproportionate. If we had other consumers of this information, it might make sense. - We could change the MSL backend to allow text to be generated out of order, so that we can decide whether to define the macro after we've generated all the function bodies. But at the moment this seems like unnecessary complexity, although it might be worth doing in the future if we had additional uses for it - say, to conditionally emit helper function definitions. Fixes #4972.
160 lines
4.5 KiB
Plaintext
160 lines
4.5 KiB
Plaintext
// language: metal1.0
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using metal::uint;
|
|
|
|
struct _mslBufferSizes {
|
|
uint size1;
|
|
uint size2;
|
|
};
|
|
|
|
struct Particle {
|
|
metal::float2 pos;
|
|
metal::float2 vel;
|
|
};
|
|
struct SimParams {
|
|
float deltaT;
|
|
float rule1Distance;
|
|
float rule2Distance;
|
|
float rule3Distance;
|
|
float rule1Scale;
|
|
float rule2Scale;
|
|
float rule3Scale;
|
|
};
|
|
typedef Particle type_3[1];
|
|
struct Particles {
|
|
type_3 particles;
|
|
};
|
|
constant uint NUM_PARTICLES = 1500u;
|
|
|
|
struct main_Input {
|
|
};
|
|
kernel void main_(
|
|
metal::uint3 global_invocation_id [[thread_position_in_grid]]
|
|
, constant SimParams& params [[buffer(0)]]
|
|
, device Particles const& particlesSrc [[buffer(1)]]
|
|
, device Particles& particlesDst [[buffer(2)]]
|
|
, constant _mslBufferSizes& _buffer_sizes [[buffer(3)]]
|
|
) {
|
|
metal::float2 vPos = {};
|
|
metal::float2 vVel = {};
|
|
metal::float2 cMass = metal::float2(0.0, 0.0);
|
|
metal::float2 cVel = metal::float2(0.0, 0.0);
|
|
metal::float2 colVel = metal::float2(0.0, 0.0);
|
|
int cMassCount = 0;
|
|
int cVelCount = 0;
|
|
metal::float2 pos = {};
|
|
metal::float2 vel = {};
|
|
uint i = 0u;
|
|
uint index = global_invocation_id.x;
|
|
if (index >= NUM_PARTICLES) {
|
|
return;
|
|
}
|
|
metal::float2 _e8 = particlesSrc.particles[index].pos;
|
|
vPos = _e8;
|
|
metal::float2 _e14 = particlesSrc.particles[index].vel;
|
|
vVel = _e14;
|
|
#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
|
|
bool loop_init = true;
|
|
LOOP_IS_REACHABLE while(true) {
|
|
if (!loop_init) {
|
|
uint _e91 = i;
|
|
i = _e91 + 1u;
|
|
}
|
|
loop_init = false;
|
|
uint _e36 = i;
|
|
if (_e36 >= NUM_PARTICLES) {
|
|
break;
|
|
}
|
|
uint _e39 = i;
|
|
if (_e39 == index) {
|
|
continue;
|
|
}
|
|
uint _e43 = i;
|
|
metal::float2 _e46 = particlesSrc.particles[_e43].pos;
|
|
pos = _e46;
|
|
uint _e49 = i;
|
|
metal::float2 _e52 = particlesSrc.particles[_e49].vel;
|
|
vel = _e52;
|
|
metal::float2 _e53 = pos;
|
|
metal::float2 _e54 = vPos;
|
|
float _e58 = params.rule1Distance;
|
|
if (metal::distance(_e53, _e54) < _e58) {
|
|
metal::float2 _e60 = cMass;
|
|
metal::float2 _e61 = pos;
|
|
cMass = _e60 + _e61;
|
|
int _e63 = cMassCount;
|
|
cMassCount = _e63 + 1;
|
|
}
|
|
metal::float2 _e66 = pos;
|
|
metal::float2 _e67 = vPos;
|
|
float _e71 = params.rule2Distance;
|
|
if (metal::distance(_e66, _e67) < _e71) {
|
|
metal::float2 _e73 = colVel;
|
|
metal::float2 _e74 = pos;
|
|
metal::float2 _e75 = vPos;
|
|
colVel = _e73 - (_e74 - _e75);
|
|
}
|
|
metal::float2 _e78 = pos;
|
|
metal::float2 _e79 = vPos;
|
|
float _e83 = params.rule3Distance;
|
|
if (metal::distance(_e78, _e79) < _e83) {
|
|
metal::float2 _e85 = cVel;
|
|
metal::float2 _e86 = vel;
|
|
cVel = _e85 + _e86;
|
|
int _e88 = cVelCount;
|
|
cVelCount = _e88 + 1;
|
|
}
|
|
}
|
|
int _e94 = cMassCount;
|
|
if (_e94 > 0) {
|
|
metal::float2 _e97 = cMass;
|
|
int _e98 = cMassCount;
|
|
metal::float2 _e102 = vPos;
|
|
cMass = (_e97 / metal::float2(static_cast<float>(_e98))) - _e102;
|
|
}
|
|
int _e104 = cVelCount;
|
|
if (_e104 > 0) {
|
|
metal::float2 _e107 = cVel;
|
|
int _e108 = cVelCount;
|
|
cVel = _e107 / metal::float2(static_cast<float>(_e108));
|
|
}
|
|
metal::float2 _e112 = vVel;
|
|
metal::float2 _e113 = cMass;
|
|
float _e116 = params.rule1Scale;
|
|
metal::float2 _e119 = colVel;
|
|
float _e122 = params.rule2Scale;
|
|
metal::float2 _e125 = cVel;
|
|
float _e128 = params.rule3Scale;
|
|
vVel = ((_e112 + (_e113 * _e116)) + (_e119 * _e122)) + (_e125 * _e128);
|
|
metal::float2 _e131 = vVel;
|
|
metal::float2 _e133 = vVel;
|
|
vVel = metal::normalize(_e131) * metal::clamp(metal::length(_e133), 0.0, 0.1);
|
|
metal::float2 _e139 = vPos;
|
|
metal::float2 _e140 = vVel;
|
|
float _e143 = params.deltaT;
|
|
vPos = _e139 + (_e140 * _e143);
|
|
float _e147 = vPos.x;
|
|
if (_e147 < -1.0) {
|
|
vPos.x = 1.0;
|
|
}
|
|
float _e153 = vPos.x;
|
|
if (_e153 > 1.0) {
|
|
vPos.x = -1.0;
|
|
}
|
|
float _e159 = vPos.y;
|
|
if (_e159 < -1.0) {
|
|
vPos.y = 1.0;
|
|
}
|
|
float _e165 = vPos.y;
|
|
if (_e165 > 1.0) {
|
|
vPos.y = -1.0;
|
|
}
|
|
metal::float2 _e174 = vPos;
|
|
particlesDst.particles[index].pos = _e174;
|
|
metal::float2 _e179 = vVel;
|
|
particlesDst.particles[index].vel = _e179;
|
|
return;
|
|
}
|