NVCC - أحجام كتلة مختلفة اعتمادًا على القوس في وقت الترجمة

Question

لدي نواة تعرض أعلى أداء ل مختلف قم بحظر الأحجام عند التشغيل على أجهزة Kepler و Fermi. أود ، في وقت الترجمة ، التحقق من هدف الهندسة المعمارية الحالية وتحديد أ THREADS_PER_BLOCK الماكرو إلى الأول) إطلاق النواة مع ؛ 2) تحديد عدد الكتل اللازمة ؛ 3) ضبط حجم الذاكرة المشتركة بشكل ثابت في النواة.

يوضح أدناه ما أحاول القيام به. لنفترض أنني أستهدف أجهزة GK104 ، وبالتالي استخدم nvcc -arch=sm_30. سيظل هذا يؤدي إلى THREADS_PER_BLOCK = 256 حيث __CUDA_ARCH__ لم يتم تعريفه لتجميع رمز المضيف. (أنا أفهم من EG هذا الجواب, ، لماذا لا يمكن أن تعمل بهذه الطريقة.)

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

يمكنني التحقق من خصائص الجهاز في وقت التشغيل واستخدام الذاكرة المشتركة الديناميكية ، ولكن أرغب في معرفة ما إذا كان يمكن ترميز هذا في وقت الترجمة دون الحاجة إلى إضافة dffermi أو -dkepler و settion يدويًا THREADS_PER_BLOCK بناء على ذلك. NB: أي مستخدمين لهذا الرمز سيقومون بتجميعه بأنفسهم ، تقريبيا بالتأكيد بالنسبة للهندسة المعمارية ، لذلك هذا ليس خيارًا غير معقول. يبدو أنه لا لزوم له في ضوء تمرير -arch= علَم.

Answer 1

nvcc لا يكتشف برنامج التحويل البرمجي وحدات معالجة الرسومات المتاحة محليًا ، فهو يستهدف دائمًا SM 1.0 افتراضيًا. وإلا فإنه يمكن أن يقدم بعض السلوك المربك تمامًا عند البناء على أنظمة مختلفة.

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