הדיאלקט Shardy (SDY) מגדיר ייצוג של חלוקת טינסורים לפי ציר ורכיבי API נוספים כדי לצרף חלוקות לטינסורים.
תפעול
sdy.all_gather
(sdy::AllGatherOp)
ביצוע תקשורת all-gather לאורך צירים
תחביר:
operation ::= `sdy.all_gather` $gathering_axes $tensor `out_sharding````=```$out_sharding attr-dict `:` type($result)
איסוף קטעי טינסור לאורך צירים שצוינו ב-gathering_axes
.
gathering_axes
היא רשימה של רשימות של צירים. הרשימה החיצונית חורגת מהמאפיינים של הטנזור. כל רשימה פנימית מציינת את הצירים שבהם צריך לבצע איסוף נפרד של המאפיין הרלוונטי. הוא יחול על חלוקת המשנה של המשתנה (tensor
) כדי לקבל את חלוקת המשנה של התוצאה (out_sharding
).
שימו לב ש-out_sharding
לא משמש לקביעת חלוקת התוצאות למחיצות. במקום זאת, חלוקת התוצאה למחיצות נקבעת לפי חלוקת המשתנה המבצע וה-gathering_axes
, ו-out_sharding
חייב להתאים לחלוקה המשוערת הזו.
דוגמה:
%1 = stablehlo.tanh(%0) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a", "b", "c"}, {}, {"d"}\]>]>} : tensor<8x8x8xf32>
%2 = sdy.all_gather [{"b", "c"}, {}, {"d"}\] %1 out_sharding=<@mesh, [{"a"}, {}, {}\]> : tensor<8x8x8xf32>
אילוצים:
- חייב לעמוד באילוצים שמפורטים ב-
Sdy_CollectiveOpInterface
. - הרכיבים ב-
gathering_axes
חייבים לעמוד באילוצים שמפורטים ב-AxisRefListAttr
. - החלת
gathering_axes
על חלוקת המשתנה לחלקים מקבלת את הערךout_sharding
.
מאפיינים: SameOperandsAndResultType
ממשקים: InferTypeOpInterface
, Sdy_CollectiveOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
gathering_axes | ::mlir::sdy::ListOfAxisRefListsAttr | רשימה של רשימות של ערכי עזר לצירים |
out_sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensor |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.all_reduce
(sdy::AllReduceOp)
ביצוע תקשורת all-reduce לאורך צירים
תחביר:
operation ::= `sdy.all_reduce` $reduction_axes $tensor `out_sharding````=```$out_sharding attr-dict `:` type($result)
הפחתת קטעי טינסור לאורך צירים שצוינו ב-reduction_axes
.
הסדר של reduction_axes
לא חשוב לתוצאה, אבל הוא יכול להשפיע על הסדר של קבוצות הרפליקה התואמות.
אילוצים:
- חייב לעמוד באילוצים שמפורטים ב-
Sdy_CollectiveOpInterface
. reduction_axes
חייב לעמוד באילוצים שמפורטים ב-AxisRefListAttr
.reduction_axes
לא יכול לחפוף לצירים של חלוקת המשנה של המשתנה.
מאפיינים: SameOperandsAndResultType
ממשקים: CollectiveOpInterface
, InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
reduction_axes | ::mlir::sdy::AxisRefListAttr | רשימת הפניות לצירים |
out_sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensor |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.all_slice
(sdy::AllSliceOp)
ביצוע פעולת חיתוך דינמי לאורך צירים
תחביר:
operation ::= `sdy.all_slice` $slicing_axes $tensor `out_sharding````=```$out_sharding attr-dict `:` type($result)
חיתוך קטעים של טינסור לאורך צירים שצוינו ב-slicing_axes
. יש דואליות אלגברית בין sdy.all_slice
לבין sdy.all_gather
.
slicing_axes
היא רשימה של רשימות של צירים. הרשימה החיצונית חורגת מהמאפיינים של הטנזור. כל רשימה פנימית מציינת את הצירים שבהם צריך לבצע חיתוך של המאפיין הרלוונטי. הוא יחול על חלוקת המשנה של המשתנה (tensor
) כדי לקבל את חלוקת המשנה של התוצאה (out_sharding
).
שימו לב ש-out_sharding
לא משמש לקביעת חלוקת התוצאות למחיצות. במקום זאת, חלוקת התוצאה למחיצות נקבעת לפי חלוקת המשתנה המבצע וה-slicing_axes
, ו-out_sharding
חייב להתאים לחלוקה המשוערת הזו.
דוגמה:
%1 = stablehlo.tanh(%0) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a"}, {}, {}\]>]>} : tensor<8x8x8xf32>
%2 = sdy.all_slice [{"b", "c"}, {}, {"d"}\] %1 out_sharding=<@mesh, [{"a", "b", "c"}, {}, {"d"}\]> : tensor<8x8x8xf32>
אילוצים:
- הרכיבים ב-
slicing_axes
חייבים לעמוד באילוצים שמפורטים ב-AxisRefListAttr
. - חייב לעמוד באילוצים שמפורטים ב-
Sdy_CollectiveOpInterface
. - החלת
slicing_axes
על חלוקת המשתנה לחלקים מקבלת את הערךout_sharding
.
מאפיינים: SameOperandsAndResultType
ממשקים: CollectiveOpInterface
, InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
slicing_axes | ::mlir::sdy::ListOfAxisRefListsAttr | רשימה של רשימות של ערכי עזר לצירים |
out_sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensor |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.all_to_all
(sdy::AllToAllOp)
ביצוע תקשורת 'הכול לכולם' לאורך צירים
תחביר:
operation ::= `sdy.all_to_all` $params $tensor `out_sharding````=```$out_sharding attr-dict `:` type($result)
לכל קבוצה של ערכי טור (tuple) מסוג (axes, src_dim, tgt_dim) ברשימת הפרמטרים, הפעולה הזו חותכת קטעי טינסור לאורך המאפיין tgt_dim
והצירים שצוינו ב-axes
, מפזרת את הקטעים האלה לאורך הצירים ומארגנת אותם ברצף לאורך המאפיין src_dim
.
הפעולה הזו היא למעשה שילוב של all-gather לאורך src_dim
ו-axes
, ואחריו all-slice לאורך tgt_dim
ו-axes
. כלומר, סיומת של המאפיין src_dim
של חלוקת הצירים בטנסור הקלט מצורפת למאפיין tgt_dim
של חלוקת הצירים בטנסור הפלט.
הפעולה 'הכול אל הכול' תחול על חלוקת המשנה של המשתנה (tensor
) כדי לקבל את חלוקת המשנה של התוצאה (out_sharding
).
שימו לב ש-out_sharding
לא משמש לקביעת חלוקת התוצאות למחיצות. במקום זאת, חלוקת התוצאות למחיצות נקבעת לפי חלוקת המשתנה src_dim
, tgt_dim
ו-axes
למחיצות, ו-out_sharding
חייב להתאים לחלוקה המשוערת הזו למחיצות.
דוגמה:
%1 = stablehlo.tanh(%0) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a", "b"}, {"c"}, {}, {}\]>]>} : tensor<8x8x4x4x32>
%2 = sdy.all_to_all [{"b"}: 0->2, {"c"}: 1->3] %1 out_sharding=<@mesh, [{"a"}, {}, {"b"}, {"c"}\]> : tensor<8x8x4x4x32>
אילוצים:
- חייב לעמוד באילוצים שמפורטים ב-
Sdy_CollectiveOpInterface
. - רשימת הפרמטרים לא יכולה להיות ריקה.
- לכל פרמטר ב-
params
:- הרכיבים ב-
axes
חייבים לעמוד באילוצים שלAxisRefAttr
. - הערכים של
src_dim
ו-tgt_dim
חייבים להיות מידות תקינות (לא שליליות וקטנות מרמת הטנזור). - כל
src_dim
אוtgt_dim
חייבים להיות ייחודיים בכל הפרמטרים. - העמודה
src_dim
חייבת להיות ממוינת בסדר עולה בכל הפרמטרים.
- הרכיבים ב-
- העברת
axes
מ-src_dim
ל-tgt_dim
בחלוקת המזהה של הפעולה מקבלת את הערךout_sharding
.
מאפיינים: SameOperandsAndResultType
ממשקים: InferTypeOpInterface
, Sdy_CollectiveOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
params | ::mlir::sdy::AlltoAllParamListAttr | רשימה של פרמטרים של כל-אל-כל |
out_sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensor |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.collective_permute
(sdy::CollectivePermuteOp)
ביצוע תקשורת של תמורה קבוצתית כדי להחליף צירים
תחביר:
operation ::= `sdy.collective_permute` $tensor `out_sharding````=```$out_sharding attr-dict `:` type($result)
שליחת מקטע של הטנסור הקלט מכל מכשיר למכשיר אחר כדי לשנות את הסדר של הצירים שמחלקים את הטנסור או להחליף אותם.
טרנספורמציה של חלוקה קולקטיבית יכולה לשנות את חלוקת המשנה של הקלט כך שכל מאפיין חייב להיות מחולק כמו קודם, כלומר הוא חייב להיות מחולק לפי צירים שמכפלת הגדלים שלהם תואמת לזו של הצירים שבהם התאמת הטנזור בעבר.
אפשר להשתמש באפשרות הזו כדי לשנות את הסדר של צירים במאפיין יחיד או במאפיינים שונים, ולהחליף צירים מחולקים לשברים בצירים עם רפליקה.
בדוגמה הבאה, הגודל של הטנסור המפוצל הוא tensor<1x4x2xf32>
, והוא נשמר על ידי המעבר הארגוני.
דוגמה:
sdy.mesh @mesh = <["a"=2, "b"=2, "c"=4, "d"=2, "e"=2, "f"=2]>
%1 = stablehlo.tanh(%0) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a", "c"}, {"f"}, {"d", "e"}\]>]>} : tensor<8x8x8xf32>
%2 = sdy.collective_permute %1 out_sharding=<@mesh, [{"c":(1)2, "b", "f"}, {"a"}, {"e", "d"}\]> : tensor<8x8x8xf32>
אילוצים:
- חייב לעמוד באילוצים שמפורטים ב-
Sdy_CollectiveOpInterface
. - אם לחלוקה לפלחים של קלט ופלט יש רשתות שונות, לאותה רשת צריכים להיות בדיוק אותם צירים וסדר שונה של מזהי המכשירים.
- לכל מאפיין, המכפלה של גדלי צירי הפיצול ב-
out_sharding
חייבת להיות זהה למכפלה של חלוקת המאפיין התואם של הפונקציה.
מאפיינים: SameOperandsAndResultType
ממשקים: CollectiveOpInterface
, InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
out_sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensor |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.constant
(sdy::ConstantOp)
פעולה קבועה
הפונקציה יוצרת טינסור output
מקבוע value
.
למידע נוסף: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#constant
דוגמה:
%output = sdy.constant dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>
מאפיינים: AlwaysSpeculatableImplTrait
ממשקים: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
השפעות: MemoryEffects::Effect{}
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
value | ::mlir::ElementsAttr | מאפיין וקטור/טנסור קבוע |
תוצאות:
תוצאה | תיאור |
---|---|
output |
טינסור בעל צורה סטטית של ערכים מכל סוג |
sdy.data_flow_edge
(sdy::DataFlowEdgeOp)
פעולת קצה של תהליך העברת הנתונים
תחביר:
operation ::= `sdy.data_flow_edge` $input (`sharding````=``` $sharding^)? attr-dict `:` type($result)
קצה של זרימת נתונים של אופרטור X מגדיר גשר בין קבוצה של מקורות (כל אחד מהם הוא אופרנד של X או אופרנד של סוגר הבלוקים של X) לבין קבוצה של יעדים (כל אחד מהם הוא תוצאה של X או ארגומנט של בלוק של X), כך שכל המקורות והיעדים צריכים להיות מחולקים לאותה דרך.
לפעולה יכולים להיות כמה צמתים של זרימת נתונים שמקבילים זה לזה.
לדוגמה:
y_0, ..., y_n = while (x_0, ..., x_n)
((pred_arg_0,... , pred_arg_n) { ... })
((body_arg_0,..., body_arg_n) {
...
return return_value_0, ..., return_value_n
})
לפעולה הזו של 'while' יש n צמתים של זרימת נתונים, והצומת ה-i של זרימת הנתונים הוא בין המקורות x_i
, return_value_i
לבין היעדים y_i
, pred_arg_i
, body_arg_i
.
פונקציית sdy.data_flow_edge
מקבלת כקלט את הבעלים של קצוות (יכול להיות כל אחד מהיעדים, אבל עדיף תוצאת פעולה ולא ארגומנט של בלוק), שאין לו שימוש אחר. הפעולה הזו לא טהורה כי היא יכולה לקבל קלט שלא היה לו שימוש מקורי.
השדה sdy.data_flow_edge
מכיל גם חלוקה אופציונלית לכל היעדים של הקצה, ויש לעדכן את החלוקה הזו במקום את החלוקה של היעדים (אם אפשר לצרף אותה) במהלך ההעברה. האפשרות הזו שימושית כשיש לפעולה הרבה צמתים, כי יעיל הרבה יותר:
- להפיץ דרך כל קצוות בנפרד.
- לעדכן את חלוקת המשנה של כל קצוות בנפרד במקום את כל היעדים בבת אחת (למשל, לפעולה יש
TensorShardingPerValueAttr
יחיד שאינו ניתן לשינוי לחלוקת המשנה של התוצאות). - מוסיפים כל צלע לרשימת המשימות בנפרד כשחלוקת המקור לשברים השתנתה.
ההעברה תפיץ את החלוקה לפלחים בין כל המקורות והיעדים של sdy.data_flow_edge
כאילו מדובר בפעולה רגילה, עם המקורות כאופרטורים והיעדים כתוצאות, וזהות sdy.op_sharding_rule
. כלומר, העברה קדימה היא מהמקורות ליעדים, והעברה אחורה היא מהיעדים למקורות.
אנחנו לא מאפשרים להגדיר את הקלט של sdy.data_flow_edge
באמצעות אופרטור SdyDialect
, כך שאפשר להניח שהוא מוגדר באמצעות אופרטור שיש לו מאפיין sdy.sharding
לא רשום.
מאפיינים: SameOperandsAndResultType
ממשקים: InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
input |
בצורת ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
בצורת ערכים מכל סוג |
sdy.manual_computation
(sdy::ManualComputationOp)
פעולה במקביל במספר מכשירים באמצעות קבוצות ידניות
תחביר:
operation ::= `sdy.manual_computation` `(`operands`)`
`in_shardings````=```custom<StrippedTensorShardingPerValueAttr>($in_shardings)
`out_shardings````=```custom<StrippedTensorShardingPerValueAttr>($out_shardings)
`manual_axes````=```$manual_axes
custom<SingleBlockRegionNoBlockId>($body)
attr-dict
`:`
functional-type(operands, results)
מעבר לאזור שנכתב בקוד מקומי לכל מכשיר עם קבוצות מפורשות, שבו הצורות הלוגיות תואמות לצורות של מאגרים פיזיים מקומיים לכל מכשיר, והקבוצות תואמות בדיוק לתקשורת פיזית בין מכשירים.
הגוף הוא מקומי ביחס ל-manual_axes. ההעברה תתבצע דרך הגוף בכל צירים חופשיים – אלה שלא נמצאים ברשימה manual_axes.
אילוצים:
- הרכיבים ב-
in_shardings
וב-out_shardings
חייבים לעמוד באילוצים שמפורטים ב-TensorShardingAttr
. - מספר הקלט/הפלט של הטנסורים הגלובליים והמקומיים של אזור הפעולה חייב להיות זהה.
- הצירים הידניים חייבים להופיע לפני כל ציר חופשי בכל חלוקה של מאפיין (dim) לפלחים.
- לא ניתן להוסיף מרווחים לצירים ידניים. כלומר, גודל המאפיין חייב להיות ניתן לחלוקה בגודל של הצירים הידניים התואמים.
- הצורות הגלובליות והמקומיות של הארגומנטים/התוצאות של אזורי הפעולה חייבות להיות זהות.
- לא מתבצעת חלוקה של צירים ידניים.
מאפיינים: IsolatedFromAbove
, RecursiveMemoryEffects
, SingleBlockImplicitTerminator<ReturnOp>
, SingleBlock
ממשקים: ShardableDataFlowOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
in_shardings | ::mlir::sdy::TensorShardingPerValueAttr | חלוקת טינסורים לפי אופרנד/תוצאה של פעולה |
out_shardings | ::mlir::sdy::TensorShardingPerValueAttr | חלוקת טינסורים לפי אופרנד/תוצאה של פעולה |
manual_axes | ::mlir::sdy::ManualAxesAttr | רשימה של צירים שבהם פעולת החישוב הידנית היא ידנית |
אופרטנדים:
אופרנד | תיאור |
---|---|
tensors |
פונקציה וריאדיאלית של טינסור מדורג של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
results |
פונקציה וריאדיאלית של טינסור מדורג של ערכים מכל סוג |
sdy.mesh
(sdy::MeshOp)
רשת בעלת שם
תחביר:
operation ::= `sdy.mesh` $sym_name `=` $mesh attr-dict
הגדרת רשת חדשה בעלת שם. לכל רשתות העמודות במודול צריך להיות אותו מספר מכשירים (למעט רשתות עמודות עם device_id יחיד).
הרשת היא פעולת Symbol
שמופיעה ב-SymbolTable
של המודול, וניתן להפנות אליה באמצעות name
שלה.
מאפיינים: HasParent<ModuleOp>
ממשקים: Symbol
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
sym_name | ::mlir::StringAttr | מאפיין מחרוזת |
mesh | ::mlir::sdy::MeshAttr | רשת של צירים ורשימת מכשירים |
sdy.named_computation
(sdy::NamedComputationOp)
פעולת חישוב בעלת שם
תחביר:
operation ::= `sdy.named_computation` `<`$name`>` `` `(` $operands `)`
(`in_shardings````=```custom<StrippedTensorShardingPerValueAttr>($in_shardings)^)?
(`out_shardings````=```custom<StrippedTensorShardingPerValueAttr>($out_shardings)^)?
custom<SingleBlockRegionNoBlockId>($body)
attr-dict
`:` functional-type($operands, results)
קיבוץ של חישוב, כלומר בלוק של פעולות, ושיוך שם. ההעברה תתבצע אל האזור או ממנו כאילו הכל היה מוטמע.
אפשר להשתמש בזה כדי לטפל בהעברה דרך הוראות קריאה לפונקציות אחרות. משתמשים ב-Shardy צריכים לכתוב תהליך ייבוא/ייצוא שממיר את פעולות הקריאה שלהם לפעולות sdy.named_computation
, ומשכפל/מעתיק את גוף הפונקציה שנקראת לגוף של named_computation
.
הסוג של כל ארגומנטים הבלוק והערכים המוחזרים באזור חייב להיות זהה לסוג של אופרטנדים וסוג התוצאות של הפעולה.
דוגמה:
%1 = sdy.named_computation<"foo">(%0) (%arg1: tensor<16x32xf32>) {
sdy.return %arg1 : tensor<16x32xf32>
} : (tensor<16x32xf32>) -> tensor<16x32xf32>
מאפיינים: IsolatedFromAbove
, RecursiveMemoryEffects
, RecursivelySpeculatableImplTrait
, SingleBlockImplicitTerminator<ReturnOp>
, SingleBlock
ממשקים: ConditionallySpeculatable
, InferTypeOpInterface
, ShardableDataFlowOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
name | ::mlir::StringAttr | מאפיין מחרוזת |
in_shardings | ::mlir::sdy::TensorShardingPerValueAttr | חלוקת טינסורים לפי אופרנד/תוצאה של פעולה |
out_shardings | ::mlir::sdy::TensorShardingPerValueAttr | חלוקת טינסורים לפי אופרנד/תוצאה של פעולה |
אופרטנדים:
אופרנד | תיאור |
---|---|
operands |
משתנה פוליגוני מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
«unnamed» | משתנה פוליגוני מכל סוג |
sdy.propagation_barrier
(sdy::PropagationBarrierOp)
פעולת מחסום להפצה
תחביר:
operation ::= `sdy.propagation_barrier` $input `allowed_direction````=```$allowed_direction attr-dict `:` type($input)
הפעולה הזו פועלת כמו פעולת זהות, ומפיקה את אותו ערך שהיא קיבלה כקלט. אבל מבחינת ההפצה, היא תאפשר לה לעבור דרכו רק בכיוון מסוים.
כך אפשר למנוע את ההפצה של חלוקות בין השימושים בתוצאה של פעולת המחסום לבין המשתנה שלה.
FORWARD
פירושו שחלוקות יכולות לזרום רק מהאופרטנד לתוצאה.BACKWARD
פירושו שחלוקות יכולות לזרום מהתוצאה לפעולה בלבד.- הערך
NONE
מציין שאין חלוקה לפלחים שיכולה להופיע בפעולה הזו. - לא ניתן לציין את
BOTH
, כי הפעולה הזו תהיה מיותרת.
מאפיינים: AlwaysSpeculatableImplTrait
, SameOperandsAndResultType
ממשקים: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
השפעות: MemoryEffects::Effect{}
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
allowed_direction | ::mlir::sdy::PropagationDirectionAttr | טיפוס מנייה (enum) של כיוון ההפצה |
אופרטנדים:
אופרנד | תיאור |
---|---|
input |
טנסור מדורג של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טנסור מדורג של ערכים מכל סוג |
sdy.reshard
(sdy::ReshardOp)
חלוקה מחדש של טנסור לחלוקה אחרת
תחביר:
operation ::= `sdy.reshard` $input $sharding attr-dict `:` type($result)
חלוקה מחדש של הטנסור הקלט לפי החלוקה שצוינה, ששונה מהחלוקה הקיימת של הטנסור הקלט.
שתי הפונקציות ShardingConstraintOp ו-ReshardOp מצרפות חלוקה ל-tensor. משך החיים שלהם הוא:
- לפני ההפצה של חלוקת המשנה, המשתמשים מוסיפים את ShardingConstraintOp.
- ההפצה של הפיצול צורכת את ShardingConstraintOp. אין ShardingConstraintOp בתוצאות של ההפצה של חלוקת המחיצות. במקום זאת, אפשר להוסיף את ReshardOp במקרה הצורך.
- חלוקה למחיצות ממירה את ReshardOp לפעולה קולקטיבית (או לפעולה של זהות). לא אמורה להיות פעולת ReshardOp בתוצאות של מחלק המחיצות.
// TODO(b/331680067). מוסיפים תבנית קנוניזציה כדי להסיר פעולות רישות מיותרות.
מאפיינים: AlwaysSpeculatableImplTrait
, SameOperandsAndResultType
ממשקים: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
השפעות: MemoryEffects::Effect{}
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
input |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.return
(sdy::ReturnOp)
הפעולה sdy.return
מסיימת את האזורים שמצורפים לפעולות מבוססות-אזור של sdy
וכל פעולה מבוססת-אזור אחרת של Shardy. היא משתנה (variadic): כארגומנטים היא מקבלת רשימה של ערכים שיכולים להיות מכל סוג (אבל מאותו סוג, למשל AnyTensor
), ולכן אפשר לעשות בה שימוש חוזר ברמות שונות של סטאק ה-IR של Shardy.
תחביר:
operation ::= `sdy.return` attr-dict ($results^ `:` type($results))?
מאפיינים: AlwaysSpeculatableImplTrait
, Terminator
ממשקים: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
השפעות: MemoryEffects::Effect{}
אופרטנדים:
אופרנד | תיאור |
---|---|
results |
משתנה פוליגוני מכל סוג |
sdy.sharding_constraint
(sdy::ShardingConstraintOp)
הגבלת טינסור לחלוקה (sharding) שצוינה
תחביר:
operation ::= `sdy.sharding_constraint` $input $sharding attr-dict `:` type($result)
הוספת חלוקה לענפים (shard) לטרנספורמציה ביניים (למשל, התוצאה של matmul) כדי לציין איך צריך לפצל את הטרנספורמציה הזו או קבוצת משנה של השימושים שלה.
אם לחלוקה יש מאפיינים פתוחים ויש לה צירים ללא אילוצים, סימן שאפשר לפצל את הטנזור עוד יותר לפי המאפיינים הפתוחים.
הפעולה הזו יכולה:
- אין להם שימוש (dangling) – כלומר, חלוקת המשנה המצורפת היא האופן שבו צריך לפצל את הטנסור של הקלט עצמו.
- יש שימושים – כלומר, חלוקת המשנה המצורפת היא האופן שבו צריך לפצל את השימושים של אופרטורי האילוצים של חלוקת המשנה, בעוד שלשימושים אחרים של הטנזור הקלט עשויה להיות חלוקת משנה שונה (אם אין שימושים אחרים לטנזור הקלט, ההתנהגות זהה לתרחיש 'אין שימושים').
מאפיינים: SameOperandsAndResultType
ממשקים: InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
sharding | ::mlir::sdy::TensorShardingAttr | חלוקה של טינסורים |
אופרטנדים:
אופרנד | תיאור |
---|---|
input |
טינסור של ערכים מכל סוג |
תוצאות:
תוצאה | תיאור |
---|---|
result |
טינסור של ערכים מכל סוג |
sdy.sharding_group
(sdy::ShardingGroupOp)
הגבלה של הטנסורים בקבוצה כך שיהיה להם אותו חלוקה לפלחים.
תחביר:
operation ::= `sdy.sharding_group` $input `group_id````=```$group_id attr-dict `:` type($input)
הפעולה הזו מספקת ממשק להקצאת טינסורים לקבוצות חלוקה (קבוצות של טינסורים שיחולו עליהן אכיפה של חלוקות זהות). במהלך ההעברה, ברגע שרכיב אחד של הקבוצה מחולק למקטעים, כל שאר המרכיבים מחולקים למקטעים באותו אופן. הפעולה הזו מקבלת את מזהה הקבוצה של הארגומנט ולא מחזירה תוצאה, אלא משנה את הייצוג הפנימי של קבוצת הפיצול כדי להוסיף את הטנסור של הקלט לקבוצה עם המזהה הנתון.
ממשקים: InferTypeOpInterface
מאפיינים:
מאפיין | סוג MLIR | תיאור |
---|---|---|
group_id | ::mlir::IntegerAttr | מאפיין של מספר שלם ללא סימן ב-64 ביט |
אופרטנדים:
אופרנד | תיאור |
---|---|
input |
טנסור מדורג של ערכים מכל סוג |
מאפיינים
AllToAllParamAttr
פרמטר 'הכול לכולם'
תחביר:
#sdy.all_to_all_param<
::llvm::ArrayRef<AxisRefAttr>, # axes
int64_t, # src_dim
int64_t # tgt_dim
>
קבוצת ערכים (tuple) שמכילה את הצירים ואת מאפייני המקור/היעד שעבורם רוצים לבצע את ההתאמה 'הכול אל הכול'.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
צירים | ::llvm::ArrayRef<AxisRefAttr> |
הצירים שבהם יתבצע החישוב של כל-אל-כל |
src_dim | int64_t |
אינדקס המאפיין 'מקור' |
tgt_dim | int64_t |
האינדקס של מאפיין היעד |
AlltoAllParamListAttr
רשימת פרמטרים של 'הכול לכולם'
תחביר:
#sdy.all_to_all_param_list<
::llvm::ArrayRef<AllToAllParamAttr> # value
>
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
ערך | ::llvm::ArrayRef<AllToAllParamAttr> |
AxisRefAttr
הפניה לציר מלא או לציר משנה מפוצל
תחביר:
#sdy.axis_ref<
::llvm::StringRef, # name
SubAxisInfoAttr # sub_axis_info
>
אילוצים:
- הערך
name
חייב להופיע במגבלהMeshAttr
. - אם השדה
sub_axis_info
קיים, הוא חייב לעמוד במגבלות שלSubAxisInfoAttr
.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
שם | ::llvm::StringRef |
שם הציר |
sub_axis_info | SubAxisInfoAttr |
מידע נוסף אם זהו ציר משנה |
AxisRefListAttr
רשימת הפניות לצירים
תחביר:
#sdy.axis_ref_list<
::llvm::ArrayRef<AxisRefAttr> # value
>
אילוצים:
- הרכיבים ב-
value
חייבים לעמוד באילוצים שלAxisRefAttr
. - אין ערכי axis-ref כפולים או צירים משניים שחופפים זה לזה.
- אין שני ערכי axis-ref סמוכים שהם צירים משניים עוקבים של אותו ציר מלא, כלומר אפשר למזג אותם לציר משני אחד או לציר המלא.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
ערך | ::llvm::ArrayRef<AxisRefAttr> |
DimMappingAttr
רשימת מדדי הגורמים של מאפיין
רשימה ריקה מציינת שמדובר במיפוי null (הניתוח או ההדפסה מתבצעים באמצעות *
), כלומר המאפיין לא ממופה לאף גורם.
אילוצים:
- יש לפחות אינדקס גורם אחד.
- אינדקסי הגורמים חייבים להיות בטווח [0,
$factor_sizes
). - אם יש כמה גורמים, אף אחד מהם לא יכול להיות בגודל 1.
- אין אינדקסים כפולים של גורמים.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
factor_indices | ::llvm::ArrayRef<int64_t> |
הגורמים שאליהם המאפיין הזה ממופה |
DimensionShardingAttr
חלוקה של מאפיינים לפלחים
רשימה של שמות צירים לחלוקה של מאפיין טינסור, מגדול לקטן, משתנה בוליאני שמציין אם אפשר לפצל את המאפיין עוד יותר ומספר שלם אופציונלי שמציין את העדיפות של חלוקת המאפיין, שתישמר במהלך ההפצה של החלוקה. רמות העדיפות מגיעות מהערות של חלוקת משתמשים, וערך נמוך יותר מציין רמת עדיפות גבוהה יותר. אם העדיפות לא מופיעה בהערה, המערכת תשתמש בעדיפות הגבוהה ביותר.
אילוצים:
- הרכיבים ב-
axes
חייבים לעמוד באילוצים שמפורטים ב-AxisRefListAttr
. - אם יש עדיפות לחלוקה של מאפיין:
- העדיפות חייבת להיות גדולה מ-0 או שווה לה.
- אם המאפיין סגור, יש לו לפחות ציר אחד.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
צירים | ::llvm::ArrayRef<AxisRefAttr> |
הפניות לצירים |
is_closed | bool |
אם אי אפשר לפצל את המאפיין הזה עוד יותר |
הקמפיין | std::optional<int64_t> |
העדיפות שמשמשת במהלך ההפצה על סמך עדיפות המשתמש |
ListOfAxisRefListsAttr
רשימת רשימות של ערכי עזר לצירים
תחביר:
#sdy.list_of_axis_ref_lists<
::llvm::ArrayRef<AxisRefListAttr> # value
>
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
ערך | ::llvm::ArrayRef<AxisRefListAttr> |
ManualAxesAttr
רשימת צירים שבהם פעולת החישוב הידנית היא ידנית
תחביר:
#sdy.manual_axes<
::llvm::ArrayRef<StringAttr> # value
>
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
ערך | ::llvm::ArrayRef<StringAttr> |
MeshAttr
רשת של צירים ורשימת מכשירים
תחביר:
#sdy.mesh<
::llvm::ArrayRef<MeshAxisAttr>, # axes
::llvm::ArrayRef<int64_t> # device_ids
>
רשת היא רשימה של צירים ורשימה אופציונלית של מזהי מכשירים שמציינת את סדר המכשירים.
אם רשימת הצירים ריקה, למערך יש ציר משתמע ללא שם בגודל 1. במקרה כזה, אם לא מציינים רשימת מזהי מכשירים, רשימת מזהי המכשירים המשתמעת היא [0]. אם מציינים רשימת מזהי מכשירים, היא חייבת להכיל מספר שלם אחד של ערך לא שלילי כלשהו. אנחנו מכנים את המצב הזה 'מקסימום חלוקה למקטעים'.
בכל המקרים שבהם לא מתבצעת חלוקה מקסימלית, אם מצוין רשימת מזהי מכשירים, המכפלה של גדלי הצירים צריכה להתאים למספר המכשירים. אם לא צוינה רשימה של מזהי מכשירים, רשימת מזהי המכשירים המשתמעת היא iota(product(axes)). כדי לפשט את התהליך, אסור גם לציין רשימת מזהי מכשירים זהה ל-iota(product(axes)); במקרה כזה, אסור לציין רשימת מזהי מכשירים.
ריכזנו כאן כמה דוגמאות לרשתות:
- רשת ריקה מייצגת רשת placeholder שאפשר להחליף במהלך ההעברה: <[]>
- רשת עם ציר ללא שם ומזהה מכשיר מפורש, שמשמשת בדרך כלל לייצוג חלוקה מקסימלית: <[], device_ids=[3]>
- רשת עם שני צירים ומזהי מכשירים משתמעים iota(6): <["a"=2, "b"=3]>
- רשת עם שני צירים ומזהי מכשירים מפורשים שמציינים את סדר המכשירים: <["a"=3, "b"=2], device_ids=[0, 2, 4, 1, 3, 5]>
אילוצים:
- אסור לכלול ב-
axes
רכיבים עם שמות כפולים. - אם מציינים את
device_ids
:- המכפלה של גדלי הצירים חייבת להתאים למספר המכשירים.
- כל האלמנטים שלו חייבים להיות לא שליליים.
- הערך של
device_ids
לא צריך להיות שווה ל-iota(product(axis_sizes))
. device_ids
הממוין חייב להיותiota(product(axis_sizes))
.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
צירים | ::llvm::ArrayRef<MeshAxisAttr> |
צירים של רשת |
device_ids | ::llvm::ArrayRef<int64_t> |
סדר מכשיר מפורש או מזהה מכשיר מקסימלי |
MeshAxisAttr
ציר בעל שם במערך
תחביר:
#sdy.mesh_axis<
::llvm::StringRef, # name
int64_t # size
>
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
שם | ::llvm::StringRef |
שם |
size | int64_t |
הגודל של הציר הזה |
OpShardingRuleAttr
הגדרה של אופן חלוקת הפעולה למחיצות.
תחביר:
#sdy.op_sharding_rule<
::llvm::ArrayRef<int64_t>, # factor_sizes
::llvm::ArrayRef<TensorMappingAttr>, # operand_mappings
::llvm::ArrayRef<TensorMappingAttr>, # result_mappings
::llvm::ArrayRef<int64_t>, # reduction_factors
::llvm::ArrayRef<int64_t>, # need_replication_factors
::llvm::ArrayRef<int64_t>, # permutation_factors
::llvm::ArrayRef<int64_t>, # blocked_propagation_factors
bool # is_custom_rule
>
כלל חלוקה מפריד את הפעולה לפי מאפיינים שונים שלה – מאפיינים, צורת אופרטורים, צורת התוצאות וכו'. לדוגמה:
%0 = stablehlo.add %arg0, %arg1 {
sdy.sharding_rule = #sdy.op_sharding_rule<
([i, j],[i, j])->([i, j])
{i=8, j=8}>
} : tensor<8x8xf32>
%1 = stablehlo.dot_general %arg2, %arg3, contracting_dims = [1] x [0] {
sdy.sharding_rule = #sdy.op_sharding_rule<
([i, k],[k, j])->([i, j])
{i=8, j=16, k=8}>
}: (tensor<8x8xf32>, tensor<8x16xf32>) -> tensor<8x16xf32>
שימו לב: אנחנו מאפשרים שימוש בגורמים בגודל 1 גם אם אי אפשר לפצל אותם, בעיקר כדי לשמור על השלמות. לפעולות רבות, כמו פעולות נקודה-לנקודה, יש מאפיינים בגודל 1 שתואמים בין אופרטנדים לתוצאות.
סוגי גורמים:
- המאפיין
reduction_factors
מכיל את המדדים של הגורמים שצריך לצמצם, כמו המאפיינים המתכווצים בפעולה של נקודה (dot). - השדה
need_replication_factors
מכיל את האינדקסים של הגורמים שדורשים שכפול מלא, כמו המאפיין הממוין בפעולת מיון. - השדה
permutation_factors
מכיל את האינדקסים של הגורמים שדורשים טרנספוזיציה קולקטיבית אם הם מחולקים למקטעים, כמו מאפייני המילוי בפעולה של מילוי. - כל שאר הגורמים נחשבים לגורמים של העברה, כלומר גורמים שלא דורשים תקשורת אם הם מחולקים לאותה דרך בכל הטנסורים שממופים אליהם.
blocked_propagation_factors
מכיל את הגורמים שאסור להפיץ ביניהם חלוקות. הוא אורתוגונלי לסוגים של הגורמים. כלומר, גורם של העברה חסומה יכול להיות כל אחד מסוגי הגורמים.
השדה is_custom_rule
מתאר אם זהו כלל שהוגדר על ידי משתמש. המשתמשים יכולים להגדיר כללי חלוקה לפלחים לשיחות בהתאמה אישית, או לשנות את כללי החלוקה לפלחים שהוגדרו מראש לפעולות הרגילות. כלל מותאם אישית תמיד נשמר ולא מוסר אף פעם.
אילוצים:
- מספר המיפויים של אופרטורים/תוצאות חייב להתאים למספר האופרטורים/התוצאות של הפעולה.
- יש לפחות מיפוי אחד (לא ניתן ליצור כלל לפעולה ללא אופרטורים/תוצאות).
- הדירוג של כל
TensorMappingAttr
תואם לדירוג של סוג הטנזור התואם. - לכל קבוצה של גורמים (
reduction_factors
,need_replication_factors
, permutation_factors
):- הרכיבים חייבים להיות בטווח [0,
$factor_sizes
]. - אין אינדקסים כפולים של גורמים בכל קבוצה ובין הקבוצות.
- הרכיבים חייבים להיות בטווח [0,
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
factor_sizes | ::llvm::ArrayRef<int64_t> |
הגדלים של כל הגורמים בכלל הזה |
operand_mappings | ::llvm::ArrayRef<TensorMappingAttr> |
מיפויים של אופרנדים |
result_mappings | ::llvm::ArrayRef<TensorMappingAttr> |
מיפויים של תוצאות |
reduction_factors | ::llvm::ArrayRef<int64_t> |
גורמים שצריך לצמצם |
need_replication_factors | ::llvm::ArrayRef<int64_t> |
גורמים שמחייבים שכפול מלא |
permutation_factors | ::llvm::ArrayRef<int64_t> |
גורמים שדורשים collective-permute |
blocked_propagation_factors | ::llvm::ArrayRef<int64_t> |
גורמים שלא מתבצעת בהם העברה של חלוקות |
is_custom_rule | bool |
אם הכלל מיועד ל-stablehlo.custom_call |
SubAxisInfoAttr
מידע על האופן שבו ציר המשנה הזה נגזר מהציר המלא
תחביר:
#sdy.sub_axis_info<
int64_t, # pre_size
int64_t # size
>
כשמפצלים ציר מלא ל-n צירי משנה, הציר עובר שינוי צורה ל-[k_1,…,k_n], וציר המשנה ה-i יכול להתבטא במכפלה של כל המימדים של הציר שמשמאל לו m=prod(k_1,...,k_(i-1))
(שנקרא גם 'מידה מוקדמת') והמידה k_i. לכן, המאפיין sub-axis-info מכיל את שני המספרים האלה, והוא מסומן באופן הבא: (m)k
עבור m לפני הגודל ו-k בגודל.
אילוצים:
- הערך של
pre-size
הוא לפחות 1. - הערך של
size
גדול מ-1. pre-size
חייב לחלק את הגודל של הציר המלא, כלומר גםpre-size
וגםsize
צריכים לחלק את הגודל של הציר המלא, והציר המשני לא יכול לחרוג מהציר המלא.- הגודל של ציר המשנה לא שווה לגודל של הציר המלא התואם. במקרה כזה, צריך להשתמש בציר המלא במקום זאת.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
pre_size | int64_t |
המכפלה של גדלי צירי המשנה שמשמאל לציר המשנה הזה |
size | int64_t |
הגודל של ציר המשנה הזה |
TensorMappingAttr
מיפויים של גורמים לכל מאפיין של טינסור.
תחביר:
#sdy.tensor_mapping<
::llvm::ArrayRef<DimMappingAttr> # dim_mappings
>
אילוצים:
- הרכיבים ב-
dim_mappings
חייבים לעמוד באילוצים ב-DimMappingAttr
. - אין אינדקסים כפולים של גורמים במאפיינים שונים.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
dim_mappings | ::llvm::ArrayRef<DimMappingAttr> |
מיפויים של מאפיינים |
TensorShardingAttr
חלוקה של טינסורים
תחביר:
#sdy.sharding<
::mlir::Attribute, # mesh_or_ref
::llvm::ArrayRef<DimensionShardingAttr>, # dim_shardings
::llvm::ArrayRef<AxisRefAttr> # replicated_axes
>
חלוקת טינסור מחולקת למערך ספציפי, וניתן להפנות רק לשמות של צירים מהמערך הזה. חלוקות המאפיינים מאפשרות לנו לדעת לאורך אילו צירים (או צירי משנה) מתבצעת חלוקה של כל מאפיין של הטנזור, מהציר הראשי ועד לציר המשני. כל שאר הצירים שלא מחלקים מאפיין לקטעים עוברים שכפול, בין שבאופן משתמע ובין שבאופן מפורש (אם הם מופיעים ברשימת הצירים המשוכפלים).
אפשר לציין את הרשת שממנה מחולק המשנה באמצעות שם של סמל, שמפנה לסמל MeshOp
תואם, או באמצעות MeshAttr
מוטמע.
אילוצים:
- הרכיבים ב-
dim_shardings
חייבים לעמוד באילוצים שמפורטים ב-DimensionShardingAttr
. - הרכיבים ב-
replicated_axes
חייבים לעמוד באילוצים שמפורטים ב-AxisRefListAttr
. - אם סוג הטנזור התואם הוא לא
ShapedType
, לרשת הפילוח צריכה להיות דרגה 0 ואין לה צירים משוחזרים. - למערך הטנזור צריך להיות דרגה.
- מספר החלוקות של המאפיינים שווה לדרגה של הטנזור.
- מאפיינים בגודל 0 לא מחולקים למקטעים.
- הפריטים ב-
replicated_axes
מסודרים לפיmesh_or_ref
(ראוAxisRefAttr::getMeshComparator
).
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
mesh_or_ref | ::mlir::Attribute |
מאפיין mesh או מאפיין עזר של סמל רשת שטוח |
dim_shardings | ::llvm::ArrayRef<DimensionShardingAttr> |
חלוקות למקטעים של מאפיינים |
replicated_axes | ::llvm::ArrayRef<AxisRefAttr> |
הפניות לצירים |
TensorShardingPerValueAttr
חלוקה של טינסורים לפי אופרטור/תוצאה של פעולה
תחביר:
#sdy.sharding_per_value<
::llvm::ArrayRef<TensorShardingAttr> # shardings
>
רשימה של TensorShardingAttr
, אחת לכל אופרנד/תוצאה של פעולה.
אילוצים:
- הרכיבים ב-
shardings
חייבים לעמוד באילוצים שלTensorShardingAttr
.
פרמטרים:
פרמטר | טיפוס C++ | תיאור |
---|---|---|
חלוקות | ::llvm::ArrayRef<TensorShardingAttr> |
חלוקה לפלחים לפי ערך |
טיפוסים בני מנייה (enum)
PropagationDirection
Propagation direction enum
מקרים:
סמל | ערך | מחרוזת |
---|---|---|
ללא | 0 |
ללא |
קדימה | 1 |
קדימה |
אחורה | 2 |
אחורה |
BOTH | 3 |
BOTH |