[msl-writer] Add support for read-write textures
Insert a fence after every `textureStore()` on a read-write texture to
ensure that writes are visible to subsequent reads.
Bug: tint:2007
Change-Id: I2fdbb272a44d4aa21e486e1d05ddd0565d064655
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/146246
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/test/tint/builtins/gen/literal/textureStore/8ebdc9.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/8ebdc9.wgsl.expected.msl
index e10f138..b8c76ce 100644
--- a/test/tint/builtins/gen/literal/textureStore/8ebdc9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/textureStore/8ebdc9.wgsl.expected.msl
@@ -1,31 +1,33 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_read_write_storage_texture;
-
-@group(1) @binding(0) var arg_0 : texture_storage_2d_array<rgba16float, read_write>;
-
-fn textureStore_8ebdc9() {
- textureStore(arg_0, vec2<u32>(1u), 1u, vec4<f32>(1.0f));
+using namespace metal;
+void textureStore_8ebdc9(texture2d_array<float, access::read_write> tint_symbol_1) {
+ tint_symbol_1.write(float4(1.0f), uint2(uint2(1u)), 1u); tint_symbol_1.fence();
}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
- textureStore_8ebdc9();
- return vec4<f32>();
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d_array<float, access::read_write> tint_symbol_2) {
+ textureStore_8ebdc9(tint_symbol_2);
+ return float4(0.0f);
}
-@fragment
-fn fragment_main() {
- textureStore_8ebdc9();
+vertex tint_symbol vertex_main(texture2d_array<float, access::read_write> tint_symbol_3 [[texture(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
}
-@compute @workgroup_size(1)
-fn compute_main() {
- textureStore_8ebdc9();
+fragment void fragment_main(texture2d_array<float, access::read_write> tint_symbol_4 [[texture(0)]]) {
+ textureStore_8ebdc9(tint_symbol_4);
+ return;
}
-Failed to generate: builtins/gen/literal/textureStore/8ebdc9.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
-enable chromium_experimental_read_write_storage_texture;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+kernel void compute_main(texture2d_array<float, access::read_write> tint_symbol_5 [[texture(0)]]) {
+ textureStore_8ebdc9(tint_symbol_5);
+ return;
+}
