Parallel Thread Execution ( PTX o NVPTX [1] ) es una máquina virtual de ejecución de subprocesos paralelos de bajo nivel y una arquitectura de conjunto de instrucciones que se utiliza en el entorno de programación Compute Unified Device Architecture ( CUDA ) de Nvidia . El compilador Nvidia CUDA (NVCC) traduce el código escrito en CUDA, un lenguaje similar a C++ , en instrucciones PTX (un lenguaje ensamblador representado como texto ASCII (Código estándar americano para el intercambio de información )), y el controlador de gráficos contiene un compilador que traduce las instrucciones PTX en código binario ejecutable, [2] que puede ejecutarse en los núcleos de procesamiento de las unidades de procesamiento gráfico (GPU) de Nvidia . La colección de compiladores GNU también tiene la capacidad básica de generar PTX en el contexto de la descarga de OpenMP . [3] El ensamblaje PTX en línea se puede utilizar en CUDA. [4]
PTX utiliza un conjunto de registros de procesador arbitrariamente grande ; la salida del compilador es casi puramente estática en forma de asignación única , con líneas consecutivas que generalmente hacen referencia a registros consecutivos. Los programas comienzan con declaraciones de la forma
.reg .u32 %r < 335 >; // declara 335 registros %r0, %r1, ..., %r334 de tipo entero de 32 bits sin signo Es un lenguaje ensamblador de tres argumentos y casi todas las instrucciones enumeran explícitamente el tipo de datos (en signo y ancho) sobre el que operan. Los nombres de registros están precedidos por un carácter % y las constantes son literales, por ejemplo:
shr .u64 %rd14 , %rd12 , 32 ; // desplaza a la derecha un entero sin signo de 64 bits desde %rd12 por 32 posiciones, resulta en %rd14 cvt .u64.u32 %rd142 , %r112 ; // convierte un entero sin signo de 32 bits a 64 bits Hay registros de predicado, pero el código compilado en el modelo de sombreado 1.0 los usa solo junto con comandos de bifurcación; la bifurcación condicional es
@ %p14 bra $label ; // rama a $label La setp.cc.typeinstrucción establece un registro de predicado en el resultado de comparar dos registros del tipo apropiado; también hay una setinstrucción, donde establece el registro de 32 bits en si el registro de 64 bits es menor o igual que el registro de 64 bits . De lo contrario, se establece en .set.le.u32.u64 %r101, %rd12, %rd28%r1010xffffffff%rd12%rd28%r1010x00000000
Existen algunos identificadores predefinidos que denotan pseudoregistros. Entre otros, %tid, %ntid, %ctaid, y %nctaidcontienen, respectivamente, índices de subprocesos, dimensiones de bloque, índices de bloque y dimensiones de cuadrícula. [5]
Los comandos load( ld) y store( st) hacen referencia a uno de varios espacios de estado distintos (bancos de memoria), por ejemplo ld.param, Hay ocho espacios de estado: [5]
.reg.sreg.const.global.local.param.shared.texLa memoria compartida se declara en el archivo PTX mediante líneas al comienzo del formulario:
.shared .align 8 .b8 pbatch_cache [ 15744 ]; // define 15 744 bytes, alineados con un límite de 8 bytes Para escribir núcleos en PTX es necesario registrar explícitamente los módulos PTX a través de la API del controlador CUDA, lo que suele ser más complicado que utilizar la API de tiempo de ejecución CUDA y el compilador CUDA de Nvidia, nvcc. El proyecto GPU Ocelot proporcionó una API para registrar módulos PTX junto con las invocaciones del núcleo de la API de tiempo de ejecución CUDA, aunque ya no se realiza un mantenimiento activo de GPU Ocelot. [6]