[tint] Add missing subgroup matrix expected files
Change-Id: I957b38db0b731fc27fde01c6f47122a9905bdaed
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/246215
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl
new file mode 100644
index 0000000..c8e40f5
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_03a4db() -> subgroup_matrix_result<i8, 8, 8> {
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_03a4db(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl
new file mode 100644
index 0000000..a6c114e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_0f7739() -> subgroup_matrix_left<i8, 8, 8> {
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_0f7739(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/122298.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/122298.wgsl.expected.wgsl
new file mode 100644
index 0000000..ffbb794
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/122298.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_122298() -> subgroup_matrix_left<i8, 8, 8> {
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_122298(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl
new file mode 100644
index 0000000..1838665
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_1b3162() -> subgroup_matrix_right<i8, 8, 8> {
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_1b3162(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/297099.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/297099.wgsl.expected.wgsl
new file mode 100644
index 0000000..1f62924
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/297099.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_297099() -> subgroup_matrix_result<i8, 8, 8> {
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_297099(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/29f533.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/29f533.wgsl.expected.wgsl
new file mode 100644
index 0000000..1dd590f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/29f533.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_29f533() -> subgroup_matrix_left<u8, 8, 8> {
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_29f533(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl
new file mode 100644
index 0000000..59248c3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_3f3203() -> subgroup_matrix_result<u8, 8, 8> {
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_3f3203(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/487a88.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/487a88.wgsl.expected.wgsl
new file mode 100644
index 0000000..308c4b5
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/487a88.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_487a88() -> subgroup_matrix_right<u8, 8, 8> {
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_487a88(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/50572b.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/50572b.wgsl.expected.wgsl
new file mode 100644
index 0000000..edb2716
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/50572b.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_50572b() -> subgroup_matrix_left<u8, 8, 8> {
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_50572b(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/567359.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/567359.wgsl.expected.wgsl
new file mode 100644
index 0000000..fecff31
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/567359.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_567359() -> subgroup_matrix_right<u8, 8, 8> {
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_567359(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl
new file mode 100644
index 0000000..dc53bf6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_58b2f2() -> subgroup_matrix_right<i8, 8, 8> {
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_58b2f2(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl
new file mode 100644
index 0000000..28c1bf7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_5e2602() -> subgroup_matrix_result<u8, 8, 8> {
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_5e2602(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/662272.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/662272.wgsl.expected.wgsl
new file mode 100644
index 0000000..dd69cdd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/662272.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_662272() -> subgroup_matrix_result<i8, 8, 8> {
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_662272(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl
new file mode 100644
index 0000000..d287b3d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_664e6e() -> subgroup_matrix_right<i8, 8, 8> {
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_664e6e(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl
new file mode 100644
index 0000000..a7580e4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_6778bb() -> subgroup_matrix_left<i8, 8, 8> {
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_6778bb(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl
new file mode 100644
index 0000000..343f35f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_86df2e() -> subgroup_matrix_right<u8, 8, 8> {
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_86df2e(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl
new file mode 100644
index 0000000..fc019d8
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_8861f0() -> subgroup_matrix_right<i8, 8, 8> {
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_8861f0(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl
new file mode 100644
index 0000000..e2aedca
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_91abc9() -> subgroup_matrix_left<i8, 8, 8> {
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_91abc9(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl
new file mode 100644
index 0000000..91b93b2
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_9bc799() -> subgroup_matrix_result<i8, 8, 8> {
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_9bc799(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl
new file mode 100644
index 0000000..b21508c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_9c4b12() -> subgroup_matrix_left<u8, 8, 8> {
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_9c4b12(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl
new file mode 100644
index 0000000..84aa29d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_ad687c() -> subgroup_matrix_left<u8, 8, 8> {
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_ad687c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/b19701.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/b19701.wgsl.expected.wgsl
new file mode 100644
index 0000000..14cb756
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/b19701.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_b19701() -> subgroup_matrix_right<i8, 8, 8> {
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_b19701(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl
new file mode 100644
index 0000000..4f55e15
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_b6f72c() -> subgroup_matrix_result<u8, 8, 8> {
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_b6f72c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl
new file mode 100644
index 0000000..62e34b1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_bcfcb6() -> subgroup_matrix_left<u8, 8, 8> {
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_bcfcb6(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl
new file mode 100644
index 0000000..313a8a1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_d72b92() -> subgroup_matrix_right<u8, 8, 8> {
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_d72b92(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/df0754.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/df0754.wgsl.expected.wgsl
new file mode 100644
index 0000000..47cf00d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/df0754.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_df0754() -> subgroup_matrix_left<i8, 8, 8> {
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_df0754(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/e40231.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/e40231.wgsl.expected.wgsl
new file mode 100644
index 0000000..fd04070
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/e40231.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_e40231() -> subgroup_matrix_result<i8, 8, 8> {
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_e40231(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl
new file mode 100644
index 0000000..fd35fe4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_e5381c() -> subgroup_matrix_result<u8, 8, 8> {
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_e5381c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl
new file mode 100644
index 0000000..16f19b4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_f364a9() -> subgroup_matrix_right<u8, 8, 8> {
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_rw.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_f364a9(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl
new file mode 100644
index 0000000..47d5199
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_f86f9f() -> subgroup_matrix_result<u8, 8, 8> {
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_ro.arg_0), 1u, true, 8u);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_f86f9f(), false, 64);
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/152780.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/152780.wgsl.expected.wgsl
new file mode 100644
index 0000000..e144480
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/152780.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_152780() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_result<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_152780();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/2025c4.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/2025c4.wgsl.expected.wgsl
new file mode 100644
index 0000000..59f1916
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/2025c4.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_2025c4() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_left<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_2025c4();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl
new file mode 100644
index 0000000..27e76c3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_43b7fd() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_right<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_43b7fd();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/443b6c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/443b6c.wgsl.expected.wgsl
new file mode 100644
index 0000000..36e6725
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/443b6c.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_443b6c() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_right<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_443b6c();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/60099d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/60099d.wgsl.expected.wgsl
new file mode 100644
index 0000000..ed0a02e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/60099d.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_60099d() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_result<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_60099d();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/85667d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/85667d.wgsl.expected.wgsl
new file mode 100644
index 0000000..e9cc969
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/85667d.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_85667d() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_right<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_85667d();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/8a88da.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/8a88da.wgsl.expected.wgsl
new file mode 100644
index 0000000..394c9df
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/8a88da.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_8a88da() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_result<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_8a88da();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/9b9975.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/9b9975.wgsl.expected.wgsl
new file mode 100644
index 0000000..4d2a83d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/9b9975.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_9b9975() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_left<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_9b9975();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl
new file mode 100644
index 0000000..b8a1fea
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_a50fc4() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_right<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_a50fc4();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl
new file mode 100644
index 0000000..153282c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_c18bf3() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_right<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c18bf3();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/c266d0.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/c266d0.wgsl.expected.wgsl
new file mode 100644
index 0000000..98f68a6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/c266d0.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_c266d0() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_left<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c266d0();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl
new file mode 100644
index 0000000..893de4a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_c9f1a5() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_result<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c9f1a5();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl
new file mode 100644
index 0000000..d9fac6a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_ca6fe7() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_left<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_ca6fe7();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl
new file mode 100644
index 0000000..243e85a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_e2ecd9() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_right<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e2ecd9();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/e4d013.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/e4d013.wgsl.expected.wgsl
new file mode 100644
index 0000000..25a25fa
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/e4d013.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_e4d013() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_result<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e4d013();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/e59232.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/e59232.wgsl.expected.wgsl
new file mode 100644
index 0000000..b43a935
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/e59232.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_e59232() {
+ subgroupMatrixStore(&(arg_0), 1u, subgroup_matrix_result<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e59232();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/f79102.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/f79102.wgsl.expected.wgsl
new file mode 100644
index 0000000..9907851
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/f79102.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_f79102() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_left<i8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_f79102();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl
new file mode 100644
index 0000000..8159895
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_f96cc1() {
+ subgroupMatrixStore(&(sb_rw.arg_0), 1u, subgroup_matrix_left<u8, 8, 8>(), true, 8u);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_f96cc1();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl
new file mode 100644
index 0000000..6893ebf
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/03a4db.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_03a4db() -> subgroup_matrix_result<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_03a4db(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl
new file mode 100644
index 0000000..4172ad3
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/0f7739.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_0f7739() -> subgroup_matrix_left<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_0f7739(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/122298.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/122298.wgsl.expected.wgsl
new file mode 100644
index 0000000..1386864
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/122298.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_122298() -> subgroup_matrix_left<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_122298(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl
new file mode 100644
index 0000000..e0a6d2d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/1b3162.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_1b3162() -> subgroup_matrix_right<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_1b3162(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/297099.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/297099.wgsl.expected.wgsl
new file mode 100644
index 0000000..2146853
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/297099.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_297099() -> subgroup_matrix_result<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_297099(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/29f533.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/29f533.wgsl.expected.wgsl
new file mode 100644
index 0000000..4995bdf
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/29f533.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_29f533() -> subgroup_matrix_left<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_29f533(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl
new file mode 100644
index 0000000..943ca9e
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/3f3203.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_3f3203() -> subgroup_matrix_result<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_3f3203(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/487a88.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/487a88.wgsl.expected.wgsl
new file mode 100644
index 0000000..5510e08
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/487a88.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_487a88() -> subgroup_matrix_right<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_487a88(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/50572b.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/50572b.wgsl.expected.wgsl
new file mode 100644
index 0000000..5406199
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/50572b.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_50572b() -> subgroup_matrix_left<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_50572b(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/567359.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/567359.wgsl.expected.wgsl
new file mode 100644
index 0000000..89d48bb
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/567359.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_567359() -> subgroup_matrix_right<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_567359(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl
new file mode 100644
index 0000000..24cec26
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/58b2f2.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_58b2f2() -> subgroup_matrix_right<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_58b2f2(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl
new file mode 100644
index 0000000..67e165f
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/5e2602.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_5e2602() -> subgroup_matrix_result<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_5e2602(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/662272.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/662272.wgsl.expected.wgsl
new file mode 100644
index 0000000..d62b709
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/662272.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_662272() -> subgroup_matrix_result<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_662272(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl
new file mode 100644
index 0000000..cce6eb5
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/664e6e.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_664e6e() -> subgroup_matrix_right<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_664e6e(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl
new file mode 100644
index 0000000..137b500
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/6778bb.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_6778bb() -> subgroup_matrix_left<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_6778bb(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl
new file mode 100644
index 0000000..a371cc5
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/86df2e.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_86df2e() -> subgroup_matrix_right<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_86df2e(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl
new file mode 100644
index 0000000..05797de
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/8861f0.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_8861f0() -> subgroup_matrix_right<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_8861f0(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl
new file mode 100644
index 0000000..05f900e
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/91abc9.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_91abc9() -> subgroup_matrix_left<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_91abc9(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl
new file mode 100644
index 0000000..b56fe22
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/9bc799.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_9bc799() -> subgroup_matrix_result<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_9bc799(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl
new file mode 100644
index 0000000..cf07ab1
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/9c4b12.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_9c4b12() -> subgroup_matrix_left<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_9c4b12(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl
new file mode 100644
index 0000000..3114628
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/ad687c.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_ad687c() -> subgroup_matrix_left<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_ad687c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/b19701.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/b19701.wgsl.expected.wgsl
new file mode 100644
index 0000000..e3af3b7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/b19701.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_b19701() -> subgroup_matrix_right<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<i8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_b19701(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl
new file mode 100644
index 0000000..8ac0c5d
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/b6f72c.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixLoad_b6f72c() -> subgroup_matrix_result<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_b6f72c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl
new file mode 100644
index 0000000..5fd7324
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/bcfcb6.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_bcfcb6() -> subgroup_matrix_left<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_bcfcb6(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl
new file mode 100644
index 0000000..50bccf8
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/d72b92.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_d72b92() -> subgroup_matrix_right<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_d72b92(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/df0754.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/df0754.wgsl.expected.wgsl
new file mode 100644
index 0000000..b295c77
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/df0754.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixLoad_df0754() -> subgroup_matrix_left<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_left<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<i8, 8, 8>>(&(arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_df0754(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/e40231.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/e40231.wgsl.expected.wgsl
new file mode 100644
index 0000000..a9f9f51
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/e40231.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<i32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_e40231() -> subgroup_matrix_result<i8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<i8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<i8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_e40231(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl
new file mode 100644
index 0000000..d42de4a
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/e5381c.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_e5381c() -> subgroup_matrix_result<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_e5381c(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl
new file mode 100644
index 0000000..95d94d0
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/f364a9.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixLoad_f364a9() -> subgroup_matrix_right<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_right<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<u8, 8, 8>>(&(sb_rw.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_f364a9(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl
new file mode 100644
index 0000000..b4f2383
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixLoad/f86f9f.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+enable chromium_experimental_subgroup_matrix;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : array<u32, 1024>;
+
+struct SB_RO {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
+
+fn subgroupMatrixLoad_f86f9f() -> subgroup_matrix_result<u8, 8, 8> {
+ var arg_1 = 1u;
+ const arg_2 = true;
+ var arg_3 = 8u;
+ var res : subgroup_matrix_result<u8, 8, 8> = subgroupMatrixLoad<subgroup_matrix_result<u8, 8, 8>>(&(sb_ro.arg_0), arg_1, arg_2, arg_3);
+ return res;
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore(&(prevent_dce), 0, subgroupMatrixLoad_f86f9f(), false, 64);
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/152780.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/152780.wgsl.expected.wgsl
new file mode 100644
index 0000000..7b638af
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/152780.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_152780() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_152780();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/2025c4.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/2025c4.wgsl.expected.wgsl
new file mode 100644
index 0000000..fc1a8db
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/2025c4.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_2025c4() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_2025c4();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl
new file mode 100644
index 0000000..e1c5c06
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/43b7fd.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_43b7fd() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_43b7fd();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/443b6c.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/443b6c.wgsl.expected.wgsl
new file mode 100644
index 0000000..90557d7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/443b6c.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_443b6c() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_443b6c();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/60099d.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/60099d.wgsl.expected.wgsl
new file mode 100644
index 0000000..2d17298
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/60099d.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_60099d() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_60099d();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/85667d.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/85667d.wgsl.expected.wgsl
new file mode 100644
index 0000000..295c176
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/85667d.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_85667d() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_85667d();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/8a88da.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/8a88da.wgsl.expected.wgsl
new file mode 100644
index 0000000..2585877
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/8a88da.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_8a88da() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_8a88da();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/9b9975.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/9b9975.wgsl.expected.wgsl
new file mode 100644
index 0000000..3975fe5
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/9b9975.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_9b9975() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_9b9975();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl
new file mode 100644
index 0000000..db3ffbc
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/a50fc4.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_a50fc4() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_a50fc4();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl
new file mode 100644
index 0000000..e80c559
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/c18bf3.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_c18bf3() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c18bf3();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/c266d0.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/c266d0.wgsl.expected.wgsl
new file mode 100644
index 0000000..a4328f9
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/c266d0.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_c266d0() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c266d0();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl
new file mode 100644
index 0000000..c488389
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/c9f1a5.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_c9f1a5() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_c9f1a5();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl
new file mode 100644
index 0000000..95f5277
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/ca6fe7.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<u32, 1024>;
+
+fn subgroupMatrixStore_ca6fe7() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_ca6fe7();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl
new file mode 100644
index 0000000..58ad502
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/e2ecd9.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_e2ecd9() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_right<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e2ecd9();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/e4d013.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/e4d013.wgsl.expected.wgsl
new file mode 100644
index 0000000..dd2e632
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/e4d013.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_e4d013() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e4d013();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/e59232.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/e59232.wgsl.expected.wgsl
new file mode 100644
index 0000000..4de49df
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/e59232.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+enable chromium_experimental_subgroup_matrix;
+
+var<workgroup> arg_0 : array<i32, 1024>;
+
+fn subgroupMatrixStore_e59232() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_result<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_e59232();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/f79102.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/f79102.wgsl.expected.wgsl
new file mode 100644
index 0000000..c9e562f
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/f79102.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<i32, 1024>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_f79102() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<i8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_f79102();
+}
diff --git a/test/tint/builtins/gen/var/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl
new file mode 100644
index 0000000..4aadf24
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupMatrixStore/f96cc1.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+enable chromium_experimental_subgroup_matrix;
+
+struct SB_RW {
+ arg_0 : array<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn subgroupMatrixStore_f96cc1() {
+ var arg_1 = 1u;
+ var arg_2 = subgroup_matrix_left<u8, 8, 8>();
+ const arg_3 = true;
+ var arg_4 = 8u;
+ subgroupMatrixStore(&(sb_rw.arg_0), arg_1, arg_2, arg_3, arg_4);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupMatrixStore_f96cc1();
+}
