This is an archive of the discontinued LLVM Phabricator instance.

Differential D115534

[mlir][spirv] Change default subgroup size
ClosedPublic

Authored by antiagainst on Dec 10 2021, 9:14 AM.

Details

Reviewers
scotttodd
mravishankar
Commits
rG1bfa40a5d6e1: [mlir][spirv] Change default subgroup size
Summary

This should really come from a matching target environment. But
as a default, it can be handy (to avoid always listing the full
resource limits attribute in IR, etc.). It's common to see 32
so use that as the subgroup size.

Depends On D115532

Diff Detail

Event Timeline

antiagainst created this revision.Dec 10 2021, 9:14 AM
Herald added a reviewer: mravishankar. · View Herald TranscriptDec 10 2021, 9:14 AM
Herald added subscribers: sdasgup3, wenzhicui, wrengr and 22 others. · View Herald Transcript
antiagainst requested review of this revision.Dec 10 2021, 9:14 AM
Herald added a project: Restricted Project. · View Herald TranscriptDec 10 2021, 9:14 AM
Herald added subscribers: stephenneuendorffer, nicolasvasilache. · View Herald Transcript
scotttodd accepted this revision.Dec 10 2021, 9:15 AM
This revision is now accepted and ready to land.Dec 10 2021, 9:15 AM
Harbormaster completed remote builds in B138690: Diff 393519.Dec 10 2021, 9:49 AM
antiagainst added a child revision: D115541: [mlir][spirv] Propagate LogicalResult in (de)serialization.Dec 10 2021, 10:58 AM
Closed by commit rG1bfa40a5d6e1: [mlir][spirv] Change default subgroup size (authored by antiagainst). · Explain WhyDec 10 2021, 4:21 PM
This revision was automatically updated to reflect the committed changes.
antiagainst added a commit: rG1bfa40a5d6e1: [mlir][spirv] Change default subgroup size.

Revision Contents

PathSize
mlir/
include/
mlir/
Dialect/
SPIRV/
IR/
3 lines

Diff 393626

mlir/include/mlir/Dialect/SPIRV/IR/TargetAndABI.td

Show First 20 Lines • Show All 72 Lines • ▼ Show 20 Linesdef SPV_ResourceLimitsAttr : StructAttr<"ResourceLimitsAttr", SPIRV_Dialect, [
// workgroup. // workgroup.
StructFieldAttr<"max_compute_workgroup_invocations", StructFieldAttr<"max_compute_workgroup_invocations",
DefaultValuedAttr<I32Attr, "128">>, DefaultValuedAttr<I32Attr, "128">>,
// The maximum size of a local compute workgroup, per dimension. // The maximum size of a local compute workgroup, per dimension.
StructFieldAttr<"max_compute_workgroup_size", StructFieldAttr<"max_compute_workgroup_size",
DefaultValuedAttr<I32ElementsAttr, "{128, 128, 64}">>, DefaultValuedAttr<I32ElementsAttr, "{128, 128, 64}">>,
// The default number of invocations in each subgroup. // The default number of invocations in each subgroup.
// 0x7FFFFFFF means unknown. StructFieldAttr<"subgroup_size", DefaultValuedAttr<I32Attr, "32">>,
StructFieldAttr<"subgroup_size", DefaultValuedAttr<I32Attr, "0x7FFFFFFF">>,
// The configurations of cooperative matrix operations // The configurations of cooperative matrix operations
// supported. Default is an empty list. // supported. Default is an empty list.
StructFieldAttr< StructFieldAttr<
"cooperative_matrix_properties_nv", "cooperative_matrix_properties_nv",
DefaultValuedAttr<SPV_CooperativeMatrixPropertiesNVArrayAttr, "{}">> DefaultValuedAttr<SPV_CooperativeMatrixPropertiesNVArrayAttr, "{}">>
]>; ]>;
#endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI #endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI