}

}

دفعة شكل مصفوفة الضرب

في مهام الترجمة، دفعة ضرب المصفوفات من الشكل التقليدي تحسب بضرب المصفوفات في أعباء العمل الأخرى أصغر من ذلك بكثير. وطول الخطوة فك محول حول شكل الحكم الإدخال. عموما أقل من 30.

كما البعد دفعة، عندما يكون حجم دفعة لسبب معين، بل هو عدد ثابت. على سبيل المثال، إذا كان حجم الدفعة 16، حجم شعاع 4، دفعة البعد هو 16 * 4 * #head (عدد رئيس متعددة headattention عادة ثمانية). مصفوفة M، K، N هو في حدود أو ل.

مشكلات في الأداء الضرب دفعة مصفوفة

أولا، أجرينا التحليل النظري للدفعة FLOP ضرب المصفوفات النواة. نتائج مثيرة للاهتمام للغاية: كل دفعة حسابيا الضرب مصفوفة مكثف تقتصر (TFLOP أقل من 1).

ثم وصفنا cuBLAS nvprof أداء دفعة من عدة الضرب شكل مصفوفة. الجدول التالي هو استخدام بعض المؤشرات NVIDIA GPU M40 (CUDA 8.0) الحصول عليها.

حتى أشكال مختلفة (متفاوتة بين M، N، K)، يتم تنفيذها وجميع maxwell_sgemmBatched_128x128_raggedMn_tn نفس العدد FLOP، وهو أكبر بكثير من القيمة النظرية. ويمكن استنتاج أن جميع هذه الأشكال المختلفة تمتلئ في نهاية المطاف إلى شكل المحدد. في جميع الأشكال، وحتى في أفضل الأحوال، FLOP التنفيذ العملي النظري فقط من 2.74 FLOP، معظم الحسابات هي زائدة عن الحاجة. وبالمثل، دعوة إلى maxwell_sgemmBatched_64x64_raggedMn_tn cuBLAS آخر نواة ظهرت أيضا في نفس الوضع.

ومن الواضح أن المصفوفة cuBLAS أدنى دفعة الضرب الكفاءة. لهذا السبب، ونحن نستخدم NMT TVM توليد دفعة فعالة نواة ضرب المصفوفات.

يتم احتساب دفعة بضرب مصفوفة

في TVM، ويحسب العادية مصفوفة دفعة بضرب البيان التالي:

# تمثيل الحساب

A = tvm.placeholder ((دفعة، M، K)، اسم = 'A')

B = tvm.placeholder ((دفعة، K، N)، اسم = 'B')

ك = tvm.reduce_axis ((0، K)، 'ك')

C = tvm.compute ((دفعة، M، N)،

امدا ب، ص، س: tvm.sum (A  * B ، المحور = ك)،

اسم = 'C')

جدولة الأمثل

بعد الإعلان عن الحساب، ونحن بحاجة إلى تصميم بعناية من المقرر أن تلعب أداء أفضل.

ضبط المعلمات / عدد المواضيع كتلة

# مؤشرات الموضوع

block_y = tvm.thread_axis ( "blockIdx.y")

block_x = tvm.thread_axis ( "blockIdx.x")

thread_y = tvm.thread_axis ((0، num_thread_y)، "threadIdx.y")

thread_x = tvm.thread_axis ((0، num_thread_x)، "threadIdx.x")

thread_yz = tvm.thread_axis ((0، vthread_y)، "vthread"، اسم = "VY")

thread_xz = tvm.thread_axis ((0، vthread_x)، "vthread"، اسم = "VX")

# بلوك التقسيم

BB، FF، MM، PP = ق .op.axis

BBFF = ق .fuse (BB، FF)

MMPP = ق .fuse (MM، PP)

من قبل، ty_block = ق .split (BBFF، عامل = num_thread_y * vthread_y)

ب س، tx_block = ق .split (MMPP، عامل = num_thread_x * vthread_x)

الصورة .bind (من قبل، block_y)

الصورة .bind (ب س، block_x)

vty، تاي = ق .split (ty_block، nparts = vthread_y)

VTX، تكساس = ق .split (tx_block، nparts = vthread_x)

الصورة .reorder (من قبل، BX، vty، VTX، تاي، تكساس)

الصورة .reorder (من قبل، ب س، تي واي، تكساس)

الصورة .bind (تاي، thread_y)

الصورة .bind (تكساس، thread_x)

الصورة .bind (vty، thread_yz)

الصورة .bind (VTX، thread_xz)

نحن مزيج من الخارجي أبعاد دفعة ضرب المصفوفات، على سبيل المثال، BB وFF-المرجع الأبعاد حساب مصفوفة في مضاعفة دفعة وأشار عادة باسم البعد "دفعة"، ونحن تقسيم الأبعاد الخارجية والداخلية من قبل عامل (number_thread * vthread).

أنها لا تتطلب مصفوفة الضرب دفعة واسطة Strided، وبالتالي يتم تعيين عدد من الظاهرية (vthready وvthreadx) المواضيع إلى 1.

العثور على أفضل مزيج من number_thread

وتستند النتائج التالية على NVIDIA GPU M40 (CUDA 8.0).

وبناء على التجارب السابقة، وسيلة للعثور على أفضل مزيج num_thread_y والعنف num_thread_x هي من خلال البحث (بحث شامل). بعد بحث شامل، للعثور على أفضل مزيج من شكل الحالية في حساب الحالي، num_thread_y = 8، num_thread_x = 32.

دمج دفعة مع عملية ضرب المصفوفات الأخرى

وcuBLAS "الصندوق الاسود" القائمة المكالمات المكتبة تستخدم عادة في حدود استراتيجيات التحسين "المرجع الانصهار". ومع ذلك، الاستخدام الفعال للدفعة نواة ضرب المصفوفات ولدت، يمكن بسهولة كسر حدود الانصهار، وليس فقط الانصهار بين العناصر الفردية، فمن الممكن الحصول على تحسين أداء أفضل.

وكما يتبين من الشكل المحتسبة بعد ضرب دفعة ويرافق دائما من قبل البث مصفوفة الأفعى أو عملية تبديل.

بواسطة "إضافة" أو "تبديل" عملية ضرب المصفوفات والتكامل دفعة يمكن أن تقلل من النواة فوق إطلاق وزائدة وقت وصول الذاكرة.

تحديد إضافة دفعة ضرب المصفوفات والبث حساب التقارب على النحو التالي:

# تمثيل الحساب

A = tvm.placeholder ((batch_size، وميزات، M، K)، اسم = 'A')

# شكل B هو (N، K) غير (K، N) لأن B ونقلها هو هذا النمط الانصهار

B = tvm.placeholder ((batch_size، وميزات، N، K)، اسم = 'B')

ENTER = tvm.placeholder ((batch_size، 1، M، N)، اسم = 'ENTER')

ك = tvm.reduce_axis ((0، K)، 'ك')

C = tvm.compute (

(Batch_size، وميزات، M، N)،

امدا YB، YF، م، س: tvm.sum (A  * B ، المحور = ك)،

اسم = 'C')

D = topi.broadcast_add (C، ENTER)

تبديل ضرب المصفوفات والانصهار البيانات دفعة تحسب على النحو التالي:

# تمثيل الحساب

A = tvm.placeholder ((batch_size، وميزات، M، K)، اسم = 'A')

B = tvm.placeholder ((batch_size، وميزات، K، N)، اسم = 'B')

ك = tvm.reduce_axis ((0، K)، 'ك')

C = tvm.compute (

(Batch_size، M، الميزات، N)،

امدا YB، م، YF، س: tvm.sum (A  * B ، المحور = ك)،

اسم = 'C')

الانصهار الأساسية الأداء

توليد أداء رمز اختبار تحديد، لتشكيل . 17 كما طول تسلسل يتم تحديد لأنه هو متوسط طول مدخلات الإنتاج لدينا.

• TF-R1.4 BatchMatmul: 513.9 لنا

• TF-R1.4 BatchMatmul + تبديل (منفصل): 541.9 لنا

• TVM BatchMatmul: 37.62 لنا

• TVM BatchMatmul + تبديل (تنصهر): 38.39 لنا

نواة التكامل الأمثل يجلب 1.7 مرة التسارع.

TensorFlow المتكاملة

في عبء العمل لدينا، يتم إدخال دفعة شكل مصفوفة الضرب محدودة، من السهل تعداد مقدما. مع هذه الشكل، محددة مسبقا، ونحن يمكن أن تولد الأمثل للغاية مقدما CUDA النواة (محسوبة شكل ثابت يمكن أن يحقق أفضل إمكانات الأمثل). وفي الوقت نفسه، فإن GM أيضا إنشاء شكل دفعة مصفوفة ليتم ضرب أكثر من نواة، وتوفير آلية للتراجع في شكل نواة ولا يتم إنشاؤها مسبقا.

ونحن سوف تدمج Tensorflow ولدت في قلب لشكل فعال ومحدد من آلية تراجعية. لقد وضعت عددا من عمليات الاندماج، مثل BatchMatMulTranspose أو BatchMatMulAdd-- استخدام TVM وقت التشغيل API بداية توليد محددة نواة أو دعوة تراجع نواة لتحديد شكل الإدخال.

تمرير الأمثل عن طريق إجراء، تلقائيا محل الدفعة matmul الأصلي + إضافة / عملية تبديل باستخدام الانصهار. وفي الوقت نفسه، جنبا إلى جنب مع FIG الأمثل أكثر عدوانية باس، حاولنا أن تولد أكثر كفاءة في الذيل طويل وضع الانصهار نواة العملية باستخدام TVM، لمواصلة تعزيز أداء النهاية.

ملخص

في علي بابا، وجدنا أن TVM هو أداة فعالة جدا لتطوير النوى GPU عالية الأداء لتلبية الاحتياجات الداخلية.

في هذا بلوق، ونحن محول نموذج على سبيل المثال، يصف استراتيجية التحسين نستخدم TVM.

أولا، مرت علينا أول تحليل مبدأ تحديد القضايا الرئيسية محول نموذج. نحن ثم استخدامها لتوليد الأمثل للغاية TVM CUDA إصدار النواة cuBLAS استبداله (حالة تسارع 13 ضعفا).

بعد ذلك، بعد استخدام الأساسية لدمج آلية الانصهار TVM دفعة قبل الضرب مصفوفة / عمليات لجلب مزيد من التحسينات الأداء (1.7 أضعاف تحسين الأداء). أداء نهاية تحسنت إلى 1.4 مرة. ولدت على أساس هذه الحبوب، قمنا بتطوير FIG باس الأمثل، واستبدال ويتم احتساب وضع انصهار النواة الأصلي تلقائيا من قبل TVM، لضمان الأمثل شفافة للمستخدمين النهائيين.

وأخيرا، يتم دمج كل هذه التحسينات في طريقة المتباعدة إلى TensorFlow، الأمر الذي يدل على الطرق المحتملة وفي مختلف أعماق TVM التكامل إطار التعلم.

حاليا لدينا التقدم في العمل - وسيتم دمج TVM إلى TensorFlow الخلفية من codegen. نأمل أن يشارك المزيد من الإنجازات في المستقبل مع المجتمع.

عن طريق: HTTP: //www.tvmlang.org

شبكة لى فنغ AI Yanxishe التشطيب مترجم.