3D Client Programmer

예전에 다른 팀과 연동하는 과정에서 매개 변수로 넘긴 값이 함수 내부로 넘어가면서 값이 혼자 변하는 버그를 경험했다.

그림 - test.cpp

그림 - main.cpp

위와 같은 상황에서 abc()를 통해 값을 넘겼을 때, 매개변수 a[3]은 모두 제대로 값이 넘어갔지만, 매개변수 b의 경우 계속해서 0의 값이 들어가는 것을 확인했었다.

이런저런 테스트 결과 "test.h"에서 abc()의 선언을 추가함으로 써 문제가 해결되었다.

이런 일이 발생한 정확한 이유는 모르겠지만, C 형태로 코딩을 할 때에는 함수의 선언이 빠졌을 때 이런 일이 발생할 수 있으므로 주의하도록 하자.

Visual Studio 2010으로 작업하는 도중에 아래 그림과 같이 Stack Overflow가 발생했었다.

그림 - 0xC00000FD: Stack oveflow

이것은 기본 배열이 0x7fffffff 크기를 초과해서 이다.

이 문제를 해결하기 위해서는 배열 대신 동적 할당을 이용해서 Heap에 저장하는 방법이 있다. (나는 이렇게 고쳤다...)

도저히 코드 수정이 불가능한 상태라면 프로젝트의 옵션을 수정하는 방법이 있다.

그것은 원래 default로 제공되는 stack의 공간을 수정하는 방법인데 원래 값은 1MB로 되어 있다. 이것을 임의의 값으로 강제로 늘려버리면 된다.

그림 - 프로젝트 설정 값 변경

링커 -> 시스템 -> 스택 예약 크기에서 0으로 되어 있는 설정 값을 1024 * 1024 * 1 = 1MB이므로 1024 * 1024 * 10 = 10485760(10MB)로 입력한다.

이렇게 코드를 변경하지 않고 수정하는 방법도 있다. 하지만, 프로젝트 옵션을 수정하는 것은 최후의 최후의 방법으로 쓰도록 하자...

* Visaul Studio 2013을 사용하기 시작했는데, 2013은 처음부터 배열의 크기가 넘어가면 빌드 에러를 띄워준다. (고맙...)

참고 : http://ocllos.tistory.com/39

우선 파사드의 의미를 사전에서 살펴보면 다음과 같다.

1. <건축> (건물의) 정면, 전면, 파사드

2. <<비유적>> (사물의) 표면, 외관, 겉보기, 허울 (또는 facade)

위의 사전적 의미에서도 알 수 있듯이 파사드 패턴이라 함은 인터페이스를 단훈화시키기 위해서 인터페이스를 변경하는데, 하나 이상의 클래스의 복잡한 인터페이스를 깔끔하면서도, 말쑥한 파사드(겉모양, 외관등)으로 덮어주는 것이다.

파사드 패턴을 다시 한번 정의하자면, 어떤 서브 시스템의 일련의 인터페이스에 대한 통합된 인터페이스를 제공함으로써 서브시스템에서 더 쉽게 사용할 수 있도록 하는 패턴이다.

그림 - 파사드 패턴

홈씨어터를 예로 들어서 설명하자면 다음과 같은 몇가지 작업을 실행해야 된다.

1. 팝콘 기계를 켠다.

2. 팝콘을 튀기기 시작한다.

3. 전등을 어둡게 조절한다.

4. 스크린을 내린다.

5. 프로젝터를 켠다.

6. 프로젝터로 DVD 신호가 입력되도록 한다.

7 프로젝터를 와이드 스크린 모드로 전환한다.

8. 앰프를 켠다.

9. 앰프의 입력을 DVD로 전환한다.

10. 앰프를 서라운드 음향 모드로 전환한다.

11. 앰프 볼륨을 중간으로 설정한다.

12. DVD 플레이를 켠다.

13. DVD 재생한다.

위와 같은 작업을 하기 위해서 사용자가 코드를 한 줄씩 모두 작업을하게 될 것이다.

만약, 이렇게 된다면 문제가 몇가지 존재하게 된다.

1. 홈씨어터의 사용법이 너무 복잡해진다.

2. 영화를 꺼야할 때는 어떻게 처리해야 할지 고민하게 된다.

3. 시스템이 업그레이드 되면 작동 방법을 배워야 한다.

이럴 때 파사드 패턴을 사용해서 복잡한 일을 간단하게 사용할 수 있다.

그림 - CD Player 구현 코드

그림 - Tuner 구현 코드

몇가지 클래스가 더 만들어져야 하지만 글이 너무 길어지기 때문에 생략한 클래스는 (Amplifier, Projecter, TheaterLight, Screen, PopcornPopper) 등이 있을 것이다.

이제 파사드 패턴을 구현한 클래스를 보도록 하자.

그림 - 파사드 패턴 클래스

위와 같이 파사드 패턴은 단순화된 인터페이스를 통해서 서브 시스템을 더 쉽게 사용할 수 있도록 하기 위한 용도로 사용된다.

덕분에 클라이언트에서는 단순한 메소드 하나만을 호출함으로 일하기도 편리하고 수원해질 수 있는 것이다.

싱글턴(Singleton)이라는 것은 디자인 패턴의 가장 기초적인 패턴이다.

하지만 이 패턴이라는 것은 사실상 딱 부러지는 코드를 제공하는 것은 아니고 코드라는게 항상 그렇지만 100% 정답은 없다.

그 중에서 싱글턴은 아래의 조건만 맞춰 진다면 개인적으로 "어떠한" 방법을 쓰든 싱글턴이라고 부르는게 맞다고 생각한다.

1. Process내에서 단 한 개만 생성되는 것을 보장하며, 그것을 표현할 수 있을 것.

2. 사용하는 시점이 생성되는 시점일 것.

싱글턴과 전역 변수가 비슷하지만 차이 점을 하나씩 비교해보자. 

싱글턴은 Process내에서 단 한개만 생성이 된다는 것을 보장함과 동시에 표현이 가능하다는 점에 있어서 전역 변수로는 표현하기 힘든 부분이라는 것에 싱글턴의 존재성을 깨닫게 된다.

아래 예제를 보자.

그림 - 싱글턴인듯 아닌듯 한 기본 예제

위의 코드는 처음 싱글턴을 공부할 때 항상 만들어보는 예제 일 것이다. 하지만 이것만으로는 이것이 단 한개 만을 가지는지는 보장이 되지 않는다. 왜냐하면 생성자와 소멸자가 private이나 protected로 보호되지 않고 public으로 선언되어 있기 때문이다.

한마디로 굳이 CreateInstance()를 사용하지 않고 직접 new나 delete로 생성 및 삭제가 외부에서 가능해지기 때문이다. 

그림 - 싱글턴 기본 예제 1

이렇게 생성자와 소멸자를 private으로 외부에서 접근하지 못하도록 보호를 한다면, 

CSingleton* pTest = new CSingleton();

delete pTest;

위의 코드는 컴파일 단계에서 에러가 날 것이다. private이기 때문에 생성도 할 수 없고, 삭제도 할 수 없기 때문이다.

하지만, 위 처럼 싱글턴을 생성하면 static 객체는 생성의 시점이 명확하게 정의되지 않기 때문에, 다른 전역 객체에서 이를 참조하고자 할 때 생성 순서 때문에 문제가 발생할 수 있다. 그래서 new를 이용하여 생성하는 것이 더 안전해 보인다.

자 이번엔 new를 이용해서 싱글턴을 만들어보자.

위의 방법은 처음에 말했던 조건에 모두 부합한다. 사용하기 위해 CreateInstance()를 하는 순간 생성되고, 원하는 시점에 삭제가 가능하다. 

만약 CreateInstance()가 동시에 여러 곳에서 호출될까봐 무서운 사람은 Singleton 전용 Mutex를 하나 만들어서 사용하는 것도 괜찮은 방법일 것이다.

여기서, 조금 생각 해볼 것은 new로 생성한 인스턴스를 굳이 삭제해야 하는 것이다. 프로그램이 종료되는 시점에는 어차피 new로 생성한 메모리까지 전부 해제가 된다. 따라서, 어차피 생성한 목적이 프로그램 종료시까지 단 하나의 객체가 동작하는 것이라면 그 크기가 어마어마한 것이 아니라면 명시적으로 메모리 해제를 수행할 필요는 없다는 것이다. 하지만, 만약의 상황을 대비해 해제 함수를 하나 만들어 두는게 더 나아 보이긴 하다.... (어려운 작업이 아니니....)

이 싱글턴은 멀티 쓰레드에서는 안전하지 않다... 

그림 - CUDA 메모리 계층 구조

그림 - CUDA 메모리 특성

액세스 속도 (느림 -> 빠름)

Global -> Local -> Shared -> Register

메모리 용량 (작음 -> 큼)

Register -> Shared -> Local -> Global

액세스 범위 (좁음 -> 넓음)

Register -> Local -> Shared -> Global

Register Memory는 온 칩 프로세서에 있는 메모리로 직접 연산을 수행하는 가장 빠른 메모리이다. CUDA C를 이용해 구현할 때도 몇 개의 레지스터를 사용하는 지와 레지스터가 어떻게 재사용 되는지 알 수 없다.

다만 커널 함수 내에서 로컬 변수를 많이 사용할수록 레지스터 사용량이 증가하고, GPU마다 레지스터의 수는 차이가 있다.

sin(), cos() 등 수학 함수를 사용할 때도 레지스터를 사용한다.

로컬 메모리(Local Memory)는 커널 함수 내에서 너무 많은 로컬 변수를 사용하거나, 배열형 변수로 큰 용량을 사용하면 프로세서 밖에 있는 DRAM에 메모리가 할당된다. 로컬 변수가 레지스터로 사용될지 로컬 메모리에 할당 될지 명확하지 않다.

Local Memory로 할당되는 경우는 다음과 같다.

1. 너무 많은 레지스터 변수를 사용했을 때

2. 너무 많은 로컬 변수를 사용했을 때

3. 로컬 변수로 배열을 사용했을 때

쉽게 말하면, 각 Thread는 자신만이 사용하기 위한 메모리이다.

공유 메모리(Shared Memory)는 CUDA의 큰 장점 중 하나이다. CUDA C 컴파일러는 공유 메모리의 변수들을 일반 변수들과는 다르게 취급한다. CUDA C 컴파일러는 GPU에서 실행되는 각 블록마다 변수의 복사본을 하나씩 생성한다. 블록 내의 각 스레드는 메모리를 서로 공유하지만, 다른 블록 내에서 보고 있는 변수의 복사본을 보거나 수정할 수는 없다. 하지만 공유 메모리는 블록 내의 스레드들이 계산하는 데 통신과 협력을 할 수 있게 함으로써 훌륭한 수단을 제공한다. 더욱이 공유메모리 버퍼는 오프칩(off-chip)인 DRAM에 상주하는 것과 달리 물리적으로 GPU 상에 상주한다. 이 때문에 공유 메모리의 접근 지연 시간은 블록당 공유 메모리를 효율적으로 생성하는 일반 버퍼, 소프트웨어로 관리되는 캐시 또는 스크래치패드(Scratchpad)보다 훨씬 짧은 경향이 있다.

GPU의 SM 안에 공유 메모리가 있고, SM은 블록 단위로 프로그램을 처리 하기 때문에 공유 메모리의 데이터는 동일한 블록 안에 있는 스레드 사이에서만 공유할 수 있다. 공유 메모리를 할당하는 방법은 정작할당, 동적할당이 있다.

쉽게 말하면, 각 Thread 간에 data 공유를 위한 메모리이다.

__shared__ int a; 

전역 메모리(Global Memory)는 비디오 카드에 장착된 DRAM 메모리를 의미한다. 전역 메모리는 GPU 칩 외부에 있기 때문에 액세스 속도는 레지스터나 공유 메모리에 비해 많이 느리지만, CPU의 메모리 보다는 빠르다.

쉽게 말하면, Block들의 집합인 Grid간에 data 공유를 위한 메모리이다.

상수 메모리(Constant Memory)는 DRAM에 있는 데이터를 읽기 전용으로 사용하며 캐시를 지원한다. NVIDIA의 그래픽 카드는 Constant Memory의 크기가 64KB이다. 최초 데이터는 DRAM에서 가져오기 때문에 속도가 좀 느리지만, 한번 캐시에 올라온 값을 반복하여 재 사용하기 때문에 속도가 빠르다. Host에서는 값을 쓰고, Device 영역에서는 값을 읽기만 할 수 있다.

쉽게 말하면, Half - Warp내의 모든 스레드가 동일한 주소의 데이터를 상수 메모리로 부터 요청하면 GPU는 단 한번의 읽기만 요청한 후 모든 스레드에게 그 데이터를 알려주는 메모리이다.

__constant__ int a;

텍스처 메모리(Texture Memory)는 상수 메모리처럼 칩에 캐싱된다. 따라서 어떤 경우엔 텍스처 메모리는 오프칩인 DRAM에 대한 메모리 요청을 줄임으로써 매우 효과적인 대역폭을 제공할 수 있다. 텍스처 캐시는 공간 구역성(spatial locality)를 자주 드러내는 메모리 접근 패턴을 가진 그래픽스 어플리케이션들을 위해 설계되었다. 

texture<자료형> a;

출처 : http://blog.naver.com/riverrun17/220420579990

CUDA 메모리에 대해 공부하다가 어떤 분이 잘 정리해 놓으셔서 긁어왔습니다....

이번에는 앞전에 했던 블록을 이용한 벡터의 합을 스레드를 이용한 벡터의 합으로 구현해보려고 한다. 

사실 코드는 몇줄 바뀌지 않는다. 단순하게 블록에서 하던 일을 스레드에서 하게 바꾸는 작업만 하면 되는 일이니까...

앞에서 작성했던 코드와 비교해서 바뀐 부분은 빨간색 네모를 확인하면 된다.

그림 - add()에서 바뀐 부분

그림 - main()에서 바뀐 부분

앞에서 하드웨어의 제약으로 인해 한 번에 실행 가능한 블록의 수는 65,535개를 초과할 수 없다고 했었다. 동일하게, 하나의 커널을 실행하는 데 있어서 블록당 스레드 수 역시 하드웨어의 제약을 받는다. 구체적으로, 이 수는 우리가 3장에서 보았던 디바이스 속성 구조체의 'maxThreadsPerBlock' 멤버의 값을 초과해서는 안된다. 현재 이용가능한 많은 그래픽스 프로세서들에서 블록당 스레드의 최대 개수는 512개이다.

그렇다면 배열의 개수가 512개보다 더 클 때 어떻게 스레드 기반으로 두 벡터의 합을 수행할 수 있는지 생각해보자. 뭐 방법은 간단하다. 스레드와 블록을 같이 사용하는 방법이다.

그림 - 스레드와 블록을 같이 사용

여기서 'blockDim'이라는 내장 변수를 사용하는데 앞에서 설명했다시피, 각 블록당 스레드의 개수 이다. 여기서는 1차원의 블록을 사용하고 있기 때문에 그냥 'blockDim.x'만 사용했고, 2차원이라면 'y'도 사용 해야 할 것이다.

또 다른 변경은 커널 실행 자체에 있다. 실행할 N개의병렬 스레드들이 여전히 필요하지만, 우리에게 부과된 512개의 스레드 제한 수를 초과하지 않도록 다수의 블록들을 통해 스레드들을 발동시키길 원한다. 하나의 해결책은 임의로 블록 크기를 고정된 스레드의 개수로 지정하는 것이다. 예를 들어, 블록당 128개의 스레드들을 사용한다고 고정시켜보자.

그리고 블록을 생성하는 개수를 정할 차례인데, 여기서 묘안은 N/128이 정수의 나눗셈이라는 것이다. 만약 N이 127이라면 N/128은 정수로 0이 된다. 즉 0개의 블록을 생성하게 되므로 스레드도 실행되지 않게 된다. 하지만, 이렇게 구현해버리면, N이 정확히 128의 배수여야만 올바른 블록의 개수가 실행될 것이다. 이것은 좋지 못하다. 이를 해결하기 위한 일반적인 요령은 다음과 같다.

그림 - 커널의 다른 실행 방법

여기서 N이 10개 이므로, 블록은 10개 스레드는 128개가 된다. 즉 1280개의 스레드가 생성되어 돌아가기 때문에 잘못된 계산이 나올수 있으므로 예외 처리를 해줘야 한다. 하지만, add() 커널 함수에서 이미 예외 처리를 해놨었다.

그림 - 예외 처리가 된 부분

tid가 위에서 계산되었을 때 N보다 클 때에는 if문에서 걸러지기 때문에 올바르게 작동하게 된다.

하지만 이렇게 구현하게 되면, 하드웨어 제약사항으로 인한 블록의 개수(65,535) * 스레드의 개수(512) 개를 초과하게 되면, 병렬 처리가 정상적으로 처리가 되지 않게 된다. 그러므로 더 수정을 해야 한다.

그림 - 더 업그레이드 된 add() 커널 함수

위와 같이 각 스레드가 현재 인덱스에서의 작업을 끝내면 그리드 내에서 작동 중인 스레드의 총 수만큼 각 인덱스를 증가시킬 필요가 있다. 즉, while() 문에서 tid 값을 위와 같이 증가 시키면 각 스레드는 유니크한 인덱스를 구할 수 있을 것이다.

그림 - 고정된 개수의 블록과 스레드를 실행하는 add() 커널 함수

위와 같이 블록의 개수(128) * 스레드의 개수(128) = 16,384개 이지만, 3만개든 10만개든 덧셈이 동시에 가능해진다.

람다(lambda)는 '람다 함수' 또는 '이름 없는 함수'라고 부르며, 그 성질은 함수 객체와 같다.

기존 C++에서 STL의 find_if, sort 등의 알고리즘을 사용할 때 특정 조건자를 사용하려면 함수 객체를 정의해야 했다.

그런데 STL 알고리즘 함수에서만 사용하려고 따로 함수 객체를 만들려니 귀찮을 수 밖에 없었다.

그래서 보통은 함수를 따로 만들거나, 함수를 정의하는 것도 귀찮아서 그냥 STL을 사용안하고 직접 컨테이너를 다루었다.

람다 덕택에 이제는 이런 수고를 할 필요가 없어졌다.

람다의 기본 문법은 다음과 같다.

다음은 람다의 간단한 사용 예다.

auto를 사용하면 람다를 변수에 대입 할 수도 있다.

람다는 일반 함수처럼 파라미터를 정의 할 수도 있다.

반환 값도 사용할 수 있다.

find_if를 쓸 때에는 더욱 유용한 것을 느낄 수 있다.

기존에는 다음과 같이 find_if 알고리즘을 이용해서 '죽은 유저'를 찾으려면 다음과 같이 함수 객체를 정의해야 했다. 

하지만 람다를 사용하게 되면, 한 줄로 간단하게 끝낼 수 있다.

* 람다는 'decltype' 과 'sizeof' 에서는 사용할 수 없다.

아마도 C++ 11에서 가장 인기 있는 기능이 auto일 것이다. auto를 가장 잘 말해주는 예는

위 처럼 써야 할 것을 C++ 11에서는 

이렇게 쓸 수 있다. 

역시 auto는 그냥 일반 자료형 보다는 STL을 사용할 때 훨씬 편하게 코딩할 수 있게 만들어 준다.

* auto는 지역 변수에서만 사용할 수 있으므로, 클래스의 멤버 변수나 전역 변수, 함수의 인자로는 사용할 수 없다.

* auto는 '변수를 정의할 때 명시적으로 형을 지정하지 않고 컴파일을 할 때 자동으로 결정해주는 키워드' 이다.