- 공유 메모리 사용법

CUDA 개발자에게 성능 향상의 관건은 공유 메모리, 상수 메모리, 레지스터 등의 프로세서 지역 메모리를 효율적으로 사용하는 것입니다. 앞 장에서 살펴본 것처럼, 전역 메모리 밴드폭이 60GB/s 라고 하지만, 실제 부동소숫점 사용시 15GF/s 밖에 안 됩니다. 높은 효율을 얻기 위해서는 지역 메모리를 재사용하는 것이 필수적입니다. CUDA 장치 개발자들은 전역 메모리의 속도지연 문제와 밴드폭 향상을 위해 노력해 왔습니다만, 그와 동시에 지역 메모리 재사용이 필요합니다.

2장에서 커널 실행을 위해 블럭 당 쓰레드 갯수와 그리드를 구성하는 블럭의 갯수 등의 실행 환경 설정을 살펴보았습니다. 블럭 내 쓰레드는 지역 멀티 프로세서 자원을 통해 서로 통신을 할 수 있다는 점이 중요합니다. 왜냐하면 CUDA 실행 모델에서 블럭은 단일 멀티 프로세서 내에서만 수행되도록 설계되었기 때문입니다. 다시 말하면, 블럭 내 공유 메모리에 씌어진 데이터는 그 블럭 내 다른 쓰레드에서도 사용이 가능하지만, 다른 블럭에 있는 쓰레드에서는 억세스가 불가능합니다. 이러한 특성의 공유 메모리는 하드웨어에서 매우 효율적으로 구현 가능하며 CUDA 개발자들에게 매우 빠른 메모리 사용을 가능하게 합니다. (몇 가지 제약에 대해서는 곧 살펴보겠습니다.)

이제 CUDA 하드웨어를 개발하는 사람들이 CUDA 소프트웨어 가발자들의 요구와 장치의 가격 사이의 균형을 맞추는 방법을 생각해 봅시다. 개발자 측면에서는 레지스터나 공유 메모리 등 대량의 지역 멀티프로세서 자원을 필요로 합니다. 소프트웨어 개발이 쉽고, 효율적인 소프트웨어를 만들 수 있으니까요. 반면 하드웨어 개발자에게는 저가격의 제품이 중요한데 불행히도 고속 지역 멀티프로세서 메모리는 비쌉니다. 저가형 CUDA 장치는 누구나 환영하는 것이므로, CUDA 지원 장치들은 다른 기능을 보유한 다양한 가격대의 제품들이 시장에 나와 있습니다. 시장은 성능 대비 가격에서 적당한 제품을 선택하게 됩니다. 기술 혁신은 매우 빠르게 진행되므로 이 방식은 매우 좋은 해결방법입니다. CUDA 지원 신제품은 항상 이전 세대보다 훨씬 강력하면서 이전 세대 제품들과 동일한 가격대에서 더 많은 효율적인 컴포넌트를 보유하게 됩니다.

그런데 이 문제는 소프트웨어 개발자 측면에서는 골치거리가 되는데, CUDA 개발자들은 이러한 다른 하드웨어 구성을 고려해야 하고 한정된 장치 자원 내에서 문제를 풀어내야 합니다. 이런 다른 구성 중에서 가장 적합한 고성능 수행 환경을 선택하는 것을 도와주는 몇 가지 도구들이 제공됩니다. 저는 CUDA occupancy calculator (http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls) 를 다운로드하여 써 보시기를 권합니다. 간단한 스프레드시트입니다. (nvcc 컴파일러는 --ptxas-options=-v 를 지정하면 이 스프레드시트에 넣어줄 커널 당 각 레지스터의 갯수 혹은 지역, 공유, 상수 메모리 사용 등의 값을 출력합니다. 매 회 반복하는 조언입니다만, 다른 구성으로 컴파일해서 그것이 성능에 미치는 영향 등을 실험해 보세요. 실행 환경 설정은 변수로 지정하는 것이니 쉽게 실험해볼 수 있습니다. 대부분 프로그램들은 최적의 실행 환경 설정을 설치와 함께 스스로 쉽게 찾을 것입니다. 그리고 CUDA 런타임 호출 cudaGetDeviceCount() 와 cudaGetDeviceProperties() 을 이용해서 시스템 내 CUDA 장치의 번호와 특성을 알아낼 수 있습니다. 이 정보를 이용하는 한 가지 방법은 최적의 실행 환경 설정의 여러 조합 중에서 찾아내거나 자동 설정을 호출하는 것입니다.

- CUDA 실행 모델

하드웨어 고유의 성능 향상을 위한 방법은 다수의 블럭을 동시에 처리하는 기능을 갖추는 것입니다. 얼마나 많은 블럭을 동시 수행할 것인지는 주어진 커널 내에서 사용되는 쓰레드 당 레지스터의 갯수와 블럭당 공유 메모리 크기에 따라 달라집니다. 단일 멀티프로세서에서 주어진 시간에 처리되는 블럭들을 활성 블럭이라 합니다. 커널의 최소 자원 사용은 각 멀티프로세서를 효율적으로 사용하게 합니다. 왜냐하면 멀티프로세서 내의 레지스터와 공유 메모리는 활성 블럭 내 쓰레드로 분배되기 때문입니다. CUDA occupancy calculator 를 이용해서 쓰레드 갯수와 활성 블럭 갯수 대비 레지스터 갯수와 공유 메모리 용량의 트레이드 오프를 살펴보세요. 자원의 올바른 조합은 커널의 성능을 크게 향상시킬 것입니다. 만약 적어도 한 블럭을 구동할 만한 레지스터나 공유 메모리 자원이 멀티프로세서 내에 없다면, 커널은 실행 불가합니다. (3회에서 소개한 cudaGetLastError() 를 이용해서 어떤 에러가 감지되는지 살펴보세요.)

각 활성 블럭들은 "워프" 라는 이름의 SIMD (Single Instruction Multiple Data) 쓰레드 그룹으로 분리됩니다. 각 워프는 동일한 갯수의 쓰레드 "워프 사이즈"를 가지고 있습니다. 이것은 멀티프로세서 내에서 SIMD 방식으로 수행됩니다. 다시 말해 워프 내 각 쓰레드는 명령 스토어 내에서 동일한 명령으로 분배됩니다. 이렇게 해서 쓰레드는 지역 혹은 전역 변수 내에 동일한 명령이나 조작을 수행하게 됩니다. SIMD 모델은 하드웨어 측면으로 보면 저가의 효율적인 방식이지만, 소프트웨어 측면에서는 조건문을 동시 처리할 수 없고 순차적으로 처리해야 합니다. 커널 내 조건문은 실행 효율에 심대한 영향을 미칠 수 있음에 주의해야 합니다. 주의해서 작성하면 대부분 잘 처리할 수 있지만 경우에 따라서는 문제가 될 수 있습니다.

활성 워프 (활성 블럭 내의 모든 워프들) 들은 시간별로 쪼개어집니다. 쓰레드 스케쥴러는 주기적으로 워프들을 차례로 이동함으로써 멀티프로세서의 계산 자원을 최대한 효율적으로 사용합니다.블럭 내 워프들의 실행 순서와 블럭의 실행순서는 따로 지정되지 않고 무작위로 진행될 수 있습니다. 하지만 쓰레드는 __syncthreads() 를 이용해서 동기화할 수 있습니다. __syncthreads() 명령 이후에 공유 (전역) 메모리에 쓰는 값만이 확실히 참조 가능하다는 점에 주의하십시오. 변수가 volatile 로 선언되지 않았다면 컴파일러는 속도 향상을 위해 메모리 입출력 최적화를 가할 수 있습니다. (다시 말해 변수가 임의로 재배치 되거나 제거될 수 있습니다.) __syncthreads() 는 조건문 내에 넣을 수 있습니다만, 전체 쓰레드 블럭에서 동일한 조건 하에서만 수행됩니다. 그렇지 않으면 코드 실행이 멈추어 버리거나 원치 않는 부작용이 발생할 수 있습니다. 다행해 __syncthreads() 는 4 클럭사이클 정도를 소비하는 부하가 적은 함수이고 따라서 어떤 쓰레드가 다른 쓰레드 수행을 대기해야 하는 경우는 없습니다. 워프의 전반부 혹은 후반부만을 뜻하는 반워프는 이번 호 후반에 소개할 협력형 메모리 억세스를 포함한 메모리 억세스에 중요한 개념입니다.

지금까지 내용을 요약해 보면 다음과 같습니다.

* 멀티프로세서 내 공유 메모리같은 자원은 한정되어 있고 소중하다.
* CUDA 장치 구성 내에서 공유 메모리 같은 한정된 멀티프로세서 자원을 효율적으로 관리하는 것은 CUDA 개발자의 숙명이다.
* (if 문과 같은) 조건문은 커널 실행에 큰 영향을 미칠 수 있으니 주의하자.
* CUDA occupancy calculator 와 nvcc 컴파일러는 실행 환경 설정을 결정하기 위한 유용한 툴이니 잘 익혀두고 사용하자.

- CUDA 메모리 모델

그림 1 은 전역 메모리와 칩 내의 다양한 메모리를 억세스하는 쓰레드의 실행을 나타내는 그림입니다.

< 그림 1 > http://www.ddj.com/architect/208401741;jsessionid=KGR5BORXP1IOPQE1GHRSKHWATMY32JVN?pgno=3

위에서 블럭 (0, 0) 과 블럭 (1, 0) 으로 표기된 각 멀티프로세서에는 다음 네 가지 종류의 메모리가 탑재됩니다.

* 쓰레드 당 지역 레지스터 한 조
* 모든 쓰레드가 공유하는 공유 메모리 영역을 구현한 병렬 데이터 캐시 혹은 공유 메모리
* 모든 쓰레드에서 공유하면서 상수 메모리 영역에서 빠른 데이터 입력을 구현한 읽기 전용 상수 캐시, 장치 메모리 내 읽기 전용 영역 (상수 메모리는 향후 다루어질 것입니다. 그 때까지는 CUDA 프로그래밍 가이드 5.1.2.2 장을 참조하세요)
* 모든 쓰레드에서 공유하면서 텍스쳐 메모리 영역에서 빠른 데이터 입력을 구현한 읽기 전용 텍스쳐 캐시, 장치 메모리 내 일기 전용 영역 (텍스쳐 메모리는 관련 글에서 다루어질 것입니다. 그 때까지는 CUDA 프로그래밍 가이드 5.1.2.3 장을 참조하세요)

멀티프로세서 내에 "지역 메모리" 라는 이름의 블럭이 그림에 들어있다는 사실에 혼돈하시면 안 됩니다. 지역 메모리는 "각 쓰레드 영역 내의 지역" 을 의미합니다. 이것은 메모리 추상화 개념이지 실제 멀티프로세서 내 하드웨어 컴포넌트가 아닙니다. 실제로, 지역 메모리는 컴파일러가 전역 메모리에서 일부 할당된 메모리이며 다른 전역 메모리 영역과 동일한 성능을 나타냅니다. 지역 메모리는 프로그래머가 쓰레드 내에서 작업 수행에 필요하게끔 컴파일러가 사용하는 것이고 더 고속의 메모리를 뜻하지 않습니다. 일반적으로 커널 내 선언된 자동 변수는 레지스터에 상주하며, 상당히 빠른 속력을 가집니다. 경우에 따라서 너무 많은 레지스터 변수가 선언되었거나, 배열 길이가 4 이상 되거나, 너무 많은 레지스터 영역을 소비하는 구조체나 배열 같은 경우, 혹은 변수가 일정한 크기로 선언되지 않았을 경우에는 컴파일러가 지역 메모리에 선언합니다.

지역 메모리가 속도를 저하시킬 수 있음에 주의하세요. -ptx 나 -keep 옵션으로 컴파일해서 얻을 수 있는 ptx 어셈블리 코드를 살펴보면 변수들이 1차 컴파일 시기에 지역 메모리에 위치하여 .local 니모닉으로 선언되어 있고 ld.local 이나 st.local 니모닉으로 억세스되고 있는지 볼 수 있습니다. 이후 컴파일 시기에도 만약 목적하는 장치에 비해 과도한 레지스터 영역을 사용하고 있다고 판단한다면 지역 메모리에 변수를 위치시킬 것입니다.

다음 회까지 occupancy calculator 를 이용해서 실행 모델과 커널 실행 구성 요소가 레지스터 갯수와 공유 메모리 소모에 어떤 영향을 미치는지 확실히 이해해 두세요.