mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-11 22:44:04 +00:00
tint/writer/msl: Generate an array<T,N> helper
And remove the WrapArraysInStructs transform. Wrapping arrays in structures becomes troublesome for `const` arrays, as currently WGSL does not allow `const` structures. MSL 2.0+ has a builtin array<> helper, but we're targetting MSL 1.2, so we have to emit our own. Fortunately, it can be done with a few lines of templated code. This produces significantly cleaner output. Change-Id: Ifc92ef21e09befa252a07c856c4b5afdc51cc2e4 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94540 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@chromium.org> Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton
committed by
Dawn LUCI CQ
Dawn LUCI CQ
parent
3c054304a8
commit
f47887d207
@@ -0,0 +1,6 @@
|
||||
var<workgroup> zero : array<array<i32, 3>, 2>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
var v = zero;
|
||||
}
|
||||
@@ -0,0 +1,20 @@
|
||||
#version 310 es
|
||||
|
||||
shared int zero[2][3];
|
||||
void tint_symbol(uint local_invocation_index) {
|
||||
{
|
||||
for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
|
||||
uint i = (idx / 3u);
|
||||
uint i_1 = (idx % 3u);
|
||||
zero[i][i_1] = 0;
|
||||
}
|
||||
}
|
||||
barrier();
|
||||
int v[2][3] = zero;
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
tint_symbol(gl_LocalInvocationIndex);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,23 @@
|
||||
groupshared int zero[2][3];
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
void main_inner(uint local_invocation_index) {
|
||||
{
|
||||
[loop] for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
|
||||
const uint i = (idx / 3u);
|
||||
const uint i_1 = (idx % 3u);
|
||||
zero[i][i_1] = 0;
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
int v[2][3] = zero;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
main_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,32 @@
|
||||
#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];
|
||||
};
|
||||
|
||||
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<tint_array<int, 3>, 2>* const tint_symbol_1) {
|
||||
for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
|
||||
uint const i = (idx / 3u);
|
||||
uint const i_1 = (idx % 3u);
|
||||
(*(tint_symbol_1))[i][i_1] = 0;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
tint_array<tint_array<int, 3>, 2> v = *(tint_symbol_1);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||
threadgroup tint_array<tint_array<int, 3>, 2> tint_symbol_2;
|
||||
tint_symbol_inner(local_invocation_index, &(tint_symbol_2));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,85 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 51
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %local_invocation_index_1
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %local_invocation_index_1 "local_invocation_index_1"
|
||||
OpName %zero "zero"
|
||||
OpName %main_inner "main_inner"
|
||||
OpName %local_invocation_index "local_invocation_index"
|
||||
OpName %idx "idx"
|
||||
OpName %v "v"
|
||||
OpName %main "main"
|
||||
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
|
||||
OpDecorate %_arr_int_uint_3 ArrayStride 4
|
||||
OpDecorate %_arr__arr_int_uint_3_uint_2 ArrayStride 12
|
||||
%uint = OpTypeInt 32 0
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
|
||||
%int = OpTypeInt 32 1
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_arr_int_uint_3 = OpTypeArray %int %uint_3
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_arr__arr_int_uint_3_uint_2 = OpTypeArray %_arr_int_uint_3 %uint_2
|
||||
%_ptr_Workgroup__arr__arr_int_uint_3_uint_2 = OpTypePointer Workgroup %_arr__arr_int_uint_3_uint_2
|
||||
%zero = OpVariable %_ptr_Workgroup__arr__arr_int_uint_3_uint_2 Workgroup
|
||||
%void = OpTypeVoid
|
||||
%11 = OpTypeFunction %void %uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%18 = OpConstantNull %uint
|
||||
%uint_6 = OpConstant %uint 6
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Workgroup_int = OpTypePointer Workgroup %int
|
||||
%36 = OpConstantNull %int
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%_ptr_Function__arr__arr_int_uint_3_uint_2 = OpTypePointer Function %_arr__arr_int_uint_3_uint_2
|
||||
%45 = OpConstantNull %_arr__arr_int_uint_3_uint_2
|
||||
%46 = OpTypeFunction %void
|
||||
%main_inner = OpFunction %void None %11
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%15 = OpLabel
|
||||
%idx = OpVariable %_ptr_Function_uint Function %18
|
||||
%v = OpVariable %_ptr_Function__arr__arr_int_uint_3_uint_2 Function %45
|
||||
OpStore %idx %local_invocation_index
|
||||
OpBranch %19
|
||||
%19 = OpLabel
|
||||
OpLoopMerge %20 %21 None
|
||||
OpBranch %22
|
||||
%22 = OpLabel
|
||||
%24 = OpLoad %uint %idx
|
||||
%26 = OpULessThan %bool %24 %uint_6
|
||||
%23 = OpLogicalNot %bool %26
|
||||
OpSelectionMerge %28 None
|
||||
OpBranchConditional %23 %29 %28
|
||||
%29 = OpLabel
|
||||
OpBranch %20
|
||||
%28 = OpLabel
|
||||
%30 = OpLoad %uint %idx
|
||||
%31 = OpUDiv %uint %30 %uint_3
|
||||
%32 = OpLoad %uint %idx
|
||||
%33 = OpUMod %uint %32 %uint_3
|
||||
%35 = OpAccessChain %_ptr_Workgroup_int %zero %31 %33
|
||||
OpStore %35 %36
|
||||
OpBranch %21
|
||||
%21 = OpLabel
|
||||
%37 = OpLoad %uint %idx
|
||||
%39 = OpIAdd %uint %37 %uint_1
|
||||
OpStore %idx %39
|
||||
OpBranch %19
|
||||
%20 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%42 = OpLoad %_arr__arr_int_uint_3_uint_2 %zero
|
||||
OpStore %v %42
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %46
|
||||
%48 = OpLabel
|
||||
%50 = OpLoad %uint %local_invocation_index_1
|
||||
%49 = OpFunctionCall %void %main_inner %50
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,6 @@
|
||||
var<workgroup> zero : array<array<i32, 3>, 2>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
var v = zero;
|
||||
}
|
||||
6
test/tint/var/initialization/workgroup/array/i32.wgsl
Normal file
6
test/tint/var/initialization/workgroup/array/i32.wgsl
Normal file
@@ -0,0 +1,6 @@
|
||||
var<workgroup> zero : array<i32, 3>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
var v = zero;
|
||||
}
|
||||
@@ -0,0 +1,19 @@
|
||||
#version 310 es
|
||||
|
||||
shared int zero[3];
|
||||
void tint_symbol(uint local_invocation_index) {
|
||||
{
|
||||
for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
|
||||
uint i = idx;
|
||||
zero[i] = 0;
|
||||
}
|
||||
}
|
||||
barrier();
|
||||
int v[3] = zero;
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
void main() {
|
||||
tint_symbol(gl_LocalInvocationIndex);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,22 @@
|
||||
groupshared int zero[3];
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint local_invocation_index : SV_GroupIndex;
|
||||
};
|
||||
|
||||
void main_inner(uint local_invocation_index) {
|
||||
{
|
||||
[loop] for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
|
||||
const uint i = idx;
|
||||
zero[i] = 0;
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
int v[3] = zero;
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
main_inner(tint_symbol.local_invocation_index);
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,31 @@
|
||||
#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];
|
||||
};
|
||||
|
||||
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<int, 3>* const tint_symbol_1) {
|
||||
for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
|
||||
uint const i = idx;
|
||||
(*(tint_symbol_1))[i] = 0;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
tint_array<int, 3> v = *(tint_symbol_1);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||
threadgroup tint_array<int, 3> tint_symbol_2;
|
||||
tint_symbol_inner(local_invocation_index, &(tint_symbol_2));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,79 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 46
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %local_invocation_index_1
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %local_invocation_index_1 "local_invocation_index_1"
|
||||
OpName %zero "zero"
|
||||
OpName %main_inner "main_inner"
|
||||
OpName %local_invocation_index "local_invocation_index"
|
||||
OpName %idx "idx"
|
||||
OpName %v "v"
|
||||
OpName %main "main"
|
||||
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
|
||||
OpDecorate %_arr_int_uint_3 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
|
||||
%int = OpTypeInt 32 1
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_arr_int_uint_3 = OpTypeArray %int %uint_3
|
||||
%_ptr_Workgroup__arr_int_uint_3 = OpTypePointer Workgroup %_arr_int_uint_3
|
||||
%zero = OpVariable %_ptr_Workgroup__arr_int_uint_3 Workgroup
|
||||
%void = OpTypeVoid
|
||||
%9 = OpTypeFunction %void %uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%16 = OpConstantNull %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Workgroup_int = OpTypePointer Workgroup %int
|
||||
%30 = OpConstantNull %int
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%_ptr_Function__arr_int_uint_3 = OpTypePointer Function %_arr_int_uint_3
|
||||
%40 = OpConstantNull %_arr_int_uint_3
|
||||
%41 = OpTypeFunction %void
|
||||
%main_inner = OpFunction %void None %9
|
||||
%local_invocation_index = OpFunctionParameter %uint
|
||||
%13 = OpLabel
|
||||
%idx = OpVariable %_ptr_Function_uint Function %16
|
||||
%v = OpVariable %_ptr_Function__arr_int_uint_3 Function %40
|
||||
OpStore %idx %local_invocation_index
|
||||
OpBranch %17
|
||||
%17 = OpLabel
|
||||
OpLoopMerge %18 %19 None
|
||||
OpBranch %20
|
||||
%20 = OpLabel
|
||||
%22 = OpLoad %uint %idx
|
||||
%23 = OpULessThan %bool %22 %uint_3
|
||||
%21 = OpLogicalNot %bool %23
|
||||
OpSelectionMerge %25 None
|
||||
OpBranchConditional %21 %26 %25
|
||||
%26 = OpLabel
|
||||
OpBranch %18
|
||||
%25 = OpLabel
|
||||
%27 = OpLoad %uint %idx
|
||||
%29 = OpAccessChain %_ptr_Workgroup_int %zero %27
|
||||
OpStore %29 %30
|
||||
OpBranch %19
|
||||
%19 = OpLabel
|
||||
%31 = OpLoad %uint %idx
|
||||
%33 = OpIAdd %uint %31 %uint_1
|
||||
OpStore %idx %33
|
||||
OpBranch %17
|
||||
%18 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
%37 = OpLoad %_arr_int_uint_3 %zero
|
||||
OpStore %v %37
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %41
|
||||
%43 = OpLabel
|
||||
%45 = OpLoad %uint %local_invocation_index_1
|
||||
%44 = OpFunctionCall %void %main_inner %45
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,6 @@
|
||||
var<workgroup> zero : array<i32, 3>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
var v = zero;
|
||||
}
|
||||
Reference in New Issue
Block a user