mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-06-02 20:51:45 +00:00
e548db90f6
This removes a lot of awkward logic from the MSL writer, and means that we now handle all module-scope variables with the same transform. Change-Id: I782e36a4b88dafbc3f8364f7caa7f95c6ae3f5f1 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67643 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
33 lines
999 B
Plaintext
33 lines
999 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;
|
|
};
|
|
|
|
void f_inner(uint local_invocation_index, threadgroup S* const tint_symbol, const constant UBO* const tint_symbol_1, device Result* const tint_symbol_2) {
|
|
for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
|
|
uint const i = idx;
|
|
(*(tint_symbol)).data.arr[i] = int();
|
|
}
|
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
(*(tint_symbol)).data.arr[(*(tint_symbol_1)).dynamic_idx] = 1;
|
|
(*(tint_symbol_2)).out = (*(tint_symbol)).data.arr[3];
|
|
}
|
|
|
|
kernel void f(const constant UBO* tint_symbol_4 [[buffer(0)]], device Result* tint_symbol_5 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
|
threadgroup S tint_symbol_3;
|
|
f_inner(local_invocation_index, &(tint_symbol_3), tint_symbol_4, tint_symbol_5);
|
|
return;
|
|
}
|
|
|