2022-02-03 18:00:04 +00:00
|
|
|
// language: metal2.0
|
2021-05-29 01:49:50 +00:00
|
|
|
#include <metal_stdlib>
|
|
|
|
#include <simd/simd.h>
|
|
|
|
|
2022-02-21 21:22:54 +00:00
|
|
|
using metal::uint;
|
|
|
|
|
2022-01-06 01:15:09 +00:00
|
|
|
struct _mslBufferSizes {
|
2022-02-21 21:22:54 +00:00
|
|
|
uint size3;
|
2022-01-06 01:15:09 +00:00
|
|
|
};
|
|
|
|
|
2021-12-27 02:57:47 +00:00
|
|
|
constexpr constant bool Foo_2 = true;
|
2021-11-09 14:24:41 +00:00
|
|
|
struct type_2 {
|
2021-05-29 02:11:39 +00:00
|
|
|
float inner[10u];
|
2021-05-29 01:49:50 +00:00
|
|
|
};
|
2021-12-27 02:57:47 +00:00
|
|
|
struct Foo {
|
|
|
|
metal::packed_float3 v3_;
|
|
|
|
float v1_;
|
|
|
|
};
|
2022-01-06 01:15:09 +00:00
|
|
|
typedef metal::float2 type_6[1];
|
2022-01-21 17:30:18 +00:00
|
|
|
struct type_8 {
|
|
|
|
metal::float4 inner[20];
|
2022-01-06 01:15:09 +00:00
|
|
|
};
|
2021-05-29 01:49:50 +00:00
|
|
|
|
2021-11-09 14:24:41 +00:00
|
|
|
kernel void main_(
|
|
|
|
threadgroup type_2& wg
|
2021-11-24 22:36:57 +00:00
|
|
|
, threadgroup metal::atomic_uint& at_1
|
2022-02-03 18:00:04 +00:00
|
|
|
, device Foo& alignment [[user(fake0)]]
|
2022-02-04 21:31:16 +00:00
|
|
|
, device type_6& dummy [[user(fake0)]]
|
|
|
|
, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
|
2021-05-29 01:49:50 +00:00
|
|
|
) {
|
2021-12-27 02:57:47 +00:00
|
|
|
float Foo_1 = 1.0;
|
2021-11-24 22:36:57 +00:00
|
|
|
bool at = true;
|
2022-01-21 17:30:18 +00:00
|
|
|
float _e9 = alignment.v1_;
|
|
|
|
wg.inner[3] = _e9;
|
|
|
|
float _e14 = metal::float3(alignment.v3_).x;
|
|
|
|
wg.inner[2] = _e14;
|
2022-02-04 21:31:16 +00:00
|
|
|
wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
|
2021-11-24 22:36:57 +00:00
|
|
|
metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
|
2021-05-29 01:49:50 +00:00
|
|
|
return;
|
|
|
|
}
|