Skip to content

Commit 5c136fd

Browse files
kvarkkvark
committed
Disallow binary mul for coop matrix, reduce test footprint
1 parent ffa39e2 commit 5c136fd

9 files changed

Lines changed: 69 additions & 113 deletions

File tree

‎naga/src/back/msl/keywords.rs‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -359,6 +359,8 @@ const RESERVED: &[&str] = &[
359359
super::writer::IMAGE_SIZE_EXTERNAL_FUNCTION,
360360
super::writer::ARGUMENT_BUFFER_WRAPPER_STRUCT,
361361
super::writer::EXTERNAL_TEXTURE_WRAPPER_STRUCT,
362+
super::writer::COOPERATIVE_LOAD_FUNCTION,
363+
super::writer::COOPERATIVE_MULTIPLY_ADD_FUNCTION,
362364
];
363365

364366
// The set of concrete integer dot product function variants.

‎naga/src/front/wgsl/lower/mod.rs‎

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -910,15 +910,6 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
910910
fn ensure_type_exists(&mut self, inner: ir::TypeInner) -> Handle<ir::Type> {
911911
self.as_global().ensure_type_exists(None, inner)
912912
}
913-
914-
fn _get_runtime_expression(&self, expr: Handle<ir::Expression>) -> &ir::Expression {
915-
match self.expr_type {
916-
ExpressionContextType::Runtime(ref ctx) => &ctx.function.expressions[expr],
917-
ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
918-
unreachable!()
919-
}
920-
}
921-
}
922913
}
923914

924915
struct ArgumentContext<'ctx, 'source> {

‎naga/src/front/wgsl/parse/mod.rs‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1475,7 +1475,7 @@ impl Parser {
14751475
Ok((ty, span))
14761476
}
14771477

1478-
/// Parses `<T,R>`, returning (T, span of T, R, span of R)
1478+
/// Parses `<T,R>`, returning (T, span of T, R)
14791479
fn cooperative_scalar_and_role<'a>(
14801480
&mut self,
14811481
lexer: &mut Lexer<'a>,

‎naga/src/valid/expression.rs‎

Lines changed: 0 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -872,21 +872,6 @@ impl super::Validator {
872872
(&Ti::Matrix { columns, .. }, &Ti::Matrix { rows, .. }) => {
873873
columns == rows
874874
}
875-
// Coop matrix * coop matrix.
876-
(
877-
&Ti::CooperativeMatrix {
878-
columns,
879-
scalar: scalar1,
880-
role: role1,
881-
..
882-
},
883-
&Ti::CooperativeMatrix {
884-
rows,
885-
scalar: scalar2,
886-
role: role2,
887-
..
888-
},
889-
) => columns == rows && scalar1 == scalar2 && role1 == role2,
890875
// Scalar * coop matrix.
891876
(&Ti::Scalar(s1), &Ti::CooperativeMatrix { scalar: s2, .. })
892877
| (&Ti::CooperativeMatrix { scalar: s1, .. }, &Ti::Scalar(s2)) => {

‎naga/tests/in/wgsl/cooperative-matrix.toml‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,8 @@
1-
targets = "IR | SPIRV | METAL | WGSL"
1+
targets = "SPIRV | METAL | WGSL"
22
god_mode = true
33

44
[spv]
5-
debug = true
5+
debug = false
66
version = [1, 4]
77

88
[msl]

‎naga/tests/in/wgsl/cooperative-matrix.wgsl‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,17 @@
1+
// type declarations with different roles
12
var<private> a: coop_mat8x8<f32, A>;
23
var<private> b: coop_mat8x8<f32, B>;
34
@group(0) @binding(0)
45
var<storage, read_write> ext: array<f32>;
56

67
@compute @workgroup_size(8, 8, 1)
78
fn main() {
9+
// loading from memory
810
var c = coopLoad<coop_mat8x8<f32, C>>(&ext[4]);
11+
// actual multiply-add
912
var d = coopMultiplyAdd(a, b, c);
13+
// storing into memory
1014
coopStore(d, &ext[0]);
15+
// operations on the type
1116
c = d;
1217
}

‎naga/tests/out/ir/wgsl-cooperative-matrix.compact.ron‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -219,6 +219,8 @@
219219
],
220220
diagnostic_filter_leaf: None,
221221
),
222+
mesh_info: None,
223+
task_payload: None,
222224
),
223225
],
224226
diagnostic_filters: [],

‎naga/tests/out/ir/wgsl-cooperative-matrix.ron‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -219,6 +219,8 @@
219219
],
220220
diagnostic_filter_leaf: None,
221221
),
222+
mesh_info: None,
223+
task_payload: None,
222224
),
223225
],
224226
diagnostic_filters: [],

‎naga/tests/out/spv/wgsl-cooperative-matrix.spvasm‎

Lines changed: 55 additions & 86 deletions
Original file line numberDiff line numberDiff line change
@@ -1,99 +1,68 @@
11
; SPIR-V
22
; Version: 1.4
33
; Generator: rspirv
4-
; Bound: 46
4+
; Bound: 45
55
OpCapability Shader
66
OpCapability CooperativeMatrixKHR
77
OpCapability VulkanMemoryModel
88
OpExtension "SPV_KHR_cooperative_matrix"
99
OpExtension "SPV_KHR_vulkan_memory_model"
1010
%1 = OpExtInstImport "GLSL.std.450"
1111
OpMemoryModel Logical Vulkan
12-
OpEntryPoint GLCompute %25 "main" %15 %18 %21
13-
OpExecutionMode %25 LocalSize 8 8 1
14-
%3 = OpString "cooperative-matrix.wgsl"
15-
OpSource Unknown 0 %3 "var<private> a: coop_mat8x8<f32, A>;
16-
var<private> b: coop_mat8x8<f32, B>;
17-
@group(0) @binding(0)
18-
var<storage, read_write> ext: array<f32>;
19-
20-
@compute @workgroup_size(8, 8, 1)
21-
fn main() {
22-
var c = coopLoad<coop_mat8x8<f32, C>>(&ext[4]);
23-
var d = coopMultiplyAdd(a, b, c);
24-
coopStore(d, &ext[0]);
25-
c = d;
26-
}
27-
"
28-
OpName %15 "a"
29-
OpName %18 "b"
30-
OpName %21 "ext"
31-
OpName %25 "main"
32-
OpName %29 "c"
33-
OpName %32 "d"
34-
OpDecorate %12 ArrayStride 4
35-
OpDecorate %21 DescriptorSet 0
36-
OpDecorate %21 Binding 0
37-
OpDecorate %22 Block
38-
OpMemberDecorate %22 0 Offset 0
12+
OpEntryPoint GLCompute %24 "main" %14 %17 %20
13+
OpExecutionMode %24 LocalSize 8 8 1
14+
OpDecorate %11 ArrayStride 4
15+
OpDecorate %20 DescriptorSet 0
16+
OpDecorate %20 Binding 0
17+
OpDecorate %21 Block
18+
OpMemberDecorate %21 0 Offset 0
3919
%2 = OpTypeVoid
40-
%4 = OpTypeFloat 32
41-
%7 = OpTypeInt 32 0
42-
%6 = OpConstant %7 3
43-
%8 = OpConstant %7 8
44-
%9 = OpConstant %7 0
45-
%5 = OpTypeCooperativeMatrixKHR %4 %6 %8 %8 %9
46-
%11 = OpConstant %7 1
47-
%10 = OpTypeCooperativeMatrixKHR %4 %6 %8 %8 %11
48-
%12 = OpTypeRuntimeArray %4
49-
%14 = OpConstant %7 2
50-
%13 = OpTypeCooperativeMatrixKHR %4 %6 %8 %8 %14
51-
%16 = OpTypePointer Private %5
52-
%17 = OpConstantNull %5
53-
%15 = OpVariable %16 Private %17
54-
%19 = OpTypePointer Private %10
55-
%20 = OpConstantNull %10
56-
%18 = OpVariable %19 Private %20
57-
%22 = OpTypeStruct %12
58-
%23 = OpTypePointer StorageBuffer %22
59-
%21 = OpVariable %23 StorageBuffer
60-
%26 = OpTypeFunction %2
61-
%27 = OpTypePointer StorageBuffer %12
62-
%30 = OpTypePointer Function %13
63-
%31 = OpConstantNull %13
64-
%33 = OpConstantNull %13
65-
%35 = OpTypePointer StorageBuffer %4
66-
%36 = OpConstant %7 4
67-
%25 = OpFunction %2 None %26
68-
%24 = OpLabel
69-
%29 = OpVariable %30 Function %31
70-
%32 = OpVariable %30 Function %33
71-
%28 = OpAccessChain %27 %21 %9
72-
OpBranch %34
73-
%34 = OpLabel
74-
OpLine %3 8 44
75-
OpLine %3 8 13
76-
%37 = OpAccessChain %35 %28 %36
77-
%38 = OpCooperativeMatrixLoadKHR %13 %37 %11 %8
78-
OpLine %3 8 5
79-
OpStore %29 %38
80-
OpLine %3 9 29
81-
%39 = OpLoad %5 %15
82-
OpLine %3 9 13
83-
%40 = OpLoad %10 %18
84-
%41 = OpLoad %13 %29
85-
%42 = OpCooperativeMatrixMulAddKHR %13 %39 %40 %41
86-
OpLine %3 9 5
87-
OpStore %32 %42
88-
OpLine %3 1 1
89-
%43 = OpLoad %13 %32
90-
OpLine %3 10 19
91-
OpLine %3 10 5
92-
%44 = OpAccessChain %35 %28 %9
93-
OpCooperativeMatrixStoreKHR %44 %43 %11 %8
94-
OpLine %3 1 1
95-
%45 = OpLoad %13 %32
96-
OpLine %3 11 5
97-
OpStore %29 %45
20+
%3 = OpTypeFloat 32
21+
%6 = OpTypeInt 32 0
22+
%5 = OpConstant %6 3
23+
%7 = OpConstant %6 8
24+
%8 = OpConstant %6 0
25+
%4 = OpTypeCooperativeMatrixKHR %3 %5 %7 %7 %8
26+
%10 = OpConstant %6 1
27+
%9 = OpTypeCooperativeMatrixKHR %3 %5 %7 %7 %10
28+
%11 = OpTypeRuntimeArray %3
29+
%13 = OpConstant %6 2
30+
%12 = OpTypeCooperativeMatrixKHR %3 %5 %7 %7 %13
31+
%15 = OpTypePointer Private %4
32+
%16 = OpConstantNull %4
33+
%14 = OpVariable %15 Private %16
34+
%18 = OpTypePointer Private %9
35+
%19 = OpConstantNull %9
36+
%17 = OpVariable %18 Private %19
37+
%21 = OpTypeStruct %11
38+
%22 = OpTypePointer StorageBuffer %21
39+
%20 = OpVariable %22 StorageBuffer
40+
%25 = OpTypeFunction %2
41+
%26 = OpTypePointer StorageBuffer %11
42+
%29 = OpTypePointer Function %12
43+
%30 = OpConstantNull %12
44+
%32 = OpConstantNull %12
45+
%34 = OpTypePointer StorageBuffer %3
46+
%35 = OpConstant %6 4
47+
%24 = OpFunction %2 None %25
48+
%23 = OpLabel
49+
%28 = OpVariable %29 Function %30
50+
%31 = OpVariable %29 Function %32
51+
%27 = OpAccessChain %26 %20 %8
52+
OpBranch %33
53+
%33 = OpLabel
54+
%36 = OpAccessChain %34 %27 %35
55+
%37 = OpCooperativeMatrixLoadKHR %12 %36 %10 %7
56+
OpStore %28 %37
57+
%38 = OpLoad %4 %14
58+
%39 = OpLoad %9 %17
59+
%40 = OpLoad %12 %28
60+
%41 = OpCooperativeMatrixMulAddKHR %12 %38 %39 %40
61+
OpStore %31 %41
62+
%42 = OpLoad %12 %31
63+
%43 = OpAccessChain %34 %27 %8
64+
OpCooperativeMatrixStoreKHR %43 %42 %10 %7
65+
%44 = OpLoad %12 %31
66+
OpStore %28 %44
9867
OpReturn
9968
OpFunctionEnd

0 commit comments

Comments
 (0)