병렬 스레드 실행

Parallel Thread Execution

병렬 스레드 실행(PTX 또는[1] NVPTX)은 Nvidia의 CUDA 프로그래밍 환경에서 사용되는 하위 수준의 병렬 스레드 실행 가상 머신 및 명령 집합 아키텍처입니다.NVCC 컴파일러는 C++와 유사한 언어인 CUDA로 작성된 코드를 PTX 명령어(ASCII 텍스트로 표시되는 어셈블리 언어)로 변환하고 그래픽 드라이버에는 PTX 명령을 Nvidia GPU의 처리 코어로 실행할 수 있는 실행 가능한 바이너리[2] 코드로 변환하는 컴파일러가 포함되어 있습니다.또한 GNU 컴파일러 컬렉션은 OpenMP [3]오프로딩 컨텍스트에서 PTX 생성을 위한 기본 기능을 가지고 있습니다.인라인 PTX 어셈블리는 [4]CUDA에서 사용할 수 있습니다.

레지스터

PTX는 임의로 큰 레지스터 세트를 사용합니다.컴파일러로부터의 출력은 거의 순수한 단일 할당 형식이며, 연속된 행은 일반적으로 연속된 레지스터를 나타냅니다.프로그램은 폼 선언으로 시작합니다.

.reg .u32 %r< >335>; // 335를 선언하다32비트 정수 타입의 레지스터 %r0, %r1, ..., %r334

이것은 3개의 인수로 이루어진 어셈블리 언어이며, 거의 모든 명령에서 동작하는 데이터 유형(부호 및 너비)이 명시적으로 나열됩니다.레지스터 이름 앞에는 % 문자가 붙고 상수는 리터럴입니다.다음은 예를 들어 다음과 같습니다. 

슈루64 %rd14, %rd12, 32; // 부호 없는 64비트 정수를 %rd12에서 32개 위치로 오른쪽으로 이동하면 %rd14가 됩니다. cvt.u64.u32 %ndpairs(%ndpause), %r112; // 부호 없는 32비트 정수를 64비트로 변환합니다.

술어 레지스터가 있지만 셰이더 모델 1.0의 컴파일된 코드에서는 branch 명령어와 함께만 사용됩니다.조건부 브랜치는 다음과 같습니다.

@%페이지 14 브래지어 $label; // $label로 분기

그setp.cc.type명령은 적절한 유형의 두 레지스터를 비교한 결과로 술어 레지스터를 설정합니다.set명령, 여기서set.le.u32.u64 %r101, %rd12, %rd2832비트 레지스터를 설정합니다.%r101로.0xffffffff64비트 레지스터가%rd1264비트 레지스터 이하입니다.%rd28.그렇지않으면%r101로 설정되어 있다.0x00000000.

의사 레지스터를 나타내는 몇 가지 정의된 식별자가 있습니다.다른 것들 중에.%tid, %ntid, %ctaid,그리고.%nctaid각각 스레드 인덱스, 블록 치수, 블록 인덱스 및 그리드 [5]치수를 포함합니다.

상태 공간

로드(ld및 스토어(st) 명령어는 예를 들어 몇 가지 다른 상태 공간(메모리 뱅크) 중 하나를 참조합니다.ld.param. 8개의 스테이트 [5]공간이 있습니다.

  • .reg: 레지스터
  • .sreg: 특별한 읽기 전용 플랫폼 고유의 레지스터
  • .const: 공유 읽기 전용 메모리
  • .global: 글로벌 메모리, 모든 스레드로 공유
  • .local: 로컬 메모리, 각 스레드 전용
  • .param: 커널에 전달된 파라미터
  • .shared: 블록 내 스레드 간에 공유되는 메모리
  • .tex: 글로벌 텍스처 메모리(사용되지 않음)

공유 메모리는 폼의 시작 부분에서 행을 통해 PTX 파일에 선언됩니다.

.공유했습니다. 정렬 8 .b8 pbatch_cache[15744]; // 8바이트 경계에 맞춰 15,744바이트를 정의합니다.

PTX에서 커널을 작성하려면 CUDA 드라이버 API를 통해 PTX 모듈을 명시적으로 등록해야 합니다.일반적으로 CUDA 런타임 API 및 NVIDIA의 CUDA 컴파일러 nvcc를 사용하는 것보다 더 번거롭습니다.GPU Ocelot 프로젝트는 PTX 모듈을 CUDA 런타임 API 커널 호출과 함께 등록하기 위한 API를 제공했지만 GPU Ocelot은 더 이상 활발하게 유지 [6]보수되지 않습니다.

「 」를 참조해 주세요.

레퍼런스

  1. ^ "User Guide for NVPTX Back-end — LLVM 7 documentation". llvm.org.
  2. ^ "CUDA Binary Utilities". docs.nvidia.com. Retrieved 2019-10-19.{{cite web}}: CS1 maint :url-status (링크)
  3. ^ "nvptx". GCC Wiki.
  4. ^ "Inline PTX Assembly in CUDA". docs.nvidia.com. Retrieved 2019-11-03.
  5. ^ a b "PTX ISA Version 2.3" (PDF).
  6. ^ "GPUOCelot: A dynamic compilation framework for PTX". github.com.{{cite web}}: CS1 maint :url-status (링크)

외부 링크