تنفيذ الخيوط المتوازية

يُعدّ تنفيذ الخيوط المتوازية ( PTX أو NVPTX ^{[ 1 ]} ) بنيةً منخفضة المستوى لآلة افتراضية ومجموعة تعليمات لتنفيذ الخيوط المتوازية ، تُستخدم في بيئة برمجة CUDA (بنية الحوسبة الموحدة للأجهزة) من Nvidia . يقوم مُصرّف Nvidia CUDA (NVCC) القائم على LLVM بترجمة التعليمات البرمجية المكتوبة بلغة OpenCL C وCUDA C/ C++ إلى تعليمات PTX (لغة وسيطة )، ويحتوي برنامج تشغيل الرسومات على مُصرّف يُترجم تعليمات PTX إلى تعليمات ثنائية قابلة للتنفيذ، والتي يمكن تشغيلها على أنوية معالجة وحدات معالجة الرسومات (GPUs) من Nvidia . ^{[ 2 ]} يمكن استخدام تجميع PTX المضمن في CUDA وOpenCL. ^{[ 3 ]}

يتمتع LLVM مع clang أيضًا بالقدرة على إنشاء ملفات PTX بناءً على توجيهات CUDA أو OpenCL C/C++ أو SYCL C++ أو OpenACC أو OpenMP . ^{[ 1 ] [ 4 ]} كما تقوم مجموعة مُجمِّعات GNU بإنشاء ملفات PTX لتفريغ العمليات بناءً على توجيهات OpenACC أو OpenMP. ^{[ 5 ]}

يتم تحديد مجموعة تعليمات PTX التي تدعمها وحدة معالجة الرسومات (GPU) معينة من خلال قدرتها الحسابية . ^{[ 1 ]}

يستخدم PTX مجموعة سجلات معالج كبيرة بشكل تعسفي ؛ ويكون ناتج المُصرّف عبارة عن شكل ثابت تقريبًا لتعيين قيمة واحدة ، حيث تشير الأسطر المتتالية عمومًا إلى سجلات متتالية. تبدأ البرامج بتصريحات على الشكل التالي:

.reg .u32 %r < 335 >; // تعريف 335 سجلًا %r0، %r1، ...، %r334 من نوع عدد صحيح غير مُوقّع 32 بت

هي لغة تجميع بثلاثة وسائط ، وتُحدد جميع التعليمات تقريبًا نوع البيانات (من حيث الإشارة والعرض) التي تعمل عليها بشكل صريح. تُسبق أسماء السجلات برمز  النسبة المئوية (%)، والثوابت تكون قيمًا حرفية، على سبيل المثال:

shr .u64 %rd14 , %rd12 , 32 ; // إزاحة عدد صحيح غير مُوقّع 64 بت من %rd12 بمقدار 32 خانة إلى اليمين، ليصبح الناتج %rd14 cvt .u64.u32 %rd142 , %r112 ; // تحويل عدد صحيح غير مُوقّع 32 بت إلى 64 بت

توجد سجلات المسند، لكن الكود المُجمَّع في نموذج التظليل 1.0 يستخدمها فقط بالتزامن مع أوامر التفرع؛ التفرع الشرطي هو

@ %p14 bra $label ; // الانتقال إلى $label

تُعيّن هذه setp.cc.typeالتعليمة سجلًا شرطيًا إلى نتيجة مقارنة سجلين من النوع المناسب، وهناك أيضًا setتعليمة أخرى تُعيّن السجل ذي 32 بت إلى قيمة معينة إذا كان السجل ذو 64 بت أقل من أو يساوي قيمة السجل ذي 64 بت . وإلا، تُعيّن قيمة السجل إلى قيمة أخرى .set.le.u32.u64%r101,%rd12,%rd28%r1010xffffffff%rd12%rd28%r1010x00000000

توجد بعض المعرفات المحددة مسبقًا التي تشير إلى السجلات الزائفة. من بينها، %tid, %ntid, %ctaidتحتوي %nctaidعلى مؤشرات الخيوط، وأبعاد الكتل، ومؤشرات الكتل، وأبعاد الشبكة، على التوالي. ^{[ 6 ]}

تشير أوامر التحميل ( ld) والتخزين ( st) إلى واحدة من عدة مساحات حالة متميزة (بنوك ذاكرة)، على سبيل المثال ld.param. هناك ثماني مساحات حالة: ^{[ 6 ] [ 7 ]}

.reg

السجلات

.sreg

سجلات خاصة، للقراءة فقط، خاصة بالمنصة

.const

ذاكرة مشتركة للقراءة فقط

.global

الذاكرة العامة، مشتركة بين جميع الخيوط

.local

ذاكرة محلية، خاصة بكل خيط

.param

المعلمات التي تم تمريرها إلى النواة

.shared

الذاكرة المشتركة بين الخيوط في كتلة واحدة

.tex

ذاكرة النسيج العامة (مهملة)

يتم تعريف الذاكرة المشتركة في ملف PTX عبر أسطر في بداية النموذج:

.shared .align 8 .b8 pbatch_cache [ 15744 ]; // تعريف 15744 بايت، محاذية لحدود 8 بايت

تتطلب كتابة النوى بلغة PTX تسجيل وحدات PTX صراحةً عبر واجهة برمجة تطبيقات برنامج تشغيل CUDA، وهو أمرٌ عادةً ما يكون أكثر تعقيدًا من استخدام واجهة برمجة تطبيقات وقت تشغيل CUDA ومترجم CUDA من Nvidia، nvcc. وقد وفّر مشروع GPU Ocelot واجهة برمجة تطبيقات لتسجيل وحدات PTX جنبًا إلى جنب مع استدعاءات نواة واجهة برمجة تطبيقات وقت تشغيل CUDA، على الرغم من أن مشروع GPU Ocelot لم يعد يُحدَّث بشكلٍ فعّال. ^{[ 8 ]}

• التمثيل الوسيط المحمول القياسي (SPIR)
• ثنائي CUDA (كوبين) - نوع من الثنائيات الدهنية

• صفحة PTX ISA على منطقة مطوري NVIDIA