[msl] Add polyfill for textureStore()
Insert a fence after texture writes to read-write storage textures to
ensure correct memory ordering.
Add MSL member function definitions for `write()` and `fence()`.
Bug: 42251016
Change-Id: I59436da345e7940c1f4d898c8fb646ad1c559c7b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/192944
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/test/tint/builtins/gen/literal/textureStore/7792fa.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/7792fa.wgsl.expected.ir.msl
index b1399f7..63afe0d 100644
--- a/test/tint/builtins/gen/literal/textureStore/7792fa.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureStore/7792fa.wgsl.expected.ir.msl
@@ -1,39 +1,29 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ texture3d<int, access::read_write> arg_0;
+};
+struct vertex_main_outputs {
+ float4 tint_symbol [[position]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:1164 internal compiler error: $B1: { # root
- %arg_0:ptr<handle, texture_storage_3d<rg32sint, read_write>, read> = var @binding_point(1, 0)
+void textureStore_7792fa(tint_module_vars_struct tint_module_vars) {
+ tint_module_vars.arg_0.write(int4(1), uint3(1u));
+ tint_module_vars.arg_0.fence();
}
-
-%textureStore_7792fa = func():void {
- $B2: {
- %3:texture_storage_3d<rg32sint, read_write> = load %arg_0
- %4:void = textureStore %3, vec3<u32>(1u), vec4<i32>(1i)
- ret
- }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+ textureStore_7792fa(tint_module_vars);
+ return float4(0.0f);
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %6:void = call %textureStore_7792fa
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(texture3d<int, access::read_write> arg_0 [[texture(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0};
+ textureStore_7792fa(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B4: {
- %8:void = call %textureStore_7792fa
- ret
- }
+kernel void compute_main(texture3d<int, access::read_write> arg_0 [[texture(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0};
+ textureStore_7792fa(tint_module_vars);
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %10:void = call %textureStore_7792fa
- ret
- }
+vertex vertex_main_outputs vertex_main(texture3d<int, access::read_write> arg_0 [[texture(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0};
+ return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
}
-
-invalid access control for storage texture
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
