1. 개요
PTX(Parallel Thread Execution)는 NVIDIA가 개발한 GPU용 병렬 스레드 실행 가상 머신(Virtual Machine) 및 그 명령어 집합 아키텍처(ISA)이다.일반적인 CPU의 기계어와 달리, PTX는 특정 세대의 GPU 하드웨어에 종속되지 않는 하드웨어 추상화 계층(Hardware Abstraction Layer)의 역할을 수행한다. 컴파일러의 중간 표현(Intermediate Representation, IR)과 같은 수준으로 볼 수 있다.
C++ 등으로 작성된 CUDA 소스 코드는 nvcc 컴파일러를 통해 우선 PTX 코드로 변환된다. 그 후, 애플리케이션이 실행될 때 그래픽 드라이버 내의 JIT(Just-In-Time) 컴파일러가 장착된 그래픽 카드의 구체적인 아키텍처(마이크로아키텍처)에 맞춰 이를 실제 기계어인 SASS(Streaming Assembler)로 최종 변환하여 실행한다.
이러한 구조 덕분에, 개발자는 새로운 GPU가 출시될 때마다 프로그램을 다시 컴파일할 필요 없이 기존 PTX 코드를 그대로 사용하여 최신 하드웨어의 이점을 누릴 수 있다.
2. 역사
| PTX ISA 버전 | CUDA 연산 능력 | CUDA 버전 | 해당 GPU 칩셋 |
| 1.0 | CUDA 1.0 | G80 | |
| 1.1 | CUDA 1.1 | ||
| 1.2 | CUDA 2.0 | ||
| 1.3 | CUDA 2.1 | ||
| 1.4 | CUDA 2.2 | ||
| 2.0 | CUDA 3.0 | GF110 | |
| 2.3 | CUDA 4.0 | ||
| 3.0 | CUDA 4.1 | GK104 | |
| 4.0 | 5.0 | CUDA 6.0 | |
| 5.0 | 6.0 | CUDA 8.0 | |
| 6.0 | 7.0 | CUDA 9.0 | |
| 6.4 | 7.5 | CUDA 10.0 | |
| 7.0 | 8.0 | CUDA 11.0 | |
| 8.0 | 9.0 | CUDA 12.0 | |
| 8.5 | 9.0 | CUDA 12.5 | |
| 8.6 | 10.1 | CUDA 12.7 | |
| 8.7 | 12.0 | CUDA 12.8 | |
| 8.8 | 12.1 | CUDA 12.9 | |
| 9.0 | 12.1 | CUDA 13.0 | |
| 9.1 | 12.1 | CUDA 13.1 |
2.1. PTX 8.0 (Hopper)
wgmma명령어 추가- DPX (Dynamic Programming X) 명령어 추가
2.2. PTX 7.0 (Ampere)
cp.async명령어 추가mbarrier명령어 추가ld.global.nc명령어 추가
2.3. PTX 6.0 (Volta)
wmma명령어 추가barrier명령어 추가neg명령어의.f16및.f16x2자료형 지원fns명령어 추가bar.warp.sync명령어 추가match.sync명령어 추가brx.idx명령어 추가
2.4. PTX 5.0 (Pascal)
dp4a,dp2a명령어 추가 (4-way/2-way dot product with accumulation)- 특수 레지스터
%clock_hi지원
2.5. PTX 3.0 (Kepler)
mad.cc,madc명령어 추가shfl명령어 추가
2.6. PTX 2.0 (Fermi)
- 단정밀도(f32) 명령어가 subnormal number를 지원
- 단정밀도
add,sub,mul이.rm및.rprounding 수식어를 지원 - 단정밀도 fused multiply-add (FMA) 명령어
fma.f32추가 - 단정밀도 및 배정밀도
div,rcp,sqrt명령어에 IEEE 754 대응 rounding 추가 testp및copysign명령어 추가
- 1.x와의 호환성을 위해
.ftz 수식어를 지원- 이전 버전에서는 배정밀도(f64)만
.rm 및 .rp 지원- 기존
mad.f32 명령어는 호환성을 위해 유지3. 프로그래밍 모델
3.1. 메모리 공간 (State Spaces)
PTX는 변수가 저장되는 메모리 공간을 명시적으로 구분한다..reg: 레지스터. 가장 빠르지만 스레드 전용 공간이다..sreg: 특수 레지스터. 읽기 전용이며 시스템 상태(스레드 ID 등)를 저장한다..global: 전역 메모리. 모든 스레드가 접근 가능하며 가장 느리다(DRAM)..shared: 공유 메모리. 하나의 스레드 블록 내에서 공유되며 빠르다(On-chip)..local: 로컬 메모리. 스레드 전용이지만 레지스터에 들어가지 못하는 데이터가 저장된다(DRAM)..const: 상수 메모리. 읽기 전용이며 캐시된다..param: 커널 함수 파라미터 전달용 공간.
3.2. 자료형
| 기본 형식 | 형식 지정자 |
| 부호 있는 정수 | .s8, .s16, .s32, .s64 |
| 부호 없는 정수 | .u8, .u16, .u32, .u64 |
| 부동소수점 | .f16, .f32, .f64 |
| 비트 (자료형 없음) | .b8, .b16, .b32, .b64 |
| Predicate | .pred |
4. 명령어 목록
5. 기타
- 실제 하드웨어에서 실행되는 기계어인 SASS는 NVIDIA가 공식적으로 문서를 제공하지 않으며, 'cuobjdump' 등의 도구를 통해서만 확인 가능하다. PTX와 SASS는 1:1로 대응되지 않는다.