2021-05-26 15:41:02 +00:00
|
|
|
#include <metal_stdlib>
|
|
|
|
|
|
|
|
using namespace metal;
|
2021-08-04 22:15:28 +00:00
|
|
|
void uses_a(threadgroup int* const tint_symbol) {
|
|
|
|
*(tint_symbol) = as_type<int>((as_type<uint>(*(tint_symbol)) + as_type<uint>(1)));
|
2021-05-26 15:41:02 +00:00
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void uses_b(threadgroup int* const tint_symbol_1) {
|
|
|
|
*(tint_symbol_1) = as_type<int>((as_type<uint>(*(tint_symbol_1)) * as_type<uint>(2)));
|
2021-05-26 15:41:02 +00:00
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void uses_a_and_b(threadgroup int* const tint_symbol_2, threadgroup int* const tint_symbol_3) {
|
|
|
|
*(tint_symbol_2) = *(tint_symbol_3);
|
2021-05-26 15:41:02 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
void no_uses() {
|
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void outer(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
|
|
|
|
*(tint_symbol_4) = 0;
|
|
|
|
uses_a(tint_symbol_4);
|
|
|
|
uses_a_and_b(tint_symbol_5, tint_symbol_4);
|
|
|
|
uses_b(tint_symbol_5);
|
2021-05-26 15:41:02 +00:00
|
|
|
no_uses();
|
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void main1_inner(uint local_invocation_index, threadgroup int* const tint_symbol_6) {
|
2021-07-30 14:08:06 +00:00
|
|
|
{
|
2022-06-01 10:08:29 +00:00
|
|
|
*(tint_symbol_6) = 0;
|
2021-06-18 22:44:31 +00:00
|
|
|
}
|
2021-06-28 15:30:57 +00:00
|
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
2021-08-04 22:15:28 +00:00
|
|
|
*(tint_symbol_6) = 42;
|
|
|
|
uses_a(tint_symbol_6);
|
|
|
|
}
|
|
|
|
|
|
|
|
kernel void main1(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
|
|
|
threadgroup int tint_symbol_7;
|
|
|
|
main1_inner(local_invocation_index, &(tint_symbol_7));
|
2021-05-26 15:41:02 +00:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void main2_inner(uint local_invocation_index_1, threadgroup int* const tint_symbol_8) {
|
2021-07-30 14:08:06 +00:00
|
|
|
{
|
2022-06-01 10:08:29 +00:00
|
|
|
*(tint_symbol_8) = 0;
|
2021-06-18 22:44:31 +00:00
|
|
|
}
|
2021-06-28 15:30:57 +00:00
|
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
2021-08-04 22:15:28 +00:00
|
|
|
*(tint_symbol_8) = 7;
|
|
|
|
uses_b(tint_symbol_8);
|
|
|
|
}
|
|
|
|
|
|
|
|
kernel void main2(uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
|
|
|
|
threadgroup int tint_symbol_9;
|
|
|
|
main2_inner(local_invocation_index_1, &(tint_symbol_9));
|
2021-05-26 15:41:02 +00:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2021-08-04 22:15:28 +00:00
|
|
|
void main3_inner(uint local_invocation_index_2, threadgroup int* const tint_symbol_10, threadgroup int* const tint_symbol_11) {
|
2021-07-30 14:08:06 +00:00
|
|
|
{
|
2022-06-01 10:08:29 +00:00
|
|
|
*(tint_symbol_10) = 0;
|
|
|
|
*(tint_symbol_11) = 0;
|
2021-06-18 22:44:31 +00:00
|
|
|
}
|
2021-06-28 15:30:57 +00:00
|
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
2021-08-04 22:15:28 +00:00
|
|
|
outer(tint_symbol_10, tint_symbol_11);
|
2021-05-26 15:41:02 +00:00
|
|
|
no_uses();
|
2021-08-04 22:15:28 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
kernel void main3(uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
|
|
|
|
threadgroup int tint_symbol_12;
|
|
|
|
threadgroup int tint_symbol_13;
|
|
|
|
main3_inner(local_invocation_index_2, &(tint_symbol_12), &(tint_symbol_13));
|
2021-05-26 15:41:02 +00:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
kernel void main4() {
|
|
|
|
no_uses();
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|