การประมวลผลแบบขนานด้วยเธรด ( PTXหรือNVPTX ^{[ 1 ]} ) เป็น เครื่องเสมือนและสถาปัตยกรรมชุดคำสั่งการประมวลผลแบบขนาน ระดับต่ำที่ใช้ใน สภาพแวดล้อมการเขียนโปรแกรม Compute Unified Device Architecture ( CUDA ) ของNvidia คอมไพเลอร์ Nvidia CUDA (NVCC) ที่ใช้ LLVM จะแปลงโค้ดที่เขียนด้วยOpenCL C และ CUDA C/ C++เป็นคำสั่ง PTX ( IL ) และไดรเวอร์กราฟิกจะมีคอมไพเลอร์ที่แปลงคำสั่ง PTX เป็นโค้ดไบนารีที่สามารถเรียกใช้งานได้ ซึ่งสามารถทำงานบนคอร์ประมวลผลของหน่วยประมวลผลกราฟิก (GPU) ของ 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 บางตัวจะถูกกำหนดโดยความสามารถในการประมวลผลของ 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 บิตเป็น 1 หากรีจิสเตอร์ 64 บิตมีค่าน้อยกว่าหรือเท่ากับรีจิสเตอร์ 64 บิตมิฉะนั้นจะถูกตั้งค่าเป็น0 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 ]; // กำหนดขนาด 15,744 ไบต์ โดยจัดเรียงให้ตรงกับขอบเขต 8 ไบต์

การเขียนเคอร์เนลใน PTX จำเป็นต้องลงทะเบียนโมดูล PTX อย่างชัดเจนผ่าน CUDA Driver API ซึ่งโดยทั่วไปจะยุ่งยากกว่าการใช้ CUDA Runtime API และคอมไพเลอร์ CUDA ของ Nvidia อย่าง nvcc โครงการ GPU Ocelot ได้จัดเตรียม API สำหรับลงทะเบียนโมดูล PTX ควบคู่ไปกับการเรียกใช้เคอร์เนล CUDA Runtime API แม้ว่า GPU Ocelot จะไม่ได้รับการบำรุงรักษาอย่างต่อเนื่องแล้วก็ตาม^{[ 8 ]}

• การแสดงผลระดับกลางแบบพกพามาตรฐาน (SPIR)
• ไบนารี CUDA (cubin) – ไบนารีแบบอ้วนชนิดหนึ่ง

• หน้า PTX ISA บน NVIDIA Developer Zone