mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-05-16 20:31:20 +00:00
128980f218
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>
28 lines
1.0 KiB
Plaintext
28 lines
1.0 KiB
Plaintext
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
|
|
template<typename T, size_t N>
|
|
struct tint_array {
|
|
const constant T& operator[](size_t i) const constant { return elements[i]; }
|
|
device T& operator[](size_t i) device { return elements[i]; }
|
|
const device T& operator[](size_t i) const device { return elements[i]; }
|
|
thread T& operator[](size_t i) thread { return elements[i]; }
|
|
const thread T& operator[](size_t i) const thread { return elements[i]; }
|
|
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
|
|
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
|
|
T elements[N];
|
|
};
|
|
|
|
tint_array<int, 4> tint_workgroupUniformLoad(threadgroup tint_array<int, 4>* const p) {
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
tint_array<int, 4> const result = *(p);
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
return result;
|
|
}
|
|
|
|
tint_array<int, 4> foo(threadgroup tint_array<int, 4>* const tint_symbol) {
|
|
return tint_workgroupUniformLoad(tint_symbol);
|
|
}
|
|
|