CUDA PTX - 2 : Inline PTX Assembly

이전 포스터 https://computing-jhson.tistory.com/15

CUDA PTX - 1 : Introduction

 

Inline PTX Assembly

저번 포스터에서 간단하게 PTX에 대해 소개하였다면, 이번 포스터에서는 NVIDIA Inline PTX Assembly 공식문서[1]를 정리해보고자 한다.

NVIDIA는 PTX ISA로 작성된 코드를 이용해 GPU data-parallel computing이 가능하도록 CUDA 프로그래밍 환경을 제공한다.  CUDA 코드에 Assembler (ASM) statements를 추가하는 방법은 다음과 같이 asm("") 문법을 사용한다. asm()을 통해 어떠한 PTX 코드도 CUDA program에서 사용할 수 있다. 컴파일 시 asm("PTX code") 구문 바로 그 위치에 PTX code가 생성된다.

// PTX 예제
__device__ kernel(int i, int j, int k) {
    ...
    // asm("template-string" : "constraint"(output) : "constraint"(input));
    asm("add.s32 %0, %1, %2;" : = "=r"(i) : "r"(j), "r"(k));
    ...
}

 

 

Parameters

asm() 구문에 CUDA 변수를 input으로 넣거나 result로 받아올 수 있다. 다음 코드 템플릿과 같이 작성하면 된다.

// asm("template-string" : "constraint"(output register) : "constraint"(input register));
asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));
// 위 코드는 다음과 같은 assembly로 번역됨
// ld.s32 r1, [j];
// ld.s32 r2, [k];
// add.s32 r3, r1, r2;
// st.s32 [i], r3;

" : : " 문법을 이용해서 output operand와 multiple input operands를 전달할 수 있다. 이때 constraint를 작성해주어야 한다. template-string에서 %n은 operands의 리스트에서 n번째 operand를 가져오도록 한다. %0은 "=r"(i)를 의미한다.

 

 

Constraints

"=" modifier는 해당 register에 값이 쓰임을 의미한다. "r"은 해당 register의 type을 의미하며, "r"은 (unsigned) int32, "h"  (unsigned) int16, "l"  (unsigned) int64, "f" float32, "d"  float64를 나타낸다. 다음 예제는 int64 값을 float32로 type conversion하는 예시이다.

asm("cvt.f32.s64 %0, %1;" : "=f"(x) : "l"(y));
// 위 코드는 다음과 같이 assembly로 번역된다
// ld.s64 rd1, [y]; signed int64 load
// cvt.f32.s64 f1, rd1; signed int64를 float32로 전환
// st.f32 [x], f1; float32 값을 address x에 저장

 

 

Register 선언 및 Namespace Conflicts

임시 값을 저장하기 위한 Register는 다음과 같이 선언할 수 있다.

asm(".reg .u32 t1;") // unsinged int32 register 선언
asm(".reg .b64 addr;") // 64bytes register 선언. 주로 address를 저장
asm(".reg .pred %p;") // branch predictor 저장 (if 문에 사용)

// Insert {} to avoid namespace conflicts
asm("{"
	".reg .u32 t2;"
    "}");

이때 같은 이름을 가지는 register는 두 번 선언될 수 없다. 만약 while문에서 같은 이름을 가진 register를 계속 선언한다면 namespace conflict가 발생해 컴파일 되지 않는다. 따라서 이럴 경우 {}로 감싸서 separate scope를 만들어 줘야 한다.

 

 

Compilier Optimization

컴파일러는 asm() 구문이 output operand만을 변경시킬 수 있다고 믿는다(No side effects). 따라서 사용되지 않는 output operand를 포함한 asm()구문은 컴파일러에 의해 제거될 수 있다. 이러한 컴파일러 최적화를 멈추기 위해서는 volatile keyword를 사용해야 한다. 컴파일러의 메모리 최적화를 막기 위해서는 다음 예제와 같이 "memory" keyword를 추가하여야 한다.

asm volatile ("mov.u32 %0, %%clock;" : "=r"(x)); // stop unused variables optimization
asm ("st.u32 [%0], %1;" : "r"(p), "r"(x) :: "memory"); // stop memory optimization

 

 

대표적인 Error 예시

/* Error 1: register constraints는 하나만 명시하여야 한다 */
asm("add.s32 %0, %1, %2;" : "=r"(i) : "rf"(j), "r"(k)); // rf -> r

/* Error 2: Scalar type 변수만 register가 읽을 수 있다 */
int4 i4;
asm("add.s32 %0, %1, %2;" : "=r"(i4) : "r"(j), "r"(k)); // struct 변수 읽기 불가

/* Error 3: Operand constraints와 operand registers의 타입이 일치하여야 한다 */
char ci;
asm("add.s32 %0,%1,%2;":"=r"(ci):"r"(j),"r"(k)); // "ci"는 char 변수이나, "r"(ci)는 int32를 읽는다.

 

 

 

Reference

[1] https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html​