18bcc9b790
This subtle bug removed any expression validation for trivially swizzled variables. Make usage suppression a more explicit concept rather than just hacking off forwarded_temporaries. There is some fallout here with loop generation since our expression invalidation is currently a bit too naive to handle loops properly. The forwarding bug masked this problem until now. If part of the loop condition is also used in the body, we end up reading an invalid expression, which in turn forces a temporary to be generated in the condition block, not good. We'll need to be smarter here ...
60 lines
1.3 KiB
Plaintext
60 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 int* _52;
|
|
device int* _55;
|
|
_52 = &_45->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;
|
|
}
|
|
}
|
|
}
|
|
|