ترشيد ما يجري في نواة Opencl البسيطة الخاصة بي فيما يتعلق بالذاكرة العالمية

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

kernel أعلاه هي إضافة ناقلات القيام بها عشر مرات لكل حلقة. لقد استخدمت دليل البرمجة ومكدس الفائض لمعرفة كيفية عمل الذاكرة العالمية ، لكن لا يزال لا يمكنني معرفة ذلك من خلال النظر إلى الكود الخاص بي إذا كنت أقوم بالوصول إلى الذاكرة العالمية بطريقة جيدة. أنا أصل إلى ذلك بطريقة متجاورة وأخمن بطريقة محاذاة. هل تقوم البطاقة بتحميل قطع 128 كيلو بايت من الذاكرة العالمية للمصفوفات A و B و C؟ هل يقوم بعد ذلك بتحميل قطع 128 كيلو بايت لكل صفيف مرة واحدة لكل 32 فهرسة GID التي تمت معالجتها؟ (4*32 = 128) يبدو أنه بعد ذلك ، لا أضيع أي عرض ترددي عالمي للذاكرة؟

راجع للشغل ، يعرض Profiler Compute كفاءة GLD و GST 1.00003 ، والتي تبدو غريبة ، اعتقدت أنه سيكون 1.0 فقط إذا تم تجميع جميع المتاجر والأحمال. كيف هو أعلى من 1.0؟

Answer 1

نعم نمط وصول الذاكرة الخاص بك هو الأمثل إلى حد كبير. يصل كل نصف النهر إلى 16 كلمة متتالية 32 بت. علاوة على ذلك ، يتم محاذاة الوصول إلى 64Byte ، نظرًا لأن المخازن المؤقتة نفسها محاذاة و startIndex لكل نصف النهر هو مضاعف 16. لذا فإن كل نصف النهر سوف يولد معاملة 64byte واحدة. لذلك يجب ألا تضيع عرض النطاق الترددي للذاكرة من خلال الوصول غير المهمل.

نظرًا لأنك طرحت أمثلة في سؤالك الأخير ، يتيح تعديل هذا الرمز لنمط الوصول الأمثل (لأن الحلقة لا تفعل أي شيء حقًا سأتجاهل ذلك):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

في البداية ، يتيح لـ SE كيف يعمل هذا على أجهزة Compute 1.3 (GT200)

بالنسبة للكتابة إلى A ، فإن هذا سيولد نمطًا غير متوقع قليلاً (باتباع الأدوات نصف التي تم تحديدها بواسطة نطاق الهوية ونمط الوصول المقابل):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

لذا ، فإننا نضيع حوالي نصف عرض النطاق الترددي (كلما كان عرض الوصول الذي تضاعف بعد ذلك لمضاعفة النصف الفردي لا يساعد كثيرًا لأنه يولد المزيد من الوصول ، وهو ما لا يضيع المزيد من البايتات إذا جاز التعبير).

بالنسبة للقراءات من B ، تصل وصول المواضيع إلى عناصر الصفيف فقط ، لذلك لكل إمكانية الوصول إلى كل نصف وصول جميع الوصول في كتلة محاذاة 128Byte (العنصر الأول هو عند حدود 128B ، لأن هذا العنصر هو مضاعف GID 16 => الفهرس هو مضاعف 32 ، لعناصر بايت 4 ، وهذا يعني أن إزاحة العنوان هو مضاعف 128 ب). يمتد AccessPattern على كتلة 128B بأكملها ، لذلك سيقوم هذا بنقل 128B لكل نصف Garp ، مرة أخرى في الخصر نصف عرض النطاق الترددي.

تنشئ القراءات من C واحدة من أسوأ سيناريوهات الحالات ، حيث يشير كل مؤشر ترابط في كتلة 128B الخاصة به ، لذلك يحتاج كل مؤشر ترابط إلى نقله الخاص ، وهو يد واحدة من سيناريو التسلسل (على الرغم من أنه ليس سيئًا تمامًا مثل الطبيعة الطبيعية ، نظرًا لأن الجهاز يجب أن يكون قادرًا على التداخل في عمليات النقل). ما هو الأسوأ هو حقيقة أن هذا سوف ينقل كتلة 32B لكل موضوع ، تضيع 7/8 من عرض النطاق الترددي (نصل إلى 4B/مؤشر الترابط ، 32B/4B = 8 ، لذلك يتم استخدام 1/8 فقط من عرض النطاق الترددي). نظرًا لأن هذا هو Accesspattern من Matrixtransposes الساذجة ، فمن المستحسن للغاية القيام بأولئك الذين يستخدمون الذاكرة المحلية (التحدث من التجربة).

حساب 1.0 (G80)

هنا النمط الوحيد الذي سيخلق وصولًا جيدًا هو الأصل ، فإن جميع الأنماط في المثال سوف تخلق وصولًا غير مصقول تمامًا ، مما يضيع 7/8 من عرض النطاق الترددي (32B نقل/مؤشر ترابط ، انظر أعلاه). بالنسبة لأجهزة G80 ، فإن كل وصول حيث لا يقوم الخيط التاسع في HANDWARP بالوصول إلى العنصر التاسع يخلق مثل هذه الوصول غير المصقول

حساب 2.0 (فيرمي)

هنا ينشئ كل وصول إلى الذاكرة 128B معاملة (كما هو ضروري لجمع جميع البيانات ، لذلك 16x128b في أسوأ الحالات) ، ولكن يتم تخزينها مؤقتًا ، مما يجعلها أقل وضوحًا حيث سيتم نقل البيانات. في الوقت الحالي ، دعنا نفترض أن ذاكرة التخزين المؤقت كبيرة بما يكفي لعقد جميع البيانات وليس هناك تعارضات ، لذلك سيتم نقل كل 128 ب cacheline مرة واحدة على الأكثر. دعنا نفترض تنفيذًا متسلسلًا للنصف ، لذلك لدينا احتلال ذاكرة التخزين المؤقت الحتمية.

ستظل الوصول إلى B دائمًا تنقل 128B كتل (لا توجد مؤشرات مؤشرات ترابط أخرى في MemoryRea corending). سيؤدي الوصول إلى C إلى إنشاء 128B نقل لكل موضوع (أسوأ AccessPattern ممكن).

للوصول إلى A هو ما يلي (تعاملهم مثل القراءات في الوقت الحالي):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

لذلك بالنسبة للصفائف الكبيرة ، فإن الوصول إلى A Will لا تضيع أي عرض النطاق الترددي تقريبًا من الناحية النظرية. على سبيل المثال ، فإن الواقع ليس بالطبع جيدًا ، لأن الوصول إلى C سوف يقوم بمضايقة ذاكرة التخزين المؤقت بشكل جيد

بالنسبة إلى البروفيلر ، أفترض أن الكفاءات التي تزيد عن 1.0 هي ببساطة نتائج لعدم دقة النقطة العائمة.

امل ان يساعد