mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-06-05 06:03:34 +00:00
89a0bde59c
Spread the array zeroing across as many workgroup invocations as possible. Bug: tint:910 Change-Id: I1cb5a6aaafd2a0a4093ea3b9797c173378bc5605 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60203 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: David Neto <dneto@google.com>
29 lines
699 B
Plaintext
29 lines
699 B
Plaintext
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
struct UBO {
|
|
/* 0x0000 */ int dynamic_idx;
|
|
};
|
|
struct tint_array_wrapper {
|
|
int arr[64];
|
|
};
|
|
struct S {
|
|
tint_array_wrapper data;
|
|
};
|
|
struct Result {
|
|
/* 0x0000 */ int out;
|
|
};
|
|
|
|
kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
|
|
threadgroup S tint_symbol_1;
|
|
for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
|
|
uint const i = idx;
|
|
tint_symbol_1.data.arr[i] = int();
|
|
}
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
tint_symbol_1.data.arr[ubo.dynamic_idx] = 1;
|
|
result.out = tint_symbol_1.data.arr[3];
|
|
return;
|
|
}
|
|
|