병렬 스레드 실행

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

레지스터

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는 각각 스레드 인덱스, 블록 차원, 블록 인덱스, 그리고 그리드 차원을 포함한다.^[5]

상태 공간

로드(ld) 및 저장(st) 명령은 여러 가지 다른 상태 공간(메모리 뱅크) 중 하나를 참조한다. 예를 들어 ld.param. 여덟 가지 상태 공간이 있다.^[5]

.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은 더 이상 활발히 유지 관리되지 않는다.^[6]

같이 보기

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

각주

↑ ^가 ^나 “User Guide for NVPTX Back-end – LLVM 7 documentation”. 《llvm.org》.
↑ “CUDA Binary Utilities”. 《docs.nvidia.com》 (미국 영어). 2019년 10월 19일에 확인함.
↑ “nvptx”. 《GCC Wiki》.
↑ “Inline PTX Assembly in CUDA”. 《docs.nvidia.com》 (미국 영어). 2019년 11월 3일에 확인함.
↑ ^가 ^나 “PTX ISA Version 2.3” (PDF).
↑ “GPUOCelot: PTX를 위한 동적 컴파일 프레임워크”. 《github.com》. 2022년 11월 7일.

외부 링크

엔비디아 개발자 존의 PTX ISA 페이지

[:0-1] 가 ^나 “User Guide for NVPTX Back-end – LLVM 7 documentation”. 《llvm.org》.

[2] “CUDA Binary Utilities”. 《docs.nvidia.com》 (미국 영어). 2019년 10월 19일에 확인함.

[3] “nvptx”. 《GCC Wiki》.

[4] “Inline PTX Assembly in CUDA”. 《docs.nvidia.com》 (미국 영어). 2019년 11월 3일에 확인함.

[ptx-isa-5] 가 ^나 “PTX ISA Version 2.3” (PDF).

[6] “GPUOCelot: PTX를 위한 동적 컴파일 프레임워크”. 《github.com》. 2022년 11월 7일.

[1]

[2]

[3]

[4]

[5]

[6]