encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-16 00:17:03 +00:00
Code Issues Packages Projects Releases Wiki Activity

tint: Add PreservePadding transform

Browse Source 
This is used to ensure that assignments to host-visible memory do not
modify padding bytes in structures and arrays. We decompose
assignments of whole structure and array types into member-wise or
element-wise copies, using helper functions.

This is used in all backends except HLSL, which already decomposes
memory accesses.

Bug: tint:1571
Change-Id: Id6de2f917fb80151cc654a7e1c8413ae956f0d61
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/112720
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price
2022-12-06 18:32:19 +00:00
committed by Dawn LUCI CQ
parent f2b8d2b1ab
commit 8753796aac
90 changed files with 5411 additions and 1972 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.glsl
View File

@@ -78,6 +78,21 @@ layout(binding = 1, std430) buffer u_block_ssbo {
S inner[4];
} s;
void assign_and_preserve_padding_1_s_X(uint dest[1], S value) {
s.inner[dest[0]].before = value.before;
s.inner[dest[0]].m = value.m;
s.inner[dest[0]].after = value.after;
}
void assign_and_preserve_padding_s(S value[4]) {
{
for(uint i = 0u; (i < 4u); i = (i + 1u)) {
uint tint_symbol[1] = uint[1](i);
assign_and_preserve_padding_1_s_X(tint_symbol, value[i]);
}
}
}
S conv_S(S_std140 val) {
return S(val.before, f16mat2(val.m_0, val.m_1), val.pad, val.pad_1, val.pad_2, val.pad_3, val.pad_4, val.pad_5, val.pad_6, val.pad_7, val.pad_8, val.pad_9, val.pad_10, val.pad_11, val.pad_12, val.after, val.pad_13, val.pad_14, val.pad_15, val.pad_16, val.pad_17, val.pad_18, val.pad_19, val.pad_20, val.pad_21, val.pad_22, val.pad_23, val.pad_24, val.pad_25, val.pad_26, val.pad_27);
}
@@ -97,8 +112,9 @@ f16mat2 load_u_inner_2_m() {
}
void f() {
s.inner = conv_arr4_S(u.inner);
s.inner[1] = conv_S(u.inner[2u]);
assign_and_preserve_padding_s(conv_arr4_S(u.inner));
uint tint_symbol_1[1] = uint[1](1u);
assign_and_preserve_padding_1_s_X(tint_symbol_1, conv_S(u.inner[2u]));
s.inner[3].m = load_u_inner_2_m();
s.inner[1].m[0] = u.inner[0u].m_1.yx;
}

16
test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
View File

@@ -22,9 +22,21 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
void assign_and_preserve_padding_1(device S* const dest, S value) {
(*(dest)).before = value.before;
(*(dest)).m = value.m;
(*(dest)).after = value.after;
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
kernel void f(device tint_array<S, 4>* tint_symbol [[buffer(1)]], const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
*(tint_symbol) = *(tint_symbol_1);
(*(tint_symbol))[1] = (*(tint_symbol_1))[2];
assign_and_preserve_padding(tint_symbol, *(tint_symbol_1));
assign_and_preserve_padding_1(&((*(tint_symbol))[1]), (*(tint_symbol_1))[2]);
(*(tint_symbol))[3].m = (*(tint_symbol_1))[2].m;
(*(tint_symbol))[1].m[0] = half2((*(tint_symbol_1))[0].m[1]).yx;
return;

262
test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.spvasm
View File

@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 104
; Bound: 142
; Schema: 0
OpCapability Shader
OpCapability Float16
@@ -26,13 +26,20 @@
OpMemberName %S 1 "m"
OpMemberName %S 2 "after"
OpName %s "s"
OpName %assign_and_preserve_padding_1_s_X "assign_and_preserve_padding_1_s_X"
OpName %dest "dest"
OpName %value "value"
OpName %assign_and_preserve_padding_s "assign_and_preserve_padding_s"
OpName %value_0 "value"
OpName %i "i"
OpName %var_for_index "var_for_index"
OpName %conv_S "conv_S"
OpName %val "val"
OpName %conv_arr4_S "conv_arr4_S"
OpName %val_0 "val"
OpName %arr "arr"
OpName %i "i"
OpName %var_for_index "var_for_index"
OpName %i_0 "i"
OpName %var_for_index_1 "var_for_index_1"
OpName %load_u_inner_2_m "load_u_inner_2_m"
OpName %f "f"
OpDecorate %u_block_std140 Block
@@ -55,6 +62,7 @@
OpDecorate %_arr_S_uint_4 ArrayStride 128
OpDecorate %s DescriptorSet 0
OpDecorate %s Binding 1
OpDecorate %_arr_uint_uint_1 ArrayStride 4
%int = OpTypeInt 32 1
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
@@ -71,110 +79,164 @@
%u_block = OpTypeStruct %_arr_S_uint_4
%_ptr_StorageBuffer_u_block = OpTypePointer StorageBuffer %u_block
%s = OpVariable %_ptr_StorageBuffer_u_block StorageBuffer
%17 = OpTypeFunction %S %S_std140
%27 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
%33 = OpConstantNull %_arr_S_uint_4
%_ptr_Function_uint = OpTypePointer Function %uint
%36 = OpConstantNull %uint
%bool = OpTypeBool
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
%49 = OpConstantNull %_arr_S_std140_uint_4
%_ptr_Function_S = OpTypePointer Function %S
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
%void = OpTypeVoid
%uint_1 = OpConstant %uint 1
%62 = OpTypeFunction %mat2v2half
%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
%17 = OpTypeFunction %void %_arr_uint_uint_1 %S
%uint_0 = OpConstant %uint 0
%26 = OpConstantNull %int
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_StorageBuffer_mat2v2half = OpTypePointer StorageBuffer %mat2v2half
%uint_2 = OpConstant %uint 2
%39 = OpTypeFunction %void %_arr_S_uint_4
%43 = OpConstantNull %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%bool = OpTypeBool
%_ptr_Function__arr_S_uint_4 = OpTypePointer Function %_arr_S_uint_4
%58 = OpConstantNull %_arr_S_uint_4
%_ptr_Function_S = OpTypePointer Function %S
%68 = OpTypeFunction %S %S_std140
%78 = OpTypeFunction %_arr_S_uint_4 %_arr_S_std140_uint_4
%_ptr_Function__arr_S_std140_uint_4 = OpTypePointer Function %_arr_S_std140_uint_4
%95 = OpConstantNull %_arr_S_std140_uint_4
%_ptr_Function_S_std140 = OpTypePointer Function %S_std140
%106 = OpTypeFunction %mat2v2half
%_ptr_Uniform_S_std140 = OpTypePointer Uniform %S_std140
%_ptr_Uniform_v2half = OpTypePointer Uniform %v2half
%void = OpTypeVoid
%78 = OpTypeFunction %void
%_ptr_StorageBuffer__arr_S_uint_4 = OpTypePointer StorageBuffer %_arr_S_uint_4
%120 = OpTypeFunction %void
%_ptr_Uniform__arr_S_std140_uint_4 = OpTypePointer Uniform %_arr_S_std140_uint_4
%int_1 = OpConstant %int 1
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%129 = OpConstantComposite %_arr_uint_uint_1 %uint_1
%int_3 = OpConstant %int 3
%_ptr_StorageBuffer_mat2v2half = OpTypePointer StorageBuffer %mat2v2half
%98 = OpConstantNull %int
%int_1 = OpConstant %int 1
%_ptr_StorageBuffer_v2half = OpTypePointer StorageBuffer %v2half
%conv_S = OpFunction %S None %17
%val = OpFunctionParameter %S_std140
%20 = OpLabel
%21 = OpCompositeExtract %int %val 0
%22 = OpCompositeExtract %v2half %val 1
%23 = OpCompositeExtract %v2half %val 2
%24 = OpCompositeConstruct %mat2v2half %22 %23
%25 = OpCompositeExtract %int %val 3
%26 = OpCompositeConstruct %S %21 %24 %25
OpReturnValue %26
OpFunctionEnd
%conv_arr4_S = OpFunction %_arr_S_uint_4 None %27
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
%30 = OpLabel
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %33
%i = OpVariable %_ptr_Function_uint Function %36
%var_for_index = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %49
OpBranch %37
%37 = OpLabel
OpLoopMerge %38 %39 None
OpBranch %40
%40 = OpLabel
%42 = OpLoad %uint %i
%43 = OpULessThan %bool %42 %uint_4
%41 = OpLogicalNot %bool %43
OpSelectionMerge %45 None
OpBranchConditional %41 %46 %45
%46 = OpLabel
OpBranch %38
%45 = OpLabel
OpStore %var_for_index %val_0
%50 = OpLoad %uint %i
%52 = OpAccessChain %_ptr_Function_S %arr %50
%54 = OpLoad %uint %i
%56 = OpAccessChain %_ptr_Function_S_std140 %var_for_index %54
%57 = OpLoad %S_std140 %56
%53 = OpFunctionCall %S %conv_S %57
OpStore %52 %53
OpBranch %39
%39 = OpLabel
%58 = OpLoad %uint %i
%60 = OpIAdd %uint %58 %uint_1
OpStore %i %60
OpBranch %37
%38 = OpLabel
%61 = OpLoad %_arr_S_uint_4 %arr
OpReturnValue %61
OpFunctionEnd
%load_u_inner_2_m = OpFunction %mat2v2half None %62
%64 = OpLabel
%69 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
%72 = OpAccessChain %_ptr_Uniform_v2half %69 %uint_1
%73 = OpLoad %v2half %72
%75 = OpAccessChain %_ptr_Uniform_v2half %69 %uint_2
%76 = OpLoad %v2half %75
%77 = OpCompositeConstruct %mat2v2half %73 %76
OpReturnValue %77
OpFunctionEnd
%f = OpFunction %void None %78
%81 = OpLabel
%83 = OpAccessChain %_ptr_StorageBuffer__arr_S_uint_4 %s %uint_0
%86 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
%87 = OpLoad %_arr_S_std140_uint_4 %86
%84 = OpFunctionCall %_arr_S_uint_4 %conv_arr4_S %87
OpStore %83 %84
%90 = OpAccessChain %_ptr_StorageBuffer_S %s %uint_0 %int_1
%92 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
%93 = OpLoad %S_std140 %92
%91 = OpFunctionCall %S %conv_S %93
OpStore %90 %91
%96 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %s %uint_0 %int_3 %uint_1
%97 = OpFunctionCall %mat2v2half %load_u_inner_2_m
OpStore %96 %97
%100 = OpAccessChain %_ptr_StorageBuffer_v2half %s %uint_0 %int_1 %uint_1 %98
%101 = OpAccessChain %_ptr_Uniform_v2half %u %uint_0 %36 %uint_2
%102 = OpLoad %v2half %101
%103 = OpVectorShuffle %v2half %102 %102 1 0
OpStore %100 %103
%assign_and_preserve_padding_1_s_X = OpFunction %void None %17
%dest = OpFunctionParameter %_arr_uint_uint_1
%value = OpFunctionParameter %S
%24 = OpLabel
%27 = OpCompositeExtract %uint %dest 0
%29 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %27 %uint_0
%30 = OpCompositeExtract %int %value 0
OpStore %29 %30
%31 = OpCompositeExtract %uint %dest 0
%33 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %s %uint_0 %31 %uint_1
%34 = OpCompositeExtract %mat2v2half %value 1
OpStore %33 %34
%35 = OpCompositeExtract %uint %dest 0
%37 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %35 %uint_2
%38 = OpCompositeExtract %int %value 2
OpStore %37 %38
OpReturn
OpFunctionEnd
%assign_and_preserve_padding_s = OpFunction %void None %39
%value_0 = OpFunctionParameter %_arr_S_uint_4
%42 = OpLabel
%i = OpVariable %_ptr_Function_uint Function %43
%var_for_index = OpVariable %_ptr_Function__arr_S_uint_4 Function %58
OpStore %i %43
OpBranch %46
%46 = OpLabel
OpLoopMerge %47 %48 None
OpBranch %49
%49 = OpLabel
%51 = OpLoad %uint %i
%52 = OpULessThan %bool %51 %uint_4
%50 = OpLogicalNot %bool %52
OpSelectionMerge %54 None
OpBranchConditional %50 %55 %54
%55 = OpLabel
OpBranch %47
%54 = OpLabel
OpStore %var_for_index %value_0
%60 = OpLoad %uint %i
%61 = OpCompositeConstruct %_arr_uint_uint_1 %60
%62 = OpLoad %uint %i
%64 = OpAccessChain %_ptr_Function_S %var_for_index %62
%65 = OpLoad %S %64
%59 = OpFunctionCall %void %assign_and_preserve_padding_1_s_X %61 %65
OpBranch %48
%48 = OpLabel
%66 = OpLoad %uint %i
%67 = OpIAdd %uint %66 %uint_1
OpStore %i %67
OpBranch %46
%47 = OpLabel
OpReturn
OpFunctionEnd
%conv_S = OpFunction %S None %68
%val = OpFunctionParameter %S_std140
%71 = OpLabel
%72 = OpCompositeExtract %int %val 0
%73 = OpCompositeExtract %v2half %val 1
%74 = OpCompositeExtract %v2half %val 2
%75 = OpCompositeConstruct %mat2v2half %73 %74
%76 = OpCompositeExtract %int %val 3
%77 = OpCompositeConstruct %S %72 %75 %76
OpReturnValue %77
OpFunctionEnd
%conv_arr4_S = OpFunction %_arr_S_uint_4 None %78
%val_0 = OpFunctionParameter %_arr_S_std140_uint_4
%81 = OpLabel
%arr = OpVariable %_ptr_Function__arr_S_uint_4 Function %58
%i_0 = OpVariable %_ptr_Function_uint Function %43
%var_for_index_1 = OpVariable %_ptr_Function__arr_S_std140_uint_4 Function %95
OpBranch %84
%84 = OpLabel
OpLoopMerge %85 %86 None
OpBranch %87
%87 = OpLabel
%89 = OpLoad %uint %i_0
%90 = OpULessThan %bool %89 %uint_4
%88 = OpLogicalNot %bool %90
OpSelectionMerge %91 None
OpBranchConditional %88 %92 %91
%92 = OpLabel
OpBranch %85
%91 = OpLabel
OpStore %var_for_index_1 %val_0
%96 = OpLoad %uint %i_0
%97 = OpAccessChain %_ptr_Function_S %arr %96
%99 = OpLoad %uint %i_0
%101 = OpAccessChain %_ptr_Function_S_std140 %var_for_index_1 %99
%102 = OpLoad %S_std140 %101
%98 = OpFunctionCall %S %conv_S %102
OpStore %97 %98
OpBranch %86
%86 = OpLabel
%103 = OpLoad %uint %i_0
%104 = OpIAdd %uint %103 %uint_1
OpStore %i_0 %104
OpBranch %84
%85 = OpLabel
%105 = OpLoad %_arr_S_uint_4 %arr
OpReturnValue %105
OpFunctionEnd
%load_u_inner_2_m = OpFunction %mat2v2half None %106
%108 = OpLabel
%111 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
%114 = OpAccessChain %_ptr_Uniform_v2half %111 %uint_1
%115 = OpLoad %v2half %114
%117 = OpAccessChain %_ptr_Uniform_v2half %111 %uint_2
%118 = OpLoad %v2half %117
%119 = OpCompositeConstruct %mat2v2half %115 %118
OpReturnValue %119
OpFunctionEnd
%f = OpFunction %void None %120
%122 = OpLabel
%126 = OpAccessChain %_ptr_Uniform__arr_S_std140_uint_4 %u %uint_0
%127 = OpLoad %_arr_S_std140_uint_4 %126
%124 = OpFunctionCall %_arr_S_uint_4 %conv_arr4_S %127
%123 = OpFunctionCall %void %assign_and_preserve_padding_s %124
%131 = OpAccessChain %_ptr_Uniform_S_std140 %u %uint_0 %uint_2
%132 = OpLoad %S_std140 %131
%130 = OpFunctionCall %S %conv_S %132
%128 = OpFunctionCall %void %assign_and_preserve_padding_1_s_X %129 %130
%134 = OpAccessChain %_ptr_StorageBuffer_mat2v2half %s %uint_0 %int_3 %uint_1
%135 = OpFunctionCall %mat2v2half %load_u_inner_2_m
OpStore %134 %135
%138 = OpAccessChain %_ptr_StorageBuffer_v2half %s %uint_0 %int_1 %uint_1 %26
%139 = OpAccessChain %_ptr_Uniform_v2half %u %uint_0 %43 %uint_2
%140 = OpLoad %v2half %139
%141 = OpVectorShuffle %v2half %140 %140 1 0
OpStore %138 %141
OpReturn
OpFunctionEnd