كيف يمكنني استخدام الذاكرة المحلية في OpenCl؟

Question

لقد كنت ألعب مع OpenCl مؤخرًا ، وأنا قادر على كتابة نواة بسيطة تستخدم الذاكرة العالمية فقط. الآن أود البدء في استخدام الذاكرة المحلية ، لكن لا يمكنني معرفة كيفية استخدامها get_local_size() و get_local_id() لحساب "قطعة" من الإخراج في وقت واحد.

على سبيل المثال ، لنفترض أنني أردت تحويل kernel من Apple Hello Hello إلى شيء يستخدم الذاكرة المحلية. كيف يمكنك أن تفعل ذلك؟ هذا هو مصدر kernel الأصلي:

__kernel square(
    __global float *input,
    __global float *output,
    const unsigned int count)
{
    int i = get_global_id(0);
    if (i < count)
        output[i] = input[i] * input[i];
}

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

Answer 1

تحقق من العينات في Nvidia أو AMD SDKs ، يجب أن توجهك في الاتجاه الصحيح. ستستخدم Matrix Transpose الذاكرة المحلية على سبيل المثال.

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

__kernel square(
    __global float *input,
    __global float *output,
    __local float *temp,
    const unsigned int count)
{
    int gtid = get_global_id(0);
    int ltid = get_local_id(0);
    if (gtid < count)
    {
        temp[ltid] = input[gtid];
        // if the threads were reading data from other threads, then we would
        // want a barrier here to ensure the write completes before the read
        output[gtid] =  temp[ltid] * temp[ltid];
    }
}

Answer 2

هناك إمكانية أخرى للقيام بذلك ، إذا كان حجم الذاكرة المحلية ثابتًا. دون استخدام مؤشر في قائمة المعلمات kernels ، يمكن الإعلان عن المخزن المؤقت المحلي داخل النواة فقط بإعلانه __local:

__local float localBuffer[1024];

هذا يزيل الكود بسبب كميات أقل من ClsetKernelarg.

Answer 3

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

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