병렬 스레드 실행

병렬 스레드 실행(Parallel Thread Execution, PTX 또는 NVPTX)은 엔비디아의 CUDA(Compute Unified Device Architecture) 프로그래밍 환경에 사용되는 저수준 병렬 스레드 실행 가상 머신 및 명령어 집합이다. 엔비디아 쿠다 컴파일러(Nvidia CUDA Compiler, NVCC)는 C++와 유사한 언어인 CUDA로 작성된 코드를 PTX 명령어(IL)로 번역하고, 그래픽 드라이버에는 PTX 명령어를 실행 가능 바이너리 코드로 번역하는 컴파일러가 포함되어 있다. 이 코드는 엔비디아 그래픽 처리 장치(GPU)의 처리 코어에서 실행할 수 있다. GNU 컴파일러 모음 및 LLVM도 PTX를 생성할 수 있다. 인라인 PTX 어셈블리는 CUDA에서 사용할 수 있다.

 레지스터 
PTX는 임의로 큰 프로세서 레지스터 세트를 사용하며, 컴파일러의 출력은 거의 순수한 정적 단일 할당 형식으로, 연속된 줄은 일반적으로 연속된 레지스터를 참조한다. 프로그램은 다음과 같은 형식의 선언으로 시작한다.
.reg .u32 %r<335>;            // 부호 없는 32비트 정수 타입의 레지스터 %r0, %r1, ..., %r334 335개 선언

이것은 세 인자 어셈블리어이며, 거의 모든 명령어가 동작하는 데이터 타입(부호 및 너비)을 명시적으로 나열한다. 레지스터 이름 앞에는 % 문자가 붙고 상수는 리터럴이다. 예를 들어:
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 명령어 또한 있다. set.le.u32.u64 %r101, %rd12, %rd28는 64비트 레지스터 %rd12가 64비트 레지스터 %rd28보다 작거나 같으면 32비트 레지스터 %r101을 0xfff로 설정한다. 그렇지 않으면 %r101은 0x00000000으로 설정된다.

의사 레지스터를 나타내는 몇 가지 미리 정의된 식별자가 있다. 그 중 %tid, %ntid, %ctaid, 그리고 %nctaid는 각각 스레드 인덱스, 블록 차원, 블록 인덱스, 그리고 그리드 차원을 포함한다.

 상태 공간 
로드(ld) 및 저장(st) 명령은 여러 가지 다른 상태 공간(메모리 뱅크) 중 하나를 참조한다. 예를 들어 ld.param.
여덟 가지 상태 공간이 있다.
 .reg  레지스터
 .sreg  특수, 읽기 전용, 플랫폼별 레지스터
 .const  공유, 읽기 전용 메모리
 .global  전역 메모리, 모든 스레드에서 공유
 .local  지역 메모리, 각 스레드에 비공개
 .param  커널에 전달되는 매개변수
 .shared  한 블록 내 스레드 간 공유되는 메모리
 .tex  전역 텍스처 메모리 (더 이상 사용되지 않음)

공유 메모리는 PTX 파일에서 다음과 같은 형식의 줄로 선언된다.
.shared .align 8 .b8 pbatch_cache[15744]; // 8바이트 경계에 정렬된 15,744바이트 정의

PTX로 커널을 작성하려면 CUDA 드라이버 API를 통해 PTX 모듈을 명시적으로 등록해야 하는데, 이는 CUDA 런타임 API와 엔비디아의 CUDA 컴파일러인 nvcc를 사용하는 것보다 일반적으로 더 번거롭다. GPU Ocelot 프로젝트는 CUDA 런타임 API 커널 호출과 함께 PTX 모듈을 등록할 수 있는 API를 제공했지만, GPU Ocelot은 더 이상 활발히 유지 관리되지 않는다.

 같이 보기 
 SPIR (Standard Portable Intermediate Representation)
 CUDA 바이너리 (큐빈) – 일종의 팻 바이너리

 각주 

 외부 링크 
 엔비디아 개발자 존의 PTX ISA 페이지

분류:엔비디아
분류:명령어 집합 구조