CUDA Instructions (1)

References

• Professional CUDA C Programming

Contents

• Floating-Point Instructions
• Intrinsic and Standard Functions
• Atomic Instructions

---

이번 포스팅에서 Instruction을 명령어로 지칭하도록 하겠습니다.

Instructions(명령어)은 프로세서에서 로직의 단일 단위를 나타냅니다. CUDA에서 작업하면서 직접적으로 명령어을 다루는 일은 드물지만, CUDA 커널 코드로부터 다른 명령어들이 언제 생성되고, 어떻게 고급 언어 기능들이 명령어로 변환되는지 이해하는 것은 중요합니다. 

기능이 같은 두 명령어 중의 하나를 선택하는 것은 성능, 정확성을 포함한 다양한 프로그램의 특성에 영향을 미칠 수 있습니다. 이는 CUDA를 엄격한 수치적 검증을 요구하는 레거시 어플리케이션에 포팅할 때 특히 중요합니다.

 

이번 포스팅에서는 CUDA 커널에 의해서 생성된 명령어에 영향을 미치는 세 가지 항목인, floating-point operation, instrinsic and standard function, 그리고 atomic operation에 대해 알아보겠습니다.

Floating-point(부동소수점) 계산은 비적분 값에 대해 동작하며 CUDA 프로그램의 정확성과 성능에 영향을 미칩니다. Instrinsic and Standard 함수는 동일한 기능의 수학적 연산을 구현하지만 정확도와 성능은 다릅니다.

Atomic operation은 여러 스레드에서 한 변수에 대한 작업을 동시에 수행할 때 정확성을 보장합니다.

 

---

Floating-Point Instructions

부동소수점 연산을 위한 IEEE 표준 754의 도입 이후에, NVIDIA를 포함한 모든 메이저 프로세서 제조사는 이 표준을 따릅니다. 이 표준은 이진 부동소수점 데이터를 아래와 같이 3개의 필드로 인코딩하도록 합니다. 

아래 포스팅에서 부동소수점에서 다룬 적이 있으니, 필요하시다면 참조하시길 바랍니다.

부동소수점 (Floating-Point)

부동소수점 (Floating-Point)

 

플랫폼 전체에서 일관된 연산을 보장하기 위해서 IEEE-754는 C 데이터 타입의 float와 double에 해당하는 32비트 및 64비트 부동소수점 포맷을 정의합니다. 이들은 아래와 같이 비트 길이가 다릅니다.

1비트의 sign s과 8비트의 exponent e, 23비트의 significand v로 구성된 32비트 부동소수점 변수는 다음과 같이 계산할 수 있습니다.

이 표준 포맷때문에 부동소수점 변수는 정수보다 더 미세한 단위로 값을 정확하게 나타낼 수 있다는 것을 이해하는 것이 중요합니다. 그러나 수치적 정확도는 제한적입니다. 부동소수점 포맷으로 정확하게 저장할 수 있는 값은 이산적이고 유한합니다. 예를 들어 다음의 코드를 살펴보겠습니다.

#include <stdio.h>

int main(int argc, char** argv)
{
    float a = 3.1415927f;
    float b = 3.1415928f;
    
    if (a == b) {
        printf("a is equal to b\n");
    }
    else {
        printf("a does not equal b\n");
    }

    return 0;
}

다음의 코드에서 a와 b의 마지막 숫자가 다르므로, 위 코드는 다음을 출력할 것이라고 추측됩니다.

'a does not equal b'

하지만, IEEE754 표준과 호환되는 아키텍처에서는 실제로 다음의 결과를 출력합니다.

'a is equal to b'

 

이 예제에서 부동소수점 변수 a와 b에 대입되는 값은 부동소수점에 의한 한정된 수의 비트로 인해 실제로는 저장될 수 없습니다. 결과적으로 이러한 값들은 저장될 수 있는 가장 가까운 값으로 반올림되며, 이로 인하여 두 값이 같다고 출력하게 됩니다.

 

올바르게 저장할 수 없는 부동소수점 값은 configurable rounding mode를 사용하여 표현 가능한 값으로 반올림됩니다. 예를 들어, 위에서 표현 불가능한 값을 가장 가까운 표현 가능한 값으로 반올림하는 기본 동작이 사용되었습니다. 반올림 모드의 다른 예로는 round-to-zero, round-up, round-down가 있습니다.

부동소수점 프로그래밍에서 또 다른 고려사항은 표현 가능한 부동소수점 값의 정밀도입니다. 아시다시피, 부동소수점 값은 정수 값보다 훨씬 미세한 값을 나타낼 수 있습니다. 하지만, 오직 이산 구간에 대해서만 값을 저장할 수 있습니다. 또한 부동소수점 값이 0에서 멀어질 수록, 표현 가능한 값들의 사이 간격도 증가합니다.

 

C 표준의 math 함수인 nextafterf를 사용하면 주어진 값에서 다음으로 높은 표현가능한 부동소수점 값을 찾을 수 있습니다. 아래 표는 몇 가지 갚에 대해 부동소수점 값과 다음으로 높은 부동소수점 값 간의 차이를 보여줍니다.

x값이 증가함에 따라서 정밀도(precision)가 크게 낮아지는 것을 볼 수 있습니다. 부동소수점 값 사이에서 이러한 큰 간격은 반올림 모드의 선택이 프로그램의 수치 결과에 극단적인 값이 발생하여 상당한 영향을 미칠 수 있습니다.

 

부동소수점 값으로 작동하는 명령어를 부동소수점 명령이라고 합니다. CUDA는 덧셈, 곱셈, 나눗셈, 뺄셈과 같은 부동소수점에 대한 모든 일반적인 산술 연산을 지원합니다.

IEEE754를 준수하는 CUDA와 기타 프로그래밍 모델은 부동소수점에서의 두 가지 정확도, 32비트와 64비트를 지원합니다. 이러한 포맷을 각각 단정밀도(single-precision), 배정밀도(double-precision)이라고 합니다. 배정밀도는 단정밀도보다 2배 더 많은 비트를 사용하기 때문에 더 많은 값을 정확하게 나타낼 수 있습니다. 즉, 배정밀도 값의 집합은 단정밀도보다 더 간격이 촘촘하고 범위가 넓습니다.

예를 들어, 방금 위에서 본 예시에서 float대신 double을 사용하면 정확한 값의 비교를 수행할 수 있습니다.

#include <stdio.h>

int main(int argc, char** argv)
{
    double a = 3.1415927;
    double b = 3.1415928;
    
    if (a == b) {
        printf("a is equal to b\n");
    }
    else {
        printf("a does not equal b\n");
    }

    return 0;
}

위 코드는

'a does not equal b'

를 출력할 것입니다.

 

---

Intrinsic and Standard Functions

단정밀도, 배정밀도 명령 이외에도 CUDA는 모든 산술 함수를 intrinsic 또는 standard 함수로 분류합니다. Standard Function은 host 및 device로부터 액세스할 수 있고, 표준화된 연산을 지원합니다. Standard 함수에는 sqrt, exp, sin과 같은 C 표준 math 라이브러리의 산술 연산이 포함됩니다. 곱셈이나 덧셈과 같은 단일 명령어 연산도 포함합니다.

 

CUDA의 intrinsic function은 오직 device 코드에서만 액세스할 수 있습니다. 프로그래밍에서 내재(intrinsic)되거나 또는 내장(build-in)된 함수는 컴파일러가 그 함수의 동작에 대한 특별한 지식을 가지고 있다는 것을 의미하며, 이는 보다 공격적인 최적화와 전문화된 명령어 생성을 가능하게 합니다. 이것이 CUDA intrinsic 함수입니다. 사실 많은 삼각 함수들은 그래픽 프로그램에서 많이 사용(변환, 회전 등을 수행하기 위해)되기 때문에 GPU의 하드웨어에서 직접 구현됩니다.

 

CUDA에서 많은 intrinsic 함수들은 standard 함수와 관련이 있는데, intrinsic 함수와 동일한 연산을 수행하는 standard 함수가 존재합니다. 예를 들어, 배정밀도 부동소수점 제곱근을 수행하기 위한 standard 함수는 sqrt입니다. 이와 동일한 기능을 수행하는 intrinsic 함수는 __dsqrt_rn 입니다.

 

intrinsic 함수는 동일한 standard  함수보다 더 적은 명령어로 수행됩니다. 결과적으로 intrinsic 함수는 동일한 기능의 standard 함수보다 빠르지만, 수치적으로 정확하지는 않습니다. Standard 함수와 intrinsic 함수를 서로 바꿔서 사용할 수 있지만, 성능 및 수치적 정확도 측면에서 다른 결과를 생성할 수 있습니다.

 

---

Atomic Instructions

Atomic Instruction(원자 연산)은 수학 연산을 수행하지만, 다른 스레드로부터 인터럽트가 없는 단일 연산을 수행합니다. 스레드가 성공적으로 atomic operation을 완료했을 때, 변수의 상태 변화가 완료되었는지 확인할 수 있습니다. Atomic operation은 여러 스레드가 서로 간섭하는 것을 방지하기 때문에 스레드에서 공유되는 데이터에 대해 read-modify-write 작업을 가능하게 합니다. GPU와 같이 동시성이 높은 환경에서는 read-modify-write 작업의 원자성을 보장하는 것이 특히 중요합니다. CUDA에서는 32비트 또는 64비트 global memory 또는 shared memory에 대한 read-modify-write atomic operation을 수행하는 atomic 함수를 지원합니다.

 

standard와 intrinsic 함수와 유사하게 각 atomic 함수는 덧셈, 곱셈, 뺄셈과 같은 기본 수학 연산을 제공합니다. 하지만 다른 명령어와는 다르게 경쟁하는 두 스레드가 공유하는 메모리 위치에서 동작할 때 정의된 동작을 가지고 있습니다.

다음의 간단한 커널을 살펴보겠습니다.

__global__
void incr(int *ptr)
{
	int temp = *ptr;
    temp = temp + 1;
    *ptr = temp;
}

이 커널은 memory location으로부터 read하고, 그 값에 1을 더한 후, 계산된 값을 다시 동일한 location으로 write합니다. 어떠한 스레드 ID도 액세스되는 memory location을 변경하는데 사용되지 않으므로, 커널의 모든 스레드는 동일한 주소에 read/write를 수행합니다.

만약 하나의 블록(32 threads)이 이 커널을 수행하면 어떻게 될까요?

32라고 생각할 수 있지만, 실제로 결과는 undefined 입니다.

문제는 동일한 메모리 위치에 하나 이상의 스레드가 writing을 하기 때문입니다. 이는 data race 또는 unsafe access to memory라고 합니다. data race는 둘 이상의 독립된 스레드의 실행이 같은 위치에 액세스하고 적어도 하나의 액세스가 그 값을 수정할 때 정의됩니다. 

 

다행히, atomic 명령어를 사용하면 이러한 동작을 피할 수 있습니다. atomic 명령어는 다음의 CUDA API를 통해 사용할 수 있습니다.

int atomicAdd(int* M, int V);

대부분 atomic 함수는 이진 함수이며, 2개의 피연산자에 대해 연산을 수행합니다. 이 함수들은 메모리 위치 M과 값 V를 입력으로 받습니다. atomic 함수와 관련된 연산은 V와 *M 메모리 위치에 이미 저장된 값에 대해 수행됩니다. 그리고 그 결과는 동일한 메모리에 다시 write 됩니다.

 

atomic 함수는 산술 함수(arithmetic functions), 비트 함수(bitwise functions), 스왑 함수(swap functions)로 나눌 수 있습니다. atomic arithmetic 함수는 대상 메모리 위치에서 간단한 연산을 수행하며, 덧셈/뺄샘/최대/최소/증가/감소와 같은 일반적인 연산을 포함합니다.

atomic bitwise 함수는 대상 메모리 위치에서 비트 연산을 수행하며 AND/OR/XOR을 포함합니다.

atomic swqp 함수는 조건부 또는 무조건적으로 메모리 위치의 값을 새로운 값과 교환합니다. atomic swap 함수는 swap의 성공 여부와 관계없이 항상 원래 대상 위치에 저장되어 있던 값을 반환합니다.

atomicExch는 무조건적으로 저장된 값을 변경합니다. atomicCAS 현재 저장된 값이 호출 스레드에 의해 지정된 값과 일치할 경우 조건부로 저장된 값을 변경합니다.

 

방금 위에서 살펴본 커널 incr는 atomicAdd 함수를 사용하여 다시 작성할 수 있습니다. atomicAdd는 자동으로 메모리 위치 M에 V의 값을 더합니다.

__global__
void incr(int *ptr)
{
	int temp = atomicAdd(ptr, 1);
}

다시 작성된 incr 커널은 ptr 위치에 저장된 값을 1 증가시키며, 증가되기 전 ptr에 저장되어 있던 값을 리턴합니다.

이렇게 변경하면, 이 커널의 동작은 이제 well-defined 입니다. 만약 32개의 스레드가 수행되면 *ptr에 저장된 값은 32가 증가되어 있습니다.

 

만약 어플리케이션에서 모든 스레드가 타겟을 성공적으로 증가시킬 필요가 없거나, 하나 또는 몇 개의 스레드에서 성공했는지 확인하려면 어떻게 해야할까요?

아래 커널을 살펴봅시다.

__global__
void check_threshold(int *arr, int threshold, int *flag)
{
	if (arr[blockDim.x * blockIdx.x + threadIdx.x] > threshold) {
    	*flag = 1;
    }
}

모든 커널이 할당된 데이터 값과 threshold를 비교합니다. 만약 그 값이 threshold보다 크다면 global flag는 set이 됩니다. 모든 스레드가 동일한 global flag에 대해 동작하므로, 여러 값이 threshold를 초과하면 플래그 변수는 안전하지 않습니다.

이는 stomicExch를 사용하여 안전하지 않은 액세스를 제거할 수 있습니다.

int atomicExch(int *M, int V);

atomicExch는 무조건적으로 M에 저장된 값을 V로 변경하고, 이전 값을 리턴합니다. check_threshold 커널을 atomicExch를 사용하여 flag에 대한 안전하지 않은 액세스를 제거하면 다음과 같습니다.

__global__
void check_threshold(int *arr, int threshold, int *flag)
{
	if (arr[blockDim.x * blockIdx.x + threadIdx.x] > threshold) {
    	atomicExch(flag, 1);
    }
}

이 예제에서 stomicExch를 사용하는 것이 꼭 필요할까요?

이 경우, 만약 안전하지 않은 액세스를 사용한다면, 여전히 적어도 하나의 스레드가 *flag에 성공적으로 wrtie한다고 보장할 수 있습니다. atomicExch를 사용하는 것은 실제로 이 커널의 동작을 변경하지는 않습니다. check_threshold 커널에서 단순히 안전하지 않은 액세스를 사용하더라도 프로그램은 올바르게 동작합니다. 사실 atomicExch 및 다른 atomic 연산을 사용하면 성능이 크게 저하될 수 있습니다.

 

atomic instructions는 GPU와 같이 매우 병렬적인 환경에서 강력합니다. 수백 또는 수천 개의 스레드가 공유하는 값에서 안전하게 동작할 수 있는 방법을 제공합니다. 하지만 intricsic 함수가 그러하듯, 추가적인 정밀도 문제로 고통받지는 않지만, atomic 연산의 사용은 성능을 심각하게 저하시킬 수 있습니다.

 

---

다음 포스팅을 통해서 이러한 명령어들을 최적화하는 방법에 대해서 알아보도록 하겠습니다.