الماكرو CONFLICT_FREE_OFFSET المستخدم في خوارزمية البادئة المتوازية من GPU Gems 3

Question

في البداية هذا هو رابط الخوارزمية:

GPU Gems 3، الفصل 39: مجموع البادئة الموازية (المسح الضوئي) مع CUDA.

لتجنب تعارضات البنوك، تتم إضافة الحشو إلى مصفوفة الذاكرة المشتركة لكل NUM_BANKS (أي 32 للأجهزة ذات القدرة الحسابية 2.x).ويتم ذلك عن طريق (كما في الشكل 39-5):

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

لا أفهم كيف يعادل ai/NUM_BANKS الماكرو:

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))  

أليس يساوي

n >> LOG_NUM_BANKS

هو موضع تقدير أي مساعدة.شكرًا

Answer 1

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

إذا كنت تريد إجراء عمليات المسح بالطريقة السهلة، فاستخدم thrust::inclusive_scan أو thrust::exclusive_scan.

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

[1] شوبهاراتا سينجوبتا، مارك هاريس، مايكل جارلاند، وجون د.أوينز.“خوارزميات المسح المتوازي الفعالة لوحدات معالجة الرسومات متعددة النواة”.في جاكوب كورزاك، ديفيد أ.بدر وجاك دونغارا، محرران، الحوسبة العلمية مع النوى المتعددة والمسرعات، تشابمان وهول/العلوم الحسابية CRC، الفصل 19، الصفحات من 413 إلى 442.تايلور وفرانسيس، يناير 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1041

[2] ميريل، د.و جريمشو، أ.المسح المتوازي لبنيات الدفق.التقرير الفني CS2009-14، قسم علوم الحاسب، جامعة فيرجينيا.ديسمبر.2009.