461f1506e7
This is not necessary, as we must emit an invalidating store before we
potentially consume an invalid expression. In fact, we're a bit
conservative here in this case for example:
int tmp = variable;
if (...)
{
variable = 10;
}
else
{
// Consuming tmp here is fine, but it was
// invalidated while emitting other branch.
// Technically, we need to study if there is an invalidating store
// in the CFG between the loading block and this block, and the other
// branch will not be a part of that analysis.
int tmp2 = tmp * tmp;
}
Fixing this case means complex CFG traversal *everywhere*, and it feels like overkill.
Fixing this exposed a bug with access chains, so fix a bug where expression dependencies were not
inherited properly in access chains. Access chains are now considered forwarded if there
is at least one dependency which is also forwarded.
61 lines
1.3 KiB
Plaintext
61 lines
1.3 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
struct foo
|
|
{
|
|
int a[128];
|
|
uint b;
|
|
float2 c;
|
|
};
|
|
|
|
struct bar
|
|
{
|
|
int d;
|
|
};
|
|
|
|
device foo* select_buffer(device foo& a, constant bar& cb)
|
|
{
|
|
return (cb.d != 0) ? &a : nullptr;
|
|
}
|
|
|
|
thread uint3* select_input(thread uint3& gl_GlobalInvocationID, thread uint3& gl_LocalInvocationID, constant bar& cb)
|
|
{
|
|
return (cb.d != 0) ? &gl_GlobalInvocationID : &gl_LocalInvocationID;
|
|
}
|
|
|
|
kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
|
{
|
|
device foo* _46 = select_buffer(buf, cb);
|
|
device foo* _45 = _46;
|
|
device foo* _48 = _45;
|
|
device int* _52;
|
|
device int* _55;
|
|
_52 = &_48->a[0u];
|
|
_55 = &buf.a[0u];
|
|
int _57;
|
|
int _58;
|
|
for (;;)
|
|
{
|
|
_57 = *_52;
|
|
_58 = *_55;
|
|
if (_57 != _58)
|
|
{
|
|
int _66 = (_57 + _58) + int((*select_input(gl_GlobalInvocationID, gl_LocalInvocationID, cb)).x);
|
|
*_52 = _66;
|
|
*_55 = _66;
|
|
_52 = &_52[1u];
|
|
_55 = &_55[1u];
|
|
continue;
|
|
}
|
|
else
|
|
{
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
|