فيرمي (البنية الدقيقة)

📅 آخر تحديث: ٢ أغسطس ٢٠٢٦
✍️ نُشر في: ٢ أغسطس ٢٠٢٦
📷 يحتوي على 4 صور
الموسوعة الحرة

صورة لإنريكو فيرمي، صاحب الفضل في ابتكار فن العمارة

فيرمي هو الاسم الرمزي لبنية معمارية دقيقة لوحدة معالجة الرسومات (GPU) طورتها شركة إنفيديا ، وطُرحت في الأسواق في أبريل 2010 كخليفة لبنية تسلا . كانت فيرمي البنية المعمارية الدقيقة الرئيسية المستخدمة في سلسلتي GeForce 400 و 500 . صُنعت جميع وحدات معالجة الرسومات فيرمي المكتبية بتقنية 40 نانومتر، بينما صُنعت وحدات معالجة الرسومات فيرمي المحمولة بتقنيتي 40 نانومتر و 28 نانومتر . تُعد فيرمي أقدم بنية معمارية دقيقة من إنفيديا تدعم واجهة برمجة تطبيقات العرض Direct3D 12 من مايكروسوفت، مستوى الميزة 11.

تبعت معالجات فيرمي معالجات كيبلر ، واستُخدمت جنبًا إلى جنب مع كيبلر في سلسلة GeForce 600 وسلسلة GeForce 700 وسلسلة GeForce 800 ، وفي الأخيرتين فقط في وحدات معالجة الرسومات المحمولة .

في سوق محطات العمل، وجدت فيرمي استخدامًا في سلسلة Quadro x000، ونماذج Quadro NVS، وفي وحدات الحوسبة Nvidia Tesla .

سُمّي هذا الطراز المعماري على اسم إنريكو فيرمي ، وهو فيزيائي إيطالي.

ملخص

Nvidia GeForce GTX 480 من سلسلة بطاقات الرسومات GeForce 400؛ الإصدار الأول الذي يتميز ببنية Fermi الدقيقة (GF100-375-A3).
الشكل 1. بنية Nvidia Fermi. الاصطلاح في الأشكال: البرتقالي - الجدولة والإرسال؛ الأخضر - التنفيذ؛ الأزرق الفاتح - السجلات والذاكرة المؤقتة.
صورة لرقاقة معالج الرسوميات GF100 الموجودة داخل بطاقات GeForce GTX 470

تتميز وحدات معالجة الرسومات فيرمي ( GPUs ) بـ 3.0 مليار ترانزستور، وتم رسم مخطط لها في الشكل 1.

  • المعالج المتعدد المتدفق (SM): يتكون من 32 نواة CUDA (انظر قسمي المعالج المتعدد المتدفق ونواة CUDA).
  • جدولة GigaThread العالمية: تقوم بتوزيع كتل الخيوط على جدولة خيوط SM وتدير تبديلات السياق بين الخيوط أثناء التنفيذ (انظر قسم جدولة Warp).
  • واجهة المضيف: تربط وحدة معالجة الرسومات بوحدة المعالجة المركزية عبر ناقل PCI-Express v2 (معدل نقل ذروة يبلغ 8  جيجابايت/ثانية).
  • ذاكرة الوصول العشوائي الديناميكية (DRAM): تدعم ما يصل إلى 6  جيجابايت من ذاكرة GDDR5 DRAM بفضل إمكانية العنونة 64 بت (انظر قسم بنية الذاكرة).
  • تردد الساعة: 1.5  جيجاهرتز (لم يتم إصداره من قبل Nvidia، ولكن تم تقديره بواسطة Insight 64).
  • ذروة الأداء: 1.5 تيرافلوب.
  • تردد الساعة العالمي للذاكرة: 2  جيجاهرتز.
  • عرض نطاق ذاكرة الوصول العشوائي الديناميكية : 192  جيجابايت/ثانية.
  • دعم فك تشفير H.264 FHD.

معالج متعدد البث

تحتوي كل وحدة معالجة متعددة (SM) على 32 نواة CUDA أحادية الدقة، و16 وحدة تحميل/تخزين، وأربع وحدات وظائف خاصة (SFUs)،  وكتلة ذاكرة عالية السرعة على الشريحة بحجم 64 كيلوبايت (انظر القسم الفرعي L1+Shared Memory) وواجهة لذاكرة التخزين المؤقت L2 (انظر القسم الفرعي L2 Cache).

تحميل/تخزين الوحدات

السماح بحساب عناوين المصدر والوجهة لـ 16 خيطًا لكل دورة ساعة. تحميل وتخزين البيانات من/إلى ذاكرة التخزين المؤقت أو ذاكرة الوصول العشوائي الديناميكية (DRAM) .

وحدات الوظائف الخاصة (SFUs)

تُنفّذ التعليمات المتسامية مثل الجيب وجيب التمام والمقلوب والجذر التربيعي. تُنفّذ كل وحدة معالجة فرعية (SFU) تعليمة واحدة لكل خيط، في كل دورة ساعة؛ بينما تُنفّذ مجموعة الخيوط على مدى ثماني دورات ساعة. يتم فصل خط أنابيب وحدة المعالجة الفرعية (SFU) عن وحدة الإرسال، مما يسمح لوحدة الإرسال بإصدار التعليمات إلى وحدات التنفيذ الأخرى أثناء انشغال وحدة المعالجة الفرعية (SFU).

نواة CUDA

وحدة الحساب والمنطق للأعداد الصحيحة (ALU)

يدعم البرنامج دقة 32 بت كاملة لجميع التعليمات، بما يتوافق مع متطلبات لغات البرمجة القياسية. كما أنه مُحسَّن لدعم 64 بت بكفاءة في نماذج محطات العمل والخوادم، ولكنه مُقيد بشكل مصطنع في إصدارات المستخدمين العاديين.

وحدة الفاصلة العائمة (FPU)

يطبق معيار IEEE 754-2008 الجديد للفاصلة العائمة، موفراً تعليمة الضرب والجمع المدمجة (FMA) لكل من العمليات الحسابية أحادية وثنائية الدقة. يمكن تنفيذ ما يصل إلى 16 عملية ضرب وجمع مدمجة ثنائية الدقة لكل وحدة معالجة متعددة (SM) في كل دورة ساعة. [ 1 ]

عملية الضرب والجمع المدمجة

تُجري عملية الضرب والجمع المدمجة (FMA) عمليتي الضرب والجمع (مثل A*B+C) بخطوة تقريب نهائية واحدة، دون فقدان الدقة في عملية الجمع. وتُعدّ FMA أكثر دقة من إجراء العمليات بشكل منفصل.

جدولة الالتواء

تستخدم بنية فيرمي مُجدول خيوط موزع ثنائي المستوى .

يمكن لكل وحدة معالجة متعددة (SM) إصدار تعليمات تستهلك أي اثنين من أعمدة التنفيذ الخضراء الأربعة الموضحة في الشكل التخطيطي 1. على سبيل المثال، يمكن لوحدة المعالجة المتعددة (SM) أن تمزج 16 عملية من النوى الـ 16 في العمود الأول مع 16 عملية من النوى الـ 16 في العمود الثاني، أو 16 عملية من وحدات التحميل/التخزين مع أربع عمليات من وحدات SFU، أو أي تركيبات أخرى يحددها البرنامج.

تتطلب عمليات الفاصلة العائمة ذات 64 بت كلاً من أول عمودين للتنفيذ، لذا فهي تعمل بنصف سرعة عمليات 32 بت.

مُجدول الالتفاف المزدوج

على مستوى المعالج المتعدد (SM)، يقوم كل مُجدول حزم بتوزيع حزم من 32 خيطًا على وحدات التنفيذ الخاصة به. يحتوي كل معالج متعدد على مُجدولَي حزم ووحدتَي إرسال تعليمات، مما يسمح بإصدار حزمتين وتنفيذهما في وقت واحد. يختار مُجدول الحزم المزدوج حزمتين، ويُصدر تعليمة واحدة من كل حزمة إلى مجموعة من 16 نواة، أو 16 وحدة تحميل/تخزين، أو 4 وحدات SFU. يمكن إصدار معظم التعليمات بشكل مزدوج؛ حيث يمكن إصدار تعليمتين للأعداد الصحيحة، أو تعليمتين للأعداد العشرية، أو مزيج من تعليمات الأعداد الصحيحة والعشرية والتحميل والتخزين ووحدات SFU في وقت واحد. لا تدعم تعليمات الدقة المزدوجة الإرسال المزدوج مع أي عملية أخرى.

أداء

تُحسب القدرة النظرية للمعالجة أحادية الدقة لوحدة معالجة الرسومات Fermi بوحدة GFLOPS كالتالي: 2 (عدد العمليات لكل تعليمة FMA لكل نواة CUDA في الدورة الواحدة) × عدد نوى CUDA × سرعة ساعة التظليل (بالجيجاهرتز). تجدر الإشارة إلى أن الجيل السابق Tesla كان قادرًا على إصدار تعليمات MAD+MUL بشكل مزدوج إلى نوى CUDA ووحدات SFU بالتوازي، لكن Fermi فقد هذه القدرة لأنه لا يستطيع إصدار سوى 32 تعليمة في الدورة الواحدة لكل وحدة معالجة متعددة، مما يُبقي نوى CUDA الـ 32 فقط مُستغلة بالكامل. [ 2 ] لذلك، لا يُمكن الاستفادة من وحدات SFU للوصول إلى أكثر من عمليتين لكل نواة CUDA في الدورة الواحدة.

تبلغ القدرة النظرية لمعالجة البيانات بدقة مزدوجة لوحدة معالجة الرسومات Fermi نصف أداء الدقة المفردة على GF100/110. ومع ذلك، عمليًا، لا تتوفر هذه القدرة بدقة مزدوجة إلا في بطاقات Quadro و Tesla الاحترافية ، بينما تُحدَّد قدرة بطاقات GeForce المخصصة للمستهلكين بشكل مصطنع إلى ثمنها. [ 3 ]

ذاكرة

ذاكرة التخزين المؤقت L1 لكل وحدة معالجة متعددة وذاكرة التخزين المؤقت L2 الموحدة التي تخدم جميع العمليات (التحميل والتخزين والنسيج).

السجلات

تحتوي كل وحدة معالجة متعددة (SM) على 32 كيلوبايت من السجلات ذات 32 بت. لكل خيط إمكانية الوصول إلى سجلاته الخاصة فقط، وليس سجلات الخيوط الأخرى. الحد الأقصى لعدد السجلات التي يمكن أن يستخدمها نواة CUDA هو 63 سجلاً. يتناقص عدد السجلات المتاحة تدريجيًا من 63 إلى 21 سجلاً مع ازدياد عبء العمل (وبالتالي متطلبات الموارد) بزيادة عدد الخيوط. تتميز السجلات بنطاق ترددي عالٍ جدًا: حوالي 8000  جيجابايت/ثانية.

ذاكرة L1+ المشتركة

ذاكرة مدمجة يمكن استخدامها إما لتخزين البيانات مؤقتًا للخيوط الفردية (تخزين السجلات/ذاكرة التخزين المؤقت من المستوى الأول) أو لمشاركة البيانات بين عدة خيوط (ذاكرة مشتركة).  يمكن تهيئة هذه الذاكرة بسعة 64 كيلوبايت إما  كذاكرة مشتركة بسعة 48 كيلوبايت مع ذاكرة  تخزين مؤقت من المستوى الأول بسعة 16 كيلوبايت، أو  كذاكرة مشتركة بسعة 16 كيلوبايت مع  ذاكرة تخزين مؤقت من المستوى الأول بسعة 48 كيلوبايت. تُمكّن الذاكرة المشتركة الخيوط داخل نفس كتلة الخيوط من التعاون، وتُسهّل إعادة استخدام البيانات المدمجة على نطاق واسع، وتقلل بشكل كبير من حركة البيانات الخارجية. يمكن للخيوط في نفس كتلة الخيوط الوصول إلى الذاكرة المشتركة. توفر هذه الذاكرة وصولًا منخفض التأخير (10-20 دورة) ونطاقًا تردديًا عاليًا جدًا (1600  جيجابايت/ثانية) لكميات معتدلة من البيانات (مثل النتائج الوسيطة في سلسلة من العمليات الحسابية، أو صف أو عمود واحد من البيانات لعمليات المصفوفات، أو سطر من الفيديو، إلخ). يقول ديفيد باترسون إن هذه الذاكرة المشتركة تستخدم فكرة لوحة التخزين المؤقتة المحلية [ 4 ].

الذاكرة المحلية

تُستخدم الذاكرة المحلية كموقع تخزين لحفظ السجلات "المُتجاوزة". يحدث تجاوز السجلات عندما يتطلب جزء من التعليمات البرمجية مساحة تخزين أكبر من المتاحة على وحدة المعالجة المتعددة (SM). تُستخدم الذاكرة المحلية فقط لبعض المتغيرات التلقائية (التي يتم تعريفها في كود الجهاز دون استخدام أي من المحددات __device__ أو __shared__ أو __constant__ ). بشكل عام، يوجد المتغير التلقائي في سجل باستثناء ما يلي: (1) المصفوفات التي لا يستطيع المُصرّف تحديد أنها مُفهرسة بقيم ثابتة؛ (2) الهياكل أو المصفوفات الكبيرة التي تستهلك مساحة كبيرة جدًا من السجلات؛ أي متغير يقرر المُصرّف نقله إلى الذاكرة المحلية عندما تستخدم النواة عددًا من السجلات يفوق المتاح على وحدة المعالجة المتعددة (SM).

ذاكرة التخزين المؤقت من المستوى الثاني

ذاكرة تخزين مؤقتة موحدة من المستوى الثاني (L2) بسعة 768 كيلوبايت، مشتركة بين 16 وحدة معالجة متعددة (SMs)، تُعنى بجميع عمليات التحميل والتخزين من وإلى الذاكرة العامة، بما في ذلك النسخ من وإلى وحدة المعالجة المركزية، بالإضافة إلى طلبات معالجة الصور. كما يُنفذ نظام ذاكرة التخزين المؤقتة من المستوى الثاني (L2) عمليات ذرية، تُستخدم لإدارة الوصول إلى البيانات التي يجب مشاركتها بين مجموعات الخيوط أو حتى النواة.

الذاكرة العالمية

يمكن الوصول إلى الذاكرة العامة (VRAM) من قبل جميع الخيوط مباشرةً، وكذلك من قبل النظام المضيف عبر ناقل PCIe. وتتميز بزمن استجابة عالٍ يتراوح بين 400 و800 دورة.

فك ضغط الفيديو / ضغط الفيديو

انظر إلى Nvidia NVDEC (المعروف سابقًا باسم NVCUVID) وكذلك Nvidia PureVideo .

لم تكن تقنية Nvidia NVENC متاحة بعد، ولكن تم تقديمها في النظام اللاحق، Kepler .

رقائق فيرمي

  • GF100
  • GF104
  • GF106
  • GF108
  • GF110
  • GF114
  • GF116
  • GF117
  • GF119

انظر أيضاً

مراجع

  1. ↑ "معمارية الحوسبة CUDA من الجيل التالي من NVIDIA: فيرمي" (ملف PDF) . 2009. تم الاطلاع عليه في 7 ديسمبر 2015 .
  2. ↑ غلاسكوفسكي، بيتر ن. (سبتمبر 2009). "فيرمي من إنفيديا: أول بنية حوسبة كاملة لوحدة معالجة الرسومات" (ملف PDF) . ص 22. تم الاطلاع عليه في 6 ديسمبر 2015. يمكن إرسال ما مجموعه 32 تعليمة من مجموعة واحدة أو مجموعتين من الالتفافات في كل دورة إلى أي مجموعتين من مجموعات التنفيذ الأربع داخل وحدة معالجة فيرمي متعددة الوحدات . 
  3. ↑ سميث، رايان (26 مارس 2010). "بطاقات NVIDIA GeForce GTX 480 وGTX 470: بعد ستة أشهر من التأخير، هل كان الانتظار يستحق كل هذا العناء؟" . AnandTech . ص 6. مؤرشف من الأصل في 2 أبريل 2010. تم الاطلاع عليه في 6 ديسمبر 2015. يُذكر أن أداء سلسلة GTX 400 في معالجة FP64 محدود بنسبة 1/8 (12.5%) من أدائها في معالجة FP32، مقارنةً بما يمكن أن يقدمه الجهاز أصلاً وهو 1/2 (50%) من أداء FP32. 
  4. ↑ باترسون، ديفيد (30 سبتمبر 2009). "أهم 10 ابتكارات في بنية NVIDIA Fermi الجديدة، وأهم 3 تحديات قادمة" (ملف PDF) . مختبر أبحاث الحوسبة المتوازية وشركة NVIDIA . تم الاطلاع عليه في 3 أكتوبر 2013 .

عام