tint: Add PreservePadding transform
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>
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
index 838c6a8..3fda2a5 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
@@ -53,6 +53,12 @@
/* 0x0248 */ tint_array<int8_t, 8> tint_pad_6;
};
+void assign_and_preserve_padding(device tint_array<float3, 2>* const dest, tint_array<float3, 2> value) {
+ for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+ (*(dest))[i] = value[i];
+ }
+}
+
kernel void tint_symbol(device S* tint_symbol_4 [[buffer(0)]]) {
(*(tint_symbol_4)).scalar_f32 = 0.0f;
(*(tint_symbol_4)).scalar_i32 = 0;
@@ -76,7 +82,7 @@
(*(tint_symbol_4)).mat4x3_f32 = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
(*(tint_symbol_4)).mat4x4_f32 = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
tint_array<float3, 2> const tint_symbol_1 = tint_array<float3, 2>{};
- (*(tint_symbol_4)).arr2_vec3_f32 = tint_symbol_1;
+ assign_and_preserve_padding(&((*(tint_symbol_4)).arr2_vec3_f32), tint_symbol_1);
Inner const tint_symbol_2 = Inner{};
(*(tint_symbol_4)).struct_inner = tint_symbol_2;
tint_array<Inner, 4> const tint_symbol_3 = tint_array<Inner, 4>{};
