👩‍🎓 🕧 🤶 كتابة مواصفات Nvidia Kepler (ثنائيات CUDA ، إصدار لغة sm

بالنسبة للغات المعالج الشائعة ، تمت كتابة الكثير من المواصفات بالفعل لـ Ghidra ، ولكن لا شيء للغات الرسومية. إنه أمر مفهوم ، لأنه يحتوي على تفاصيله الخاصة: المسندات والثوابت التي يتم من خلالها تمرير المعلمات ، بما في ذلك الأشياء الأخرى الموروثة من التظليل. بالإضافة إلى ذلك ، فإن التنسيق المستخدم لتخزين الرمز غالبًا ما يكون ملكية ، وتحتاج إلى عكسه بنفسك.

في هذه المقالة ، باستخدام مثالين ، سنكتشف ما هو.

البرنامج الأول هو أبسط محور (التناظرية من العالم مرحبا GPGPU). والثاني يساعد على فهم تنفيذ الشروط والقفزات على GPU ، منذ ذلك الحين كل شيء مختلف هناك.

تستخدم جميع لغات Nvidia ترميزًا نهائيًا صغيرًا ، لذا انسخ على الفور وحدات البايت من محرر السداسي في بعض المفكرة (على سبيل المثال ، Notepad ++) بترتيب عكسي من 8 قطع (طول التعليمات ثابت هنا). ثم ، من خلال الآلة الحاسبة للمبرمجين (تلك التي من Microsoft مناسبة) ، نترجمها إلى كود ثنائي. بعد ذلك ، نبحث عن التطابقات ، ونؤلف قناع التعليمات ، ثم المعاملات. لفك الشفرة والبحث عن القناع ، تم استخدام محرر السداسي ومحرر cuobjdump ، وأحيانًا يكون المجمع مطلوبًا ، كما هو الحال في AMDGPU (لأن أداة التفكيك غير متاحة هناك ، ولكن هذا موضوع لمقال منفصل). إنه يعمل على هذا النحو: نحاول عكس جميع البتات المشبوهة بالتسلسل في الآلة الحاسبة ، ثم نحصل على قيمة سداسية عشرية جديدة للبايتات ، ونستبدلها في ثنائي تم تجميعه عبر nvcc أو مجمع ، إذا كان موجودًا ، وهذا ليس هو الحال دائمًا.ثم من خلال cuobjdump نتحقق.

أنشر شفرة المصدر بالتنسيق (بشكل أساسي في C ، بدون الإيجابيات و OOP من أجل اتصال أوثق برمز GPU الخاص بالجهاز) ، ثم تشويش + بايت في وقت واحد ، لأن هذا أكثر ملاءمة ، لا تحتاج إلى التبديل.

نقوم بنسخه إلى axpy.cu وتجميعه من خلال cmd: nvcc axpy.cu --cubin - gpu-architecture sm_30

ملف ELF الناتج المسمى axpy.cubin غير مجمّع في نفس المكان: cuobjdump axpy.cubin -sass

مثال 1:

__global__ void axpy(float param_1, float* param_2, float* param_3) {
unsigned int uVar1 = threadIdx.x;
param_2[uVar1] = param_1 * param_3[uVar1];
}

تفريغ

/*0000*/
/* 0x22c04282c2804307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ MOV32I R5, 0x4;
/* 0x1800000010015de2 */
/*0020*/ ISCADD R2.CC, R0, c[0x0][0x150], 0x2;
/* 0x4001400540009c43 */
/*0030*/ LD.E R2, [R2];
/* 0x8400000000209c85 */
/*0038*/ ISCADD R4.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520011c43 */
/*0040*/
/* 0x20000002e04283f7 */
/*0048*/ IMAD.U32.U32.HI.X R5, R0, R5, c[0x0][0x14c];
/* 0x208a800530015c43 */
/*0050*/ FMUL R0, R2, c[0x0][0x140];
/* 0x5800400500201c00 */
/*0058*/ ST.E [R4], R0;
/* 0x9400000000401c85 */
/*0060*/ EXIT;
/* 0x8000000000001de7 */
/*0068*/ BRA 0x68;
/* 0x4003ffffe0001de7 */
/*0070*/ NOP;
/* 0x4000000000001de4 */
/*0078*/ NOP;
/* 0x4000000000001de4 */

نتيجة فك

void axpy(float param_1,float *param_2,float *param_3) {
  uint uVar1;
  
  uVar1 = *&threadIdx.x;
  param_2[uVar1] = param_3[uVar1] * param_1;
  return;
}

مثال 2:

__global__ void predicates(float* param_1, float* param_2) {
    unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
    if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
    else param_2[uVar1] = uVar1;
}

تفريغ

/*0000*/
/* 0x2272028042823307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ S2R R3, SR_CTAID.X;
/* 0x2c0000009400dc04 */
/*0020*/ IMAD R0, R3, c[0x0][0x28], R0;
/* 0x20004000a0301ca3 */
/*0028*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0030*/ IADD32I R2, R0, -0x6;
/* 0x0bffffffe8009c02 */
/*0038*/ I2F.F32.U32 R4, R0;
/* 0x1800000001211c04 */
/*0040*/
/* 0x22c042e04282c2c7 */
/*0048*/ ISETP.GE.U32.AND P0, PT, R2, 0x4, PT;
/* 0x1b0ec0001021dc03 */
/*0050*/ @P0 ISCADD R2.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520008043 */
/*0058*/ @P0 IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x14c];
/* 0x208680053000c043 */
/*0060*/ @P0 ST.E [R2], R4;
/* 0x9400000000210085 */
/*0068*/ @P0 EXIT;
/* 0x80000000000001e7 */
/*0070*/ ISCADD R2.CC, R0, c[0x0][0x140], 0x2;
/* 0x4001400500009c43 */
/*0078*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0080*/
/* 0x2000000002e04287 */
/*0088*/ IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x144];
/* 0x208680051000dc43 */
/*0090*/ ST.E [R2], R4;
/* 0x9400000000211c85 */
/*0098*/ EXIT;
/* 0x8000000000001de7 */
/*00a0*/ BRA 0xa0;
/* 0x4003ffffe0001de7 */
/*00a8*/ NOP;
/* 0x4000000000001de4 */
/*00b0*/ NOP;
/* 0x4000000000001de4 */
/*00b8*/ NOP;
/* 0x4000000000001de4 */

نتيجة فك

void predicates(float *param_1,float *param_2) {
  uint uVar1;
  
  uVar1 = *&blockIdx.x * (int)_DAT_constants_00000028 + *&threadIdx.x;
  if (uVar1 - 6 < 4) {
    param_1[uVar1] = (float)uVar1;
    return;
  }
  param_2[uVar1] = (float)uVar1;
  return;
}

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

بشكل عام ، القاعدة هي هذه - لاختبار الواجهة الأمامية ، نأخذ أي مثال بسيط (بحد أدنى من التحسينات الممكنة) مناسب أولاً (إعادة إنتاج الأخطاء). بالنسبة للباقي ، فإن الشفرة المترجمة ستكون لها بالفعل تحسينات (أو تصححها بطريقة أو بأخرى من خلال إعادة البناء). ولكن في الوقت الحالي ، فإن المهمة الرئيسية هي على الأقل مجرد رمز صحيح يفعل نفس الشيء مثل رمز الجهاز. هذا هو "نمذجة البرمجيات". "نمذجة البرمجيات" نفسها لا تعني إعادة هيكلة ، ترجمة C إلى C ++ ، استعادة الطبقات ، وأكثر من ذلك مثل تحديد القوالب.

الآن نحن نبحث عن أنماط فن الإستذكار والمعاملات والمعدلات.

للقيام بذلك ، قارن البتات (في التمثيل الثنائي) بين التعليمات المشبوهة (أو السلاسل ، إذا كان من الأسهل الاتصال بها بهذه الطريقة). يمكنك أيضًا استخدام ما ينشره المستخدمون الآخرون في أسئلتهم حول stackoverflow مثل "ساعدني في فهم رمز ثنائي / ساس / آلة" ، واستخدام البرامج التعليمية (بما في ذلك باللغة الصينية) والموارد الأخرى. لذلك ، يتم تخزين رقم العملية الرئيسي في البتات 58-63 ، ولكن هناك أيضًا بتات إضافية 0-4 (يميزون التعليمات "I2F" ، "ISETP" ، "MOV32I") ، في مكان ما بدلاً من 0-2 (للإهمال ، 3- 4 بتات في تعليمات فارغة ، في المواصفات تم تمييزها على أنها "UNK").

بالنسبة للسجلات والأرقام الثابتة ، يمكنك تجربة أداة التفكيك من أجل العثور على جميع البتات التي تؤثر على إخراج التفريغ ، مثل تلك الموضوعة تحت المفسد. جميع الحقول التي تمكنت من العثور عليها موجودة في المواصفات على Github ، ملف CUDA.slaspec ، قسم الرمز المميز.

ثم تحتاج إلى التوصل إلى عناوين للسجلات ، ومرة أخرى على Github. هذا ضروري لأن على المستوى الجزئي ، تسجل سليغ كمتغيرات عالمية في الفضاء بنوع "register_space" ، ولكن منذ نظرًا لأن مساحتها ليست مميزة "لا يمكن الاستدلال عليها" (وبالتأكيد لا يمكن أن تكون كذلك) ، فإنها تصبح في المترجم إما المتغيرات المحلية (غالبًا ما تكون مع الإصلاح "Var" ، ولكن أحيانًا تكون البادئة "المحلية" مثل) أو المعلمات (" param_ "). لم تكن SP في متناول اليد أبدًا ، فهي مطلوبة رسميًا في الغالب لضمان عمل المترجم. هناك حاجة إلى جهاز كمبيوتر (شيء مثل IP من x86) للمضاهاة.

ثم هناك سجلات أصلية ، شيء يشبه الأعلام ، ولكن "غرض عام" أكثر من غرض محدد مسبقًا ، مثل الفائض ، (في) المساواة إلى الصفر ، إلخ.

ثم ، سجل قفل لنمذجة مجموعة من التعليمات ISCADD .CC و IMAD.HI ، لأن ينفذ أولها في تنفيذي عملية العد لنفسه وللثانية ، لتجنب نقل جزء من المجموع إلى 4 بايت العليا ، حيث هذا سوف يفسد فك التجميع. ولكن بعد ذلك تحتاج إلى قفل التسجيل التالي حتى تكتمل عملية IMAD.HI. شيء مشابه ، أي كان التناقض بين الوثائق الرسمية والناتج المتوقع لجهاز فك الشفرة موجودًا بالفعل في SPU لنفس Ghidra.

ثم هناك سجلات خاصة ، يتم تنفيذها حتى الآن من خلال cpool. في المستقبل ، أخطط لاستبدالها بأحرف محددة افتراضيًا لنوع من المساحة "المستدرجة". هذه هي نفس threadIdx و blockIdx.

ثم نربط المتغيرات بحقول res ، par0 ، par1 ، par2 ، res. ثم هناك جداول فرعية ، وبعدها - ما كان كل شيء - الجداول الرئيسية (الجذر) مع التعليمات الرئيسية.

هنا تحتاج إلى اتباع صارم لتنسيق "mnemonics-operands" ، ومع ذلك ، يتم منح إعفاء للمعدِّلات ، التي ، مع ذلك ، يجب إرفاقها بالذاكرة أو بالجزء مع المعاملات. لا يُسمح بأي تنسيقات أخرى ، حتى نفس Hexagon DSP asm سيتعين تكييفها مع هذه البنية ، والتي ، مع ذلك ، ليست صعبة للغاية.

ستكون الخطوة الأخيرة هي كتابة تنفيذ للحصول على تعليمات بلغة البرامج الثابتة Pcode. الشيء الوحيد الذي أود ملاحظته من المثال الأول هو تعليمات ISCADD .CC و IMAD.HI ، حيث يأخذ الأول منهم مؤشرًا لتسجيله وإلغاء الإشارة إليه كمؤشرات إلى 8 بايت بدلاً من 4. ويتم ذلك عن قصد من أجل التكيف بشكل أفضل مع المحلل اللغوي وسلوكه ، على الرغم مما هو مكتوب في وثائق نفيديا حول تحويل جزء من المبلغ.

في المثال الثاني ، من الأفضل التحقق من إعدادات أداة فك الشفرة مقابل النقش "تبسيط التنبؤ". النقطة هي أن المسندات هي شرط واحد ومتشابه لتعليمات مختلفة ، في الواقع ، لا شيء أكثر من "SIMD" المعروف أو التماثلية التالية. أولئك. إذا تم تعيين البت المسند ، يتم تنفيذ التعليمات ، علاوة على ذلك ، على التوالي.

تحتاج أيضًا إلى التعود على كتابة تنفيذ على الفور لكل تعليمة ، وليس فقط نموذجًا أوليًا (استذكار المُعامِلات) ، لأن هناك أيضًا أداة فك الشفرة ومحاكي ومحللات أخرى.

ولكن بشكل عام ، تعد كتابة تنفيذ في Pcode مهمة أسهل من كتابة قواعد اللغة لمفكك تشفير بايت. اتضح بسرعة لإصلاح التنفيذ لبعض التعليمات المعقدة من x86 (وليس فقط) ، وذلك بفضل لغة وسيطة مريحة للغاية ، وسط واحد واحد (محسن) ، واثنين من الخلفية (بشكل أساسي C ؛ كبديل - Java / C # ، أشبه بالأخيرة ، على سبيل المثال. (حيث يظهر الانتقال من وقت لآخر ، ولكن ليس فاصل مصنف).

في المقالات التالية ، قد تكون هناك أيضًا واجهات أمامية للغات المدارة مثل DXBC و SPIR-V ، وسوف يستخدمون الواجهة الخلفية Java / C #. ولكن حتى الآن ، يتم التخطيط فقط لرموز الجهاز ، لأنه تتطلب الرموز الفرعية نهجًا خاصًا. مساعدة Ghidra

مشروع : Pcode الزلاجة

كتابة مواصفات Nvidia Kepler (ثنائيات CUDA ، إصدار لغة sm_30) لـ Ghidra

More articles: