Generate missing hlsl e2e tests
Add missing hlsl-fxc tests. All subgroup operations are
unsupported before Shader Model 6.0 so they are skipped.
Bug: 354738715
Change-Id: I6ceaf8c4358e5ad947a117fa0c2d278ceafdc6e4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/202695
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Natalie Chouinard <chouinard@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Auto-Submit: Natalie Chouinard <chouinard@google.com>
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..427c73a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadBroadcast_0464d1() {
+ vector<float16_t, 2> res = QuadReadLaneAt((float16_t(1.0h)).xx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_0464d1());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_0464d1());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..131acfa
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadBroadcast_0639ea() {
+ int res = QuadReadLaneAt(1, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_0639ea()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_0639ea()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0870451
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadBroadcast_0cc513() {
+ float3 res = QuadReadLaneAt((1.0f).xxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0cc513()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0cc513()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7a95b72
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadBroadcast_0e0e6e() {
+ int3 res = QuadReadLaneAt((1).xxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0e0e6e()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0e0e6e()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4b61ccf
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadBroadcast_2d0b7d() {
+ uint4 res = QuadReadLaneAt((1u).xxxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadBroadcast_2d0b7d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadBroadcast_2d0b7d());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..feadca4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadBroadcast_355db5() {
+ float3 res = QuadReadLaneAt((1.0f).xxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_355db5()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_355db5()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..502aa25
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadBroadcast_3c3824() {
+ vector<float16_t, 2> res = QuadReadLaneAt((float16_t(1.0h)).xx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_3c3824());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_3c3824());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..86329d0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadBroadcast_4d9898() {
+ vector<float16_t, 4> res = QuadReadLaneAt((float16_t(1.0h)).xxxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_4d9898());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_4d9898());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d937563
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadBroadcast_641316() {
+ uint2 res = QuadReadLaneAt((1u).xx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadBroadcast_641316());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadBroadcast_641316());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6cabe82
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadBroadcast_704803() {
+ int3 res = QuadReadLaneAt((1).xxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_704803()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_704803()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..de1b947
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadBroadcast_76f499() {
+ int4 res = QuadReadLaneAt((1).xxxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_76f499()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_76f499()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e943315
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadBroadcast_78129b() {
+ float16_t res = QuadReadLaneAt(float16_t(1.0h), 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_78129b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_78129b());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..29caeee
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadBroadcast_796753() {
+ vector<float16_t, 3> res = QuadReadLaneAt((float16_t(1.0h)).xxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_796753());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_796753());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e6d76d7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadBroadcast_820991() {
+ float4 res = QuadReadLaneAt((1.0f).xxxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_820991()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_820991()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2ccad8d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadBroadcast_960c6b() {
+ float res = QuadReadLaneAt(1.0f, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_960c6b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_960c6b()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..62abb74
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadBroadcast_9d802c() {
+ float4 res = QuadReadLaneAt((1.0f).xxxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_9d802c()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_9d802c()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..adf330d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadBroadcast_a2d2b4() {
+ uint res = QuadReadLaneAt(1u, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadBroadcast_a2d2b4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadBroadcast_a2d2b4());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..08876a1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadBroadcast_ae401e() {
+ uint3 res = QuadReadLaneAt((1u).xxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadBroadcast_ae401e());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadBroadcast_ae401e());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fe631fe
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadBroadcast_b68331() {
+ uint4 res = QuadReadLaneAt((1u).xxxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadBroadcast_b68331());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadBroadcast_b68331());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2c0afb7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadBroadcast_bed00b() {
+ int4 res = QuadReadLaneAt((1).xxxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_bed00b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_bed00b()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..dc4d251
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadBroadcast_c0e704() {
+ int2 res = QuadReadLaneAt((1).xx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_c0e704()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_c0e704()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6ff0ed8
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadBroadcast_cd3624() {
+ float2 res = QuadReadLaneAt((1.0f).xx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cd3624()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cd3624()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4a247d4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadBroadcast_cebc6a() {
+ float16_t res = QuadReadLaneAt(float16_t(1.0h), 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_cebc6a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_cebc6a());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8607af7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadBroadcast_cfbf48() {
+ float2 res = QuadReadLaneAt((1.0f).xx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cfbf48()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cfbf48()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ffd0d86
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadBroadcast_e6d39d() {
+ float res = QuadReadLaneAt(1.0f, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_e6d39d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_e6d39d()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3ff1e1b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadBroadcast_e6d948() {
+ uint res = QuadReadLaneAt(1u, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadBroadcast_e6d948());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadBroadcast_e6d948());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fe339cc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadBroadcast_e7c301() {
+ vector<float16_t, 4> res = QuadReadLaneAt((float16_t(1.0h)).xxxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_e7c301());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_e7c301());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1004505
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadBroadcast_ef7d5d() {
+ vector<float16_t, 3> res = QuadReadLaneAt((float16_t(1.0h)).xxx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_ef7d5d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_ef7d5d());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6b2de11
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadBroadcast_f1e8ec() {
+ uint3 res = QuadReadLaneAt((1u).xxx, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadBroadcast_f1e8ec());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadBroadcast_f1e8ec());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..87c3166
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadBroadcast_f5f923() {
+ int2 res = QuadReadLaneAt((1).xx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_f5f923()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_f5f923()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7ee8efb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadBroadcast_f60448() {
+ uint2 res = QuadReadLaneAt((1u).xx, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadBroadcast_f60448());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadBroadcast_f60448());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cef1f25
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadBroadcast_f9d579() {
+ int res = QuadReadLaneAt(1, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_f9d579()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_f9d579()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1ba078f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapDiagonal_15ac75() {
+ vector<float16_t, 2> res = QuadReadAcrossDiagonal((float16_t(1.0h)).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapDiagonal_15ac75());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapDiagonal_15ac75());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..000d799
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapDiagonal_2be5e7() {
+ float16_t res = QuadReadAcrossDiagonal(float16_t(1.0h));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapDiagonal_2be5e7());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapDiagonal_2be5e7());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..031c449
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapDiagonal_331804() {
+ float4 res = QuadReadAcrossDiagonal((1.0f).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_331804()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_331804()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e093560
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapDiagonal_348173() {
+ uint2 res = QuadReadAcrossDiagonal((1u).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapDiagonal_348173());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapDiagonal_348173());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0107499
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapDiagonal_486196() {
+ float res = QuadReadAcrossDiagonal(1.0f);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_486196()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_486196()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3a93f97
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapDiagonal_730e40() {
+ uint res = QuadReadAcrossDiagonal(1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapDiagonal_730e40());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapDiagonal_730e40());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..af7706b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapDiagonal_8077c8() {
+ float2 res = QuadReadAcrossDiagonal((1.0f).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_8077c8()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_8077c8()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fb4be26
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapDiagonal_856536() {
+ uint3 res = QuadReadAcrossDiagonal((1u).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapDiagonal_856536());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapDiagonal_856536());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..713a44d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapDiagonal_9ccb38() {
+ int res = QuadReadAcrossDiagonal(1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_9ccb38()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_9ccb38()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fa4be0a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapDiagonal_a090b0() {
+ int2 res = QuadReadAcrossDiagonal((1).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_a090b0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_a090b0()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c3b3869
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapDiagonal_a665b1() {
+ int4 res = QuadReadAcrossDiagonal((1).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_a665b1()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_a665b1()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a080bdc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapDiagonal_a82e1d() {
+ int3 res = QuadReadAcrossDiagonal((1).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_a82e1d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_a82e1d()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5f55a68
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapDiagonal_af19a5() {
+ vector<float16_t, 4> res = QuadReadAcrossDiagonal((float16_t(1.0h)).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapDiagonal_af19a5());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapDiagonal_af19a5());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5f2bab2
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapDiagonal_b905fc() {
+ float3 res = QuadReadAcrossDiagonal((1.0f).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_b905fc()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_b905fc()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..30e5efc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapDiagonal_c31636() {
+ uint4 res = QuadReadAcrossDiagonal((1u).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapDiagonal_c31636());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapDiagonal_c31636());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2c23281
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapDiagonal_e4bec8() {
+ vector<float16_t, 3> res = QuadReadAcrossDiagonal((float16_t(1.0h)).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapDiagonal_e4bec8());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapDiagonal_e4bec8());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..baec144
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapX_02834c() {
+ vector<float16_t, 4> res = QuadReadAcrossX((float16_t(1.0h)).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapX_02834c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapX_02834c());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5015f25
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapX_053f3b() {
+ int3 res = QuadReadAcrossX((1).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_053f3b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_053f3b()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..bc4e19c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapX_07f1fc() {
+ uint4 res = QuadReadAcrossX((1u).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapX_07f1fc());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapX_07f1fc());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1e3ee38
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapX_150d6f() {
+ float3 res = QuadReadAcrossX((1.0f).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_150d6f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_150d6f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d331d5a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapX_19f8ce() {
+ uint2 res = QuadReadAcrossX((1u).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapX_19f8ce());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapX_19f8ce());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..599a853
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapX_1e1086() {
+ int res = QuadReadAcrossX(1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_1e1086()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_1e1086()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0731839
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapX_69af6a() {
+ float4 res = QuadReadAcrossX((1.0f).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_69af6a()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_69af6a()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5de71a6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapX_8203ad() {
+ uint res = QuadReadAcrossX(1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapX_8203ad());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapX_8203ad());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7e4720a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapX_879738() {
+ float2 res = QuadReadAcrossX((1.0f).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_879738()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_879738()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..714787c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapX_9bea80() {
+ float res = QuadReadAcrossX(1.0f);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_9bea80()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_9bea80()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7ce77ba
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapX_a4e103() {
+ float16_t res = QuadReadAcrossX(float16_t(1.0h));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapX_a4e103());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapX_a4e103());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..456db0b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapX_b1a5fe() {
+ int2 res = QuadReadAcrossX((1).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_b1a5fe()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_b1a5fe()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..da74984
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapX_bc2013() {
+ vector<float16_t, 3> res = QuadReadAcrossX((float16_t(1.0h)).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapX_bc2013());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapX_bc2013());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c98c509
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapX_bddb9f() {
+ uint3 res = QuadReadAcrossX((1u).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapX_bddb9f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapX_bddb9f());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..80d4eb7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapX_d60cec() {
+ vector<float16_t, 2> res = QuadReadAcrossX((float16_t(1.0h)).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapX_d60cec());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapX_d60cec());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3c94ef7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapX_edfa1f() {
+ int4 res = QuadReadAcrossX((1).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_edfa1f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_edfa1f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..72aecf6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapY_06a67c() {
+ uint3 res = QuadReadAcrossY((1u).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapY_06a67c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapY_06a67c());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3fe2af4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapY_0c4938() {
+ uint res = QuadReadAcrossY(1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapY_0c4938());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapY_0c4938());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5ac9f13
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapY_0d05a8() {
+ int2 res = QuadReadAcrossY((1).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_0d05a8()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_0d05a8()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b3469ec
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapY_14bb9a() {
+ int4 res = QuadReadAcrossY((1).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_14bb9a()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_14bb9a()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..dfde4ce
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapY_1f1a06() {
+ float2 res = QuadReadAcrossY((1.0f).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_1f1a06()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_1f1a06()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e931439
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapY_264908() {
+ vector<float16_t, 3> res = QuadReadAcrossY((float16_t(1.0h)).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapY_264908());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapY_264908());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..aa707db
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapY_5b2e67() {
+ vector<float16_t, 4> res = QuadReadAcrossY((float16_t(1.0h)).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapY_5b2e67());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapY_5b2e67());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ffb46b0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapY_6f6bc9() {
+ float res = QuadReadAcrossY(1.0f);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_6f6bc9()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_6f6bc9()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4bf24e7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapY_9277e9() {
+ float16_t res = QuadReadAcrossY(float16_t(1.0h));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapY_9277e9());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapY_9277e9());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7382420
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapY_94ab6d() {
+ int res = QuadReadAcrossY(1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_94ab6d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_94ab6d()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..feeb9d4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapY_a27e1c() {
+ uint2 res = QuadReadAcrossY((1u).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapY_a27e1c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapY_a27e1c());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b0bb8c9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapY_a50fcb() {
+ vector<float16_t, 2> res = QuadReadAcrossY((float16_t(1.0h)).xx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapY_a50fcb());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapY_a50fcb());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0485f39
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapY_b9d9e7() {
+ float4 res = QuadReadAcrossY((1.0f).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_b9d9e7()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_b9d9e7()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1d5c49b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapY_bb697b() {
+ uint4 res = QuadReadAcrossY((1u).xxxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapY_bb697b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapY_bb697b());
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5567c2e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapY_be4e72() {
+ int3 res = QuadReadAcrossY((1).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_be4e72()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_be4e72()));
+}
+
diff --git a/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b5fdea52
--- /dev/null
+++ b/test/tint/builtins/gen/literal/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapY_d1ab4d() {
+ float3 res = QuadReadAcrossY((1.0f).xxx);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_d1ab4d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_d1ab4d()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupElect/3943d6.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/subgroupElect/3943d6.wgsl.expected.ir.msl
new file mode 100644
index 0000000..836711c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupElect/3943d6.wgsl.expected.ir.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupElect_3943d6() {
+ bool res = simd_is_first();
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupElect_3943d6();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupElect_3943d6();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6d67c76
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleDown_10eb45() {
+ vector<float16_t, 4> res = WaveReadLaneAt((float16_t(1.0h)).xxxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleDown_10eb45());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleDown_10eb45());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c61d1dc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleDown_1b530f() {
+ int3 res = WaveReadLaneAt((1).xxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_1b530f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_1b530f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f7425dd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleDown_257ff0() {
+ float4 res = WaveReadLaneAt((1.0f).xxxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_257ff0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_257ff0()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cc79f2f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleDown_313d9b() {
+ int4 res = WaveReadLaneAt((1).xxxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_313d9b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_313d9b()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3f1ef46
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleDown_57b1e8() {
+ vector<float16_t, 2> res = WaveReadLaneAt((float16_t(1.0h)).xx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleDown_57b1e8());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleDown_57b1e8());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..be9ad07
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleDown_5d8b9f() {
+ float3 res = WaveReadLaneAt((1.0f).xxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_5d8b9f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_5d8b9f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0e7a9e1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleDown_63fdb0() {
+ vector<float16_t, 3> res = WaveReadLaneAt((float16_t(1.0h)).xxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleDown_63fdb0());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleDown_63fdb0());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..98abf96
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleDown_642789() {
+ uint3 res = WaveReadLaneAt((1u).xxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleDown_642789());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleDown_642789());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f68ca01
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleDown_7a0cf5() {
+ float2 res = WaveReadLaneAt((1.0f).xx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_7a0cf5()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_7a0cf5()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6cabb92
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleDown_7f8886() {
+ float res = WaveReadLaneAt(1.0f, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_7f8886()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_7f8886()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fba126c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleDown_9c6714() {
+ float16_t res = WaveReadLaneAt(float16_t(1.0h), (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleDown_9c6714());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleDown_9c6714());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d5a3bba
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleDown_b41899() {
+ int2 res = WaveReadLaneAt((1).xx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_b41899()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_b41899()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..992849d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleDown_c9f1c4() {
+ uint2 res = WaveReadLaneAt((1u).xx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleDown_c9f1c4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleDown_c9f1c4());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..12d2030
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleDown_d269eb() {
+ int res = WaveReadLaneAt(1, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_d269eb()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_d269eb()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..91af3e5
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleDown_d46304() {
+ uint4 res = WaveReadLaneAt((1u).xxxx, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleDown_d46304());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleDown_d46304());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4ad4841
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleDown_d90c2f() {
+ uint res = WaveReadLaneAt(1u, (WaveGetLaneIndex() + 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleDown_d90c2f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleDown_d90c2f());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b9a91b6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleUp_0990cd() {
+ vector<float16_t, 4> res = WaveReadLaneAt((float16_t(1.0h)).xxxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleUp_0990cd());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleUp_0990cd());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c5c5823
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleUp_1bb93f() {
+ int res = WaveReadLaneAt(1, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_1bb93f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_1bb93f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d8186b2
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleUp_23c7ca() {
+ float res = WaveReadLaneAt(1.0f, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_23c7ca()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_23c7ca()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e680449
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleUp_3242a6() {
+ uint res = WaveReadLaneAt(1u, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleUp_3242a6());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleUp_3242a6());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..583567f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleUp_33d495() {
+ float4 res = WaveReadLaneAt((1.0f).xxxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_33d495()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_33d495()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ebd66c0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleUp_3e609f() {
+ int4 res = WaveReadLaneAt((1).xxxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_3e609f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_3e609f()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5a1cb15
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleUp_58de69() {
+ uint2 res = WaveReadLaneAt((1u).xx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleUp_58de69());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleUp_58de69());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1202e70
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleUp_868e52() {
+ vector<float16_t, 3> res = WaveReadLaneAt((float16_t(1.0h)).xxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleUp_868e52());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleUp_868e52());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a7c833b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleUp_87c9d6() {
+ float3 res = WaveReadLaneAt((1.0f).xxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_87c9d6()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_87c9d6()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0663a7b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleUp_88eb07() {
+ uint4 res = WaveReadLaneAt((1u).xxxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleUp_88eb07());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleUp_88eb07());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..adbacb8
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleUp_8a63f3() {
+ int3 res = WaveReadLaneAt((1).xxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_8a63f3()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_8a63f3()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2935c56
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleUp_a2075a() {
+ vector<float16_t, 2> res = WaveReadLaneAt((float16_t(1.0h)).xx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleUp_a2075a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleUp_a2075a());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4f52f78
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleUp_abaea0() {
+ uint3 res = WaveReadLaneAt((1u).xxx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleUp_abaea0());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleUp_abaea0());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..32e77e5
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleUp_b58804() {
+ float2 res = WaveReadLaneAt((1.0f).xx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_b58804()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_b58804()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0df5a72
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleUp_bbf7f4() {
+ float16_t res = WaveReadLaneAt(float16_t(1.0h), (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleUp_bbf7f4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleUp_bbf7f4());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e368ac3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleUp_db5bcb() {
+ int2 res = WaveReadLaneAt((1).xx, (WaveGetLaneIndex() - 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_db5bcb()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_db5bcb()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..bd9fe7f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleXor_071aa0() {
+ int2 res = WaveReadLaneAt((1).xx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_071aa0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_071aa0()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fe88293
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleXor_08f588() {
+ uint4 res = WaveReadLaneAt((1u).xxxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleXor_08f588());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleXor_08f588());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..795d16e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleXor_1d36b6() {
+ float res = WaveReadLaneAt(1.0f, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_1d36b6()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_1d36b6()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5051174
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleXor_1e247f() {
+ vector<float16_t, 2> res = WaveReadLaneAt((float16_t(1.0h)).xx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleXor_1e247f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleXor_1e247f());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2e160b3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleXor_1f2590() {
+ vector<float16_t, 3> res = WaveReadLaneAt((float16_t(1.0h)).xxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleXor_1f2590());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleXor_1f2590());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9bb71bb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleXor_2e033d() {
+ vector<float16_t, 4> res = WaveReadLaneAt((float16_t(1.0h)).xxxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleXor_2e033d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleXor_2e033d());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3000d5b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleXor_445e83() {
+ int res = WaveReadLaneAt(1, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_445e83()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_445e83()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e72c01a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleXor_7435fe() {
+ int3 res = WaveReadLaneAt((1).xxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_7435fe()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_7435fe()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..782801a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleXor_80b6e9() {
+ uint res = WaveReadLaneAt(1u, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleXor_80b6e9());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleXor_80b6e9());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8be3f27
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleXor_9f945a() {
+ uint3 res = WaveReadLaneAt((1u).xxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleXor_9f945a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleXor_9f945a());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3c6eb90
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleXor_bdddba() {
+ int4 res = WaveReadLaneAt((1).xxxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_bdddba()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_bdddba()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cab7a6b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleXor_c88290() {
+ float4 res = WaveReadLaneAt((1.0f).xxxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_c88290()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_c88290()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..406512c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleXor_caa816() {
+ float3 res = WaveReadLaneAt((1.0f).xxx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_caa816()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_caa816()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ee9ec99
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleXor_d224ab() {
+ float16_t res = WaveReadLaneAt(float16_t(1.0h), (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleXor_d224ab());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleXor_d224ab());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..990908a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleXor_e3c10b() {
+ uint2 res = WaveReadLaneAt((1u).xx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleXor_e3c10b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleXor_e3c10b());
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e5b3655
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleXor_f7b453() {
+ float2 res = WaveReadLaneAt((1.0f).xx, (WaveGetLaneIndex() ^ 1u));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_f7b453()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_f7b453()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5d0ea92
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/0464d1.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadBroadcast_0464d1() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_0464d1());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_0464d1());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c6b7e37
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/0639ea.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadBroadcast_0639ea() {
+ int arg_0 = 1;
+ int res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_0639ea()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_0639ea()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ff11d382
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/0cc513.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadBroadcast_0cc513() {
+ float3 arg_0 = (1.0f).xxx;
+ float3 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0cc513()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0cc513()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5b927d5
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/0e0e6e.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadBroadcast_0e0e6e() {
+ int3 arg_0 = (1).xxx;
+ int3 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0e0e6e()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_0e0e6e()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..68b0e10
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/2d0b7d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadBroadcast_2d0b7d() {
+ uint4 arg_0 = (1u).xxxx;
+ uint4 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadBroadcast_2d0b7d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadBroadcast_2d0b7d());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..51b8fc3
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/355db5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadBroadcast_355db5() {
+ float3 arg_0 = (1.0f).xxx;
+ float3 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_355db5()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_355db5()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ebce2a0
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/3c3824.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadBroadcast_3c3824() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_3c3824());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadBroadcast_3c3824());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..2bec8fb
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/4d9898.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadBroadcast_4d9898() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_4d9898());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_4d9898());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7829328
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/641316.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadBroadcast_641316() {
+ uint2 arg_0 = (1u).xx;
+ uint2 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadBroadcast_641316());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadBroadcast_641316());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c06795a
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/704803.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadBroadcast_704803() {
+ int3 arg_0 = (1).xxx;
+ int3 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_704803()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadBroadcast_704803()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3a283d7
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/76f499.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadBroadcast_76f499() {
+ int4 arg_0 = (1).xxxx;
+ int4 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_76f499()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_76f499()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..de543ed8
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/78129b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadBroadcast_78129b() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_78129b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_78129b());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7a08768
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/796753.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadBroadcast_796753() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_796753());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_796753());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c2b1408
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/820991.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadBroadcast_820991() {
+ float4 arg_0 = (1.0f).xxxx;
+ float4 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_820991()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_820991()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7f72a49
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/960c6b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadBroadcast_960c6b() {
+ float arg_0 = 1.0f;
+ float res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_960c6b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_960c6b()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..81fe53b
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/9d802c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadBroadcast_9d802c() {
+ float4 arg_0 = (1.0f).xxxx;
+ float4 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_9d802c()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_9d802c()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b5d6440
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/a2d2b4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadBroadcast_a2d2b4() {
+ uint arg_0 = 1u;
+ uint res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadBroadcast_a2d2b4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadBroadcast_a2d2b4());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..39d49ea
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/ae401e.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadBroadcast_ae401e() {
+ uint3 arg_0 = (1u).xxx;
+ uint3 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadBroadcast_ae401e());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadBroadcast_ae401e());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..119d752
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/b68331.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadBroadcast_b68331() {
+ uint4 arg_0 = (1u).xxxx;
+ uint4 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadBroadcast_b68331());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadBroadcast_b68331());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8bc33a6
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/bed00b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadBroadcast_bed00b() {
+ int4 arg_0 = (1).xxxx;
+ int4 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_bed00b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadBroadcast_bed00b()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..535ca8e
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/c0e704.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadBroadcast_c0e704() {
+ int2 arg_0 = (1).xx;
+ int2 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_c0e704()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_c0e704()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f15a165
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/cd3624.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadBroadcast_cd3624() {
+ float2 arg_0 = (1.0f).xx;
+ float2 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cd3624()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cd3624()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4e2ced9
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/cebc6a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadBroadcast_cebc6a() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_cebc6a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadBroadcast_cebc6a());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..877f9ee
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/cfbf48.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadBroadcast_cfbf48() {
+ float2 arg_0 = (1.0f).xx;
+ float2 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cfbf48()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_cfbf48()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..42541c6
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/e6d39d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadBroadcast_e6d39d() {
+ float arg_0 = 1.0f;
+ float res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_e6d39d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_e6d39d()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..af5f53f
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/e6d948.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadBroadcast_e6d948() {
+ uint arg_0 = 1u;
+ uint res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadBroadcast_e6d948());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadBroadcast_e6d948());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..49dea4ce
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/e7c301.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadBroadcast_e7c301() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_e7c301());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadBroadcast_e7c301());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ab93ab5
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/ef7d5d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadBroadcast_ef7d5d() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_ef7d5d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadBroadcast_ef7d5d());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6d8c1a4
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/f1e8ec.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadBroadcast_f1e8ec() {
+ uint3 arg_0 = (1u).xxx;
+ uint3 res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadBroadcast_f1e8ec());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadBroadcast_f1e8ec());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8c1abc4
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/f5f923.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadBroadcast_f5f923() {
+ int2 arg_0 = (1).xx;
+ int2 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_f5f923()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadBroadcast_f5f923()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..80672ca
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/f60448.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadBroadcast_f60448() {
+ uint2 arg_0 = (1u).xx;
+ uint2 res = QuadReadLaneAt(arg_0, 1u);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadBroadcast_f60448());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadBroadcast_f60448());
+}
+
diff --git a/test/tint/builtins/gen/var/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..26e5a26
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadBroadcast/f9d579.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadBroadcast_f9d579() {
+ int arg_0 = 1;
+ int res = QuadReadLaneAt(arg_0, 1);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_f9d579()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadBroadcast_f9d579()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a093c52
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/15ac75.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapDiagonal_15ac75() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapDiagonal_15ac75());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapDiagonal_15ac75());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a9d857c
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/2be5e7.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapDiagonal_2be5e7() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapDiagonal_2be5e7());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapDiagonal_2be5e7());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..969c1ab
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/331804.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapDiagonal_331804() {
+ float4 arg_0 = (1.0f).xxxx;
+ float4 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_331804()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_331804()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..00bd53e29
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/348173.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapDiagonal_348173() {
+ uint2 arg_0 = (1u).xx;
+ uint2 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapDiagonal_348173());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapDiagonal_348173());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9d6715c
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/486196.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapDiagonal_486196() {
+ float arg_0 = 1.0f;
+ float res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_486196()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_486196()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..75e07e8
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/730e40.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapDiagonal_730e40() {
+ uint arg_0 = 1u;
+ uint res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapDiagonal_730e40());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapDiagonal_730e40());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7d520d1
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/8077c8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapDiagonal_8077c8() {
+ float2 arg_0 = (1.0f).xx;
+ float2 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_8077c8()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_8077c8()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..25866a0
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/856536.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapDiagonal_856536() {
+ uint3 arg_0 = (1u).xxx;
+ uint3 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapDiagonal_856536());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapDiagonal_856536());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a2d06f6
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/9ccb38.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapDiagonal_9ccb38() {
+ int arg_0 = 1;
+ int res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_9ccb38()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapDiagonal_9ccb38()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f8861fa
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a090b0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapDiagonal_a090b0() {
+ int2 arg_0 = (1).xx;
+ int2 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_a090b0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapDiagonal_a090b0()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f6fd704
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a665b1.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapDiagonal_a665b1() {
+ int4 arg_0 = (1).xxxx;
+ int4 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_a665b1()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapDiagonal_a665b1()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ab49048
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/a82e1d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapDiagonal_a82e1d() {
+ int3 arg_0 = (1).xxx;
+ int3 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_a82e1d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_a82e1d()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..12ecb02
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/af19a5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapDiagonal_af19a5() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapDiagonal_af19a5());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapDiagonal_af19a5());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..83c8cfd
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/b905fc.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapDiagonal_b905fc() {
+ float3 arg_0 = (1.0f).xxx;
+ float3 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_b905fc()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapDiagonal_b905fc()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0f3a528
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/c31636.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapDiagonal_c31636() {
+ uint4 arg_0 = (1u).xxxx;
+ uint4 res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapDiagonal_c31636());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapDiagonal_c31636());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9f098f1
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapDiagonal/e4bec8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapDiagonal_e4bec8() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> res = QuadReadAcrossDiagonal(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapDiagonal_e4bec8());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapDiagonal_e4bec8());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0f3fd1d
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/02834c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapX_02834c() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapX_02834c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapX_02834c());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e2460f2
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/053f3b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapX_053f3b() {
+ int3 arg_0 = (1).xxx;
+ int3 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_053f3b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_053f3b()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fbcb010
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/07f1fc.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapX_07f1fc() {
+ uint4 arg_0 = (1u).xxxx;
+ uint4 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapX_07f1fc());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapX_07f1fc());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ab4edf2
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/150d6f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapX_150d6f() {
+ float3 arg_0 = (1.0f).xxx;
+ float3 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_150d6f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapX_150d6f()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7c9f44e
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/19f8ce.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapX_19f8ce() {
+ uint2 arg_0 = (1u).xx;
+ uint2 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapX_19f8ce());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapX_19f8ce());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ccda9d2
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/1e1086.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapX_1e1086() {
+ int arg_0 = 1;
+ int res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_1e1086()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_1e1086()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1a2271a
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/69af6a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapX_69af6a() {
+ float4 arg_0 = (1.0f).xxxx;
+ float4 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_69af6a()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_69af6a()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1d01356
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/8203ad.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapX_8203ad() {
+ uint arg_0 = 1u;
+ uint res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapX_8203ad());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapX_8203ad());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8d1e799
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/879738.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapX_879738() {
+ float2 arg_0 = (1.0f).xx;
+ float2 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_879738()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_879738()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4d1a9fa
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/9bea80.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapX_9bea80() {
+ float arg_0 = 1.0f;
+ float res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_9bea80()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapX_9bea80()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..15cb00b
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/a4e103.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapX_a4e103() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapX_a4e103());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapX_a4e103());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..58c08c1
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/b1a5fe.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapX_b1a5fe() {
+ int2 arg_0 = (1).xx;
+ int2 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_b1a5fe()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapX_b1a5fe()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0ce9e57
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/bc2013.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapX_bc2013() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapX_bc2013());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapX_bc2013());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..124b900
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/bddb9f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapX_bddb9f() {
+ uint3 arg_0 = (1u).xxx;
+ uint3 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapX_bddb9f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapX_bddb9f());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8ed8370
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/d60cec.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapX_d60cec() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapX_d60cec());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapX_d60cec());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e157b31
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapX/edfa1f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapX_edfa1f() {
+ int4 arg_0 = (1).xxxx;
+ int4 res = QuadReadAcrossX(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_edfa1f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapX_edfa1f()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..65d370a
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/06a67c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 quadSwapY_06a67c() {
+ uint3 arg_0 = (1u).xxx;
+ uint3 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, quadSwapY_06a67c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, quadSwapY_06a67c());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d87d2dc
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/0c4938.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint quadSwapY_0c4938() {
+ uint arg_0 = 1u;
+ uint res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, quadSwapY_0c4938());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, quadSwapY_0c4938());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..63a844e
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/0d05a8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 quadSwapY_0d05a8() {
+ int2 arg_0 = (1).xx;
+ int2 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_0d05a8()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_0d05a8()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..730b966
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/14bb9a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 quadSwapY_14bb9a() {
+ int4 arg_0 = (1).xxxx;
+ int4 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_14bb9a()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_14bb9a()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..29e3150
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/1f1a06.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 quadSwapY_1f1a06() {
+ float2 arg_0 = (1.0f).xx;
+ float2 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_1f1a06()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(quadSwapY_1f1a06()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6fa8b32
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/264908.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> quadSwapY_264908() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapY_264908());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, quadSwapY_264908());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..aa690379
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/5b2e67.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> quadSwapY_5b2e67() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapY_5b2e67());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, quadSwapY_5b2e67());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..df1e46c
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/6f6bc9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float quadSwapY_6f6bc9() {
+ float arg_0 = 1.0f;
+ float res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_6f6bc9()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_6f6bc9()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ed4fc57
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/9277e9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t quadSwapY_9277e9() {
+ float16_t arg_0 = float16_t(1.0h);
+ float16_t res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapY_9277e9());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, quadSwapY_9277e9());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9302009
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/94ab6d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int quadSwapY_94ab6d() {
+ int arg_0 = 1;
+ int res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_94ab6d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(quadSwapY_94ab6d()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..16458ec
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/a27e1c.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 quadSwapY_a27e1c() {
+ uint2 arg_0 = (1u).xx;
+ uint2 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, quadSwapY_a27e1c());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, quadSwapY_a27e1c());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..800493f
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/a50fcb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> quadSwapY_a50fcb() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapY_a50fcb());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, quadSwapY_a50fcb());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c34dcbd
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/b9d9e7.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 quadSwapY_b9d9e7() {
+ float4 arg_0 = (1.0f).xxxx;
+ float4 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_b9d9e7()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(quadSwapY_b9d9e7()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..dfeef1b
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/bb697b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 quadSwapY_bb697b() {
+ uint4 arg_0 = (1u).xxxx;
+ uint4 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, quadSwapY_bb697b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, quadSwapY_bb697b());
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1422d57
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/be4e72.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 quadSwapY_be4e72() {
+ int3 arg_0 = (1).xxx;
+ int3 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_be4e72()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_be4e72()));
+}
+
diff --git a/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c16714e
--- /dev/null
+++ b/test/tint/builtins/gen/var/quadSwapY/d1ab4d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 quadSwapY_d1ab4d() {
+ float3 arg_0 = (1.0f).xxx;
+ float3 res = QuadReadAcrossY(arg_0);
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_d1ab4d()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(quadSwapY_d1ab4d()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupElect/3943d6.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/subgroupElect/3943d6.wgsl.expected.ir.msl
new file mode 100644
index 0000000..836711c
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupElect/3943d6.wgsl.expected.ir.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupElect_3943d6() {
+ bool res = simd_is_first();
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupElect_3943d6();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupElect_3943d6();
+}
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..deacce5
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/10eb45.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleDown_10eb45() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 4> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 4> res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleDown_10eb45());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleDown_10eb45());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..97ac3c7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/1b530f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleDown_1b530f() {
+ int3 arg_0 = (1).xxx;
+ uint arg_1 = 1u;
+ int3 v = arg_0;
+ uint v_1 = arg_1;
+ int3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_1b530f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_1b530f()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d6ee63b
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/257ff0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleDown_257ff0() {
+ float4 arg_0 = (1.0f).xxxx;
+ uint arg_1 = 1u;
+ float4 v = arg_0;
+ uint v_1 = arg_1;
+ float4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_257ff0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_257ff0()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b2a70a3
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/313d9b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleDown_313d9b() {
+ int4 arg_0 = (1).xxxx;
+ uint arg_1 = 1u;
+ int4 v = arg_0;
+ uint v_1 = arg_1;
+ int4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_313d9b()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleDown_313d9b()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..bbfb17b
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/57b1e8.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleDown_57b1e8() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ uint arg_1 = 1u;
+ vector<float16_t, 2> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 2> res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleDown_57b1e8());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleDown_57b1e8());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1861fe3
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/5d8b9f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleDown_5d8b9f() {
+ float3 arg_0 = (1.0f).xxx;
+ uint arg_1 = 1u;
+ float3 v = arg_0;
+ uint v_1 = arg_1;
+ float3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_5d8b9f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleDown_5d8b9f()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..583d3b7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/63fdb0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleDown_63fdb0() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 3> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 3> res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleDown_63fdb0());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleDown_63fdb0());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fb39605
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/642789.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleDown_642789() {
+ uint3 arg_0 = (1u).xxx;
+ uint arg_1 = 1u;
+ uint3 v = arg_0;
+ uint v_1 = arg_1;
+ uint3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleDown_642789());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleDown_642789());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6c7010a
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/7a0cf5.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleDown_7a0cf5() {
+ float2 arg_0 = (1.0f).xx;
+ uint arg_1 = 1u;
+ float2 v = arg_0;
+ uint v_1 = arg_1;
+ float2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_7a0cf5()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_7a0cf5()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c19db9d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/7f8886.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleDown_7f8886() {
+ float arg_0 = 1.0f;
+ uint arg_1 = 1u;
+ float v = arg_0;
+ uint v_1 = arg_1;
+ float res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_7f8886()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_7f8886()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4cc67ea
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/9c6714.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleDown_9c6714() {
+ float16_t arg_0 = float16_t(1.0h);
+ uint arg_1 = 1u;
+ float16_t v = arg_0;
+ uint v_1 = arg_1;
+ float16_t res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleDown_9c6714());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleDown_9c6714());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..0f84db7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/b41899.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleDown_b41899() {
+ int2 arg_0 = (1).xx;
+ uint arg_1 = 1u;
+ int2 v = arg_0;
+ uint v_1 = arg_1;
+ int2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_b41899()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleDown_b41899()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..bd502fe
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/c9f1c4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleDown_c9f1c4() {
+ uint2 arg_0 = (1u).xx;
+ uint arg_1 = 1u;
+ uint2 v = arg_0;
+ uint v_1 = arg_1;
+ uint2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleDown_c9f1c4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleDown_c9f1c4());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..079e8d1
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/d269eb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleDown_d269eb() {
+ int arg_0 = 1;
+ uint arg_1 = 1u;
+ int v = arg_0;
+ uint v_1 = arg_1;
+ int res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_d269eb()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleDown_d269eb()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e7178e7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/d46304.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleDown_d46304() {
+ uint4 arg_0 = (1u).xxxx;
+ uint arg_1 = 1u;
+ uint4 v = arg_0;
+ uint v_1 = arg_1;
+ uint4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleDown_d46304());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleDown_d46304());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9432a22
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleDown/d90c2f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleDown_d90c2f() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint v = arg_0;
+ uint v_1 = arg_1;
+ uint res = WaveReadLaneAt(v, (WaveGetLaneIndex() + v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleDown_d90c2f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleDown_d90c2f());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ba6cc92
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/0990cd.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleUp_0990cd() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 4> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 4> res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleUp_0990cd());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleUp_0990cd());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e0542aa
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/1bb93f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleUp_1bb93f() {
+ int arg_0 = 1;
+ uint arg_1 = 1u;
+ int v = arg_0;
+ uint v_1 = arg_1;
+ int res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_1bb93f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_1bb93f()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..5c606d4
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/23c7ca.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleUp_23c7ca() {
+ float arg_0 = 1.0f;
+ uint arg_1 = 1u;
+ float v = arg_0;
+ uint v_1 = arg_1;
+ float res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_23c7ca()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleUp_23c7ca()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6639655
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/3242a6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleUp_3242a6() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint v = arg_0;
+ uint v_1 = arg_1;
+ uint res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleUp_3242a6());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleUp_3242a6());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e5c3cea
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/33d495.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleUp_33d495() {
+ float4 arg_0 = (1.0f).xxxx;
+ uint arg_1 = 1u;
+ float4 v = arg_0;
+ uint v_1 = arg_1;
+ float4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_33d495()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_33d495()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c5b8190
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/3e609f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleUp_3e609f() {
+ int4 arg_0 = (1).xxxx;
+ uint arg_1 = 1u;
+ int4 v = arg_0;
+ uint v_1 = arg_1;
+ int4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_3e609f()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleUp_3e609f()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cf0f346
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/58de69.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleUp_58de69() {
+ uint2 arg_0 = (1u).xx;
+ uint arg_1 = 1u;
+ uint2 v = arg_0;
+ uint v_1 = arg_1;
+ uint2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleUp_58de69());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleUp_58de69());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..9a8707d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/868e52.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleUp_868e52() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 3> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 3> res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleUp_868e52());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleUp_868e52());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d411650
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/87c9d6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleUp_87c9d6() {
+ float3 arg_0 = (1.0f).xxx;
+ uint arg_1 = 1u;
+ float3 v = arg_0;
+ uint v_1 = arg_1;
+ float3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_87c9d6()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_87c9d6()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4daf379
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/88eb07.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleUp_88eb07() {
+ uint4 arg_0 = (1u).xxxx;
+ uint arg_1 = 1u;
+ uint4 v = arg_0;
+ uint v_1 = arg_1;
+ uint4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleUp_88eb07());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleUp_88eb07());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cd034dc
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/8a63f3.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleUp_8a63f3() {
+ int3 arg_0 = (1).xxx;
+ uint arg_1 = 1u;
+ int3 v = arg_0;
+ uint v_1 = arg_1;
+ int3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_8a63f3()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleUp_8a63f3()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..60aefa7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/a2075a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleUp_a2075a() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ uint arg_1 = 1u;
+ vector<float16_t, 2> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 2> res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleUp_a2075a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleUp_a2075a());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..de93727
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/abaea0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleUp_abaea0() {
+ uint3 arg_0 = (1u).xxx;
+ uint arg_1 = 1u;
+ uint3 v = arg_0;
+ uint v_1 = arg_1;
+ uint3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleUp_abaea0());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleUp_abaea0());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..87b55a7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/b58804.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleUp_b58804() {
+ float2 arg_0 = (1.0f).xx;
+ uint arg_1 = 1u;
+ float2 v = arg_0;
+ uint v_1 = arg_1;
+ float2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_b58804()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_b58804()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..79d2663
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/bbf7f4.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleUp_bbf7f4() {
+ float16_t arg_0 = float16_t(1.0h);
+ uint arg_1 = 1u;
+ float16_t v = arg_0;
+ uint v_1 = arg_1;
+ float16_t res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleUp_bbf7f4());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleUp_bbf7f4());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..ce1593e
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleUp/db5bcb.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleUp_db5bcb() {
+ int2 arg_0 = (1).xx;
+ uint arg_1 = 1u;
+ int2 v = arg_0;
+ uint v_1 = arg_1;
+ int2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() - v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_db5bcb()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleUp_db5bcb()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f7b3685
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/071aa0.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int2 subgroupShuffleXor_071aa0() {
+ int2 arg_0 = (1).xx;
+ uint arg_1 = 1u;
+ int2 v = arg_0;
+ uint v_1 = arg_1;
+ int2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_071aa0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_071aa0()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..cb7e26b
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/08f588.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint4 subgroupShuffleXor_08f588() {
+ uint4 arg_0 = (1u).xxxx;
+ uint arg_1 = 1u;
+ uint4 v = arg_0;
+ uint v_1 = arg_1;
+ uint4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, subgroupShuffleXor_08f588());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, subgroupShuffleXor_08f588());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..3189c03
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/1d36b6.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float subgroupShuffleXor_1d36b6() {
+ float arg_0 = 1.0f;
+ uint arg_1 = 1u;
+ float v = arg_0;
+ uint v_1 = arg_1;
+ float res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_1d36b6()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_1d36b6()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..309cb4d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/1e247f.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 2> subgroupShuffleXor_1e247f() {
+ vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ uint arg_1 = 1u;
+ vector<float16_t, 2> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 2> res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleXor_1e247f());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 2> >(0u, subgroupShuffleXor_1e247f());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..f9be9f7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/1f2590.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 3> subgroupShuffleXor_1f2590() {
+ vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 3> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 3> res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleXor_1f2590());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 3> >(0u, subgroupShuffleXor_1f2590());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..d148eff
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/2e033d.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+vector<float16_t, 4> subgroupShuffleXor_2e033d() {
+ vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ uint arg_1 = 1u;
+ vector<float16_t, 4> v = arg_0;
+ uint v_1 = arg_1;
+ vector<float16_t, 4> res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleXor_2e033d());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<vector<float16_t, 4> >(0u, subgroupShuffleXor_2e033d());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..fd11aac
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/445e83.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupShuffleXor_445e83() {
+ int arg_0 = 1;
+ uint arg_1 = 1u;
+ int v = arg_0;
+ uint v_1 = arg_1;
+ int res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_445e83()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupShuffleXor_445e83()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..bfb42f2
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/7435fe.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int3 subgroupShuffleXor_7435fe() {
+ int3 arg_0 = (1).xxx;
+ uint arg_1 = 1u;
+ int3 v = arg_0;
+ uint v_1 = arg_1;
+ int3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_7435fe()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_7435fe()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..c98ce5d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/80b6e9.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint subgroupShuffleXor_80b6e9() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint v = arg_0;
+ uint v_1 = arg_1;
+ uint res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, subgroupShuffleXor_80b6e9());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, subgroupShuffleXor_80b6e9());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..1adf2d4
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/9f945a.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint3 subgroupShuffleXor_9f945a() {
+ uint3 arg_0 = (1u).xxx;
+ uint arg_1 = 1u;
+ uint3 v = arg_0;
+ uint v_1 = arg_1;
+ uint3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, subgroupShuffleXor_9f945a());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, subgroupShuffleXor_9f945a());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6740408
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/bdddba.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int4 subgroupShuffleXor_bdddba() {
+ int4 arg_0 = (1).xxxx;
+ uint arg_1 = 1u;
+ int4 v = arg_0;
+ uint v_1 = arg_1;
+ int4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_bdddba()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_bdddba()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..8e52073
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/c88290.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float4 subgroupShuffleXor_c88290() {
+ float4 arg_0 = (1.0f).xxxx;
+ uint arg_1 = 1u;
+ float4 v = arg_0;
+ uint v_1 = arg_1;
+ float4 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_c88290()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store4(0u, asuint(subgroupShuffleXor_c88290()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..b3118ff
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/caa816.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float3 subgroupShuffleXor_caa816() {
+ float3 arg_0 = (1.0f).xxx;
+ uint arg_1 = 1u;
+ float3 v = arg_0;
+ uint v_1 = arg_1;
+ float3 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_caa816()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store3(0u, asuint(subgroupShuffleXor_caa816()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..64adc30
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/d224ab.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float16_t subgroupShuffleXor_d224ab() {
+ float16_t arg_0 = float16_t(1.0h);
+ uint arg_1 = 1u;
+ float16_t v = arg_0;
+ uint v_1 = arg_1;
+ float16_t res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleXor_d224ab());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store<float16_t>(0u, subgroupShuffleXor_d224ab());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..92a0c1a
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/e3c10b.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+uint2 subgroupShuffleXor_e3c10b() {
+ uint2 arg_0 = (1u).xx;
+ uint arg_1 = 1u;
+ uint2 v = arg_0;
+ uint v_1 = arg_1;
+ uint2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, subgroupShuffleXor_e3c10b());
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, subgroupShuffleXor_e3c10b());
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..6a8b230
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupShuffleXor/f7b453.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,22 @@
+SKIP: Wave ops not supported before SM 6.0
+
+
+RWByteAddressBuffer prevent_dce : register(u0);
+float2 subgroupShuffleXor_f7b453() {
+ float2 arg_0 = (1.0f).xx;
+ uint arg_1 = 1u;
+ float2 v = arg_0;
+ uint v_1 = arg_1;
+ float2 res = WaveReadLaneAt(v, (WaveGetLaneIndex() ^ v_1));
+ return res;
+}
+
+void fragment_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_f7b453()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store2(0u, asuint(subgroupShuffleXor_f7b453()));
+}
+
