James Price 128980f218 tint: Add support for workgroupUniformLoad
Accept any type in the intrinsics definition, and then manually
validate that there are no atomics in the type. Add manual E2E tests
for composite types.

Use the BuiltinPolyfill transform to implement it for all backends.

Update the uniformity analysis with special-case tags for the builtin.

Fixed: tint:1780
Change-Id: I95786dff4df70a0b16ed1c53b853b5d0ec6bc501
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/114862
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>
2023-01-06 02:25:06 +00:00

29 lines
806 B
Plaintext

#include <metal_stdlib>
using namespace metal;
int tint_workgroupUniformLoad(threadgroup int* const p) {
threadgroup_barrier(mem_flags::mem_threadgroup);
int const result = *(p);
threadgroup_barrier(mem_flags::mem_threadgroup);
return result;
}
void workgroupUniformLoad_9d33de(threadgroup int* const tint_symbol) {
int res = tint_workgroupUniformLoad(tint_symbol);
}
void compute_main_inner(uint local_invocation_index, threadgroup int* const tint_symbol_1) {
{
*(tint_symbol_1) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
workgroupUniformLoad_9d33de(tint_symbol_1);
}
kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup int tint_symbol_2;
compute_main_inner(local_invocation_index, &(tint_symbol_2));
return;
}