티스토리 툴바


블로그 이미지
궁수자리

Rss feed Tistory
프로그래밍 이야기 2011/03/12 00:00

키넥트 카메라 캘리브레이션 하기 (Kinect Camera Calibration)

키넥트의 카메라는 컬러 영상을 획득하는 RGB 카메라와 Depth 영상을 만들어내는 IR 카메라로 구성되어 있다.  Kinect가 IR 카메라로 깊이 (Depth) 영상을 얻어내는 방법이 Time of Flight 방식인 걸로 아는 분들이 많은데, 실제로는 패턴 (Structured light)을 대상물에 투영하고 스테레오 매칭 (Stereo matching)을 통해 깊이값을 계산하는 것으로 보인다. 일단 IR 영상을 보면 키넥트의 emitter 에서 수많은 점들을 환경에 투영하는 것을 볼 수 있는데, 깊이 값이 나오는 경우는 투영된 두 인접한 점 패턴들이 서로 구분이 가능한 거리만큼 대상물이 카메라와 멀리 떨어져 있어야 하기 때문이다. 만일 ToF 방식이라면 가까워도 깊이값이 나와야 할 것인데, 실제로는 그렇지 않다. 

영상을 비교해보면 RGB 카메라의 화각 (FOV)이 IR 카메라의 화각보다 더 넓으며 컬러 영상이 깊이 영상과 동일한 장면을 찍고 있는 것이 아니라는 것을 알 수 있다. 아래 사진은 컬러 영상과 깊이 영상을 겹쳐놓은 것인데, 손가락의 위치나 팔 등이 전혀 일치하지 않는 것을 볼 수 있다. 

깊이 영상과 컬러 영상 사이에는 1:1 픽셀 매칭이 되지 않는다.

이미 인터넷의 수 많은 동영상들에서 볼 수 있듯이, 키넥트를 제스쳐 인식 등의 상호작용 분야에  사용하려고만 한다면 depth 영상을 사용하는 데 별로 큰 문제는 없다. 그러나 컬러 영상과 연동해서 사용하려고 한다면 좀 문제가 된다. 예를 들어, 컬러 영상에서 특정 위치의 3D 정보를 얻어오고 싶은 경우, 동일한 위치에 있는 깊이 영상의 픽셀이 가진 값을 읽어서 사용하는 것은 옳지않은 방법이다. 왜냐하면 두 카메라의 다른 특성으로 인해 깊이 영상과 컬러 영상 간에 1:1로 픽셀이 대응하는 것이 아니기 때문이다.  따라서 컬러 영상에서 특정 픽셀이 갖는 깊이 값을 알아내기 위해서는 두 카메라 사이의 캘리브레이션이 반드시 필요하다. 

카메라 캘리브레이션은 잘 알려져 있는 체스 보드 패턴을 이용해 쉽게 할 수 있다. 

먼저 체스 보드 패턴을 프린트해서 출력하고 IR 카메라와 RGB 카메라로 패턴의 영상을 획득한다. 이 때 두 카메라가 동일한 장면을 찍도록 하는 것이 나중에 확인을 하는 데 수월하다. libfreenect 라이브러리는 IR 영상과 컬러 영상을 동시에 얻을 수는 없으므로, 패턴을 움직여서 고정시킨 후 각각 카메라에서 따로 영상을 저장해 주어야 한다.  IR 카메라로 영상을 획득할 때는 패턴을 투영하는 Emitter 부분을 가려주어야 패턴 영상을 제대로 얻을 수 있으며, IR 카메라는 형광등 조명 아래서는 너무 어둡게 영상이 찍히므로 백열등을 하나 준비해서 켜주는 것이 좋다.  또한 IR 영상의 해상도는 640x488 라는 것에 주의하자.

체스보드 패턴을 이용해 찍은 IR 영상과 컬러 영상

IR Emitter 를 가리지 않으면 점 패턴으로 인해 제대로된 영상을 얻을 수 없다

영상 획득이 끝나면 각각의 카메라에 대한 캘리브레이션을 수행한다. 캘리브레이션은 OpenCV 함수들을 이용해 직접 프로그램을 만들어서 하거나, MATLAB ToolboxGML Camera calibration 프로그램을 써서 할 수 있다. 캘리브레이션이 끝나면 각 카메라의 내부 파라미터 행렬 K_ir과 K_rgb 및 두 카메라의 distortion parameters 를 얻는다. 

GML Camera Calibration Toolbox

각 카메라의 특성 외에도 두 카메라 사이의 관계를 알아내야 한다.  두 카메라 사이의 관계는 3x3 Rotation 행렬과 3x1 Translation 벡터로 표현되는데, 이 관계는 앞서 획득한 두 카메라의 영상들로부터 알아낼 수 있다.  두 카메라 사이가 대략 2.5cm 정도 떨어져 있고 거의 평행 하므로 Rotation 행렬은 거의 Identity 행렬에 가깝게 나오게 되며, Translation 벡터도 주로 X축 방향으로만 값을 가지게 된다. OpenCV를 사용하는 경우 cvStereoCalibrate 함수를 이용하면 구할 수 있다. 

두 카메라의 특성 및 두 카메라 사이의 관계를 알았으니, 이제 컬러 영상의 깊이 값을 얻는 과정을 알아보자. 컬러 영상의 깊이 값을 얻는 과정은 의외로 단순하다. 

우리는 키넥트가 제공하는 깊이 영상으로부터 깊이 영상의 한 픽셀이 3D 공간상에서 어느 위치에 있는지를 back-projection을 통해 알아낼 수 있다. 이 3D 공간 상의 점은 IR 카메라의 좌표계에 있는데, 이것을 RGB 카메라의 좌표계로 가져오기 위해서 앞서 구한 Rotation 행렬과 Translation 벡터를 사용한다. RGB 카메라의 좌표계로 변환된 점의 Z 값은 해당 점을 RGB 영상으로 다시 투영한 픽셀에 대응하는 깊이 값이 된다. 

다시 말하면 아래와 같은 관계다.  

For a pixel p_ir in IR image, 

P_ir = inv(K) * p_ir 
P_rgb = R * P_ir + t
p_rgb = K_rgb * P_rgb
p_rgb의 depth = P_rgb의 Z 값

P_ir : IR 카메라의 좌표계로 표현된 3D 공간의 점
R, t : IR 카메라의 좌표계에서 RGB 카메라의 좌표계로 변환하는 행렬 및 벡터 
P_rgb : RGB 카메라의 좌표계로 표현된 3D 공간의 점 
p_rgb : P_rgb가 RGB 영상으로 투영된 픽셀. 

위에서 투영 전 후에 homogeneous coordinate로 바꾸는 과정은 생략했다. 두 개 이상의 3D 점들이 RGB 영상의 동일한 픽셀로 투영되는 경우도 있는데, 이럴 때는 가까운 것을 취하면 된다. 

아래 그림은 위에 설명한 방식으로 계산한 컬러 영상의 깊이 영상이다. 주변에 존재하는 검은 부분은 RGB 카메라와 IR 카메라의 화각 차이로 RGB 카메라가 더 넓은 영역을 보기 때문에 생기게 된다 (그 부분은 깊이 값이 존재하지 않음). 

계산된 컬러 영상의 깊이 영상

 RGB 카메라에 대해 계산한 깊이 영상과 RGB 영상을 겹쳐서 보면 캘리브레이션 이전과는 다르게 픽셀과 깊이 값이 잘 맞는 것을 볼 수 있다. 깊이값을 컬러영상에 매핑하는 과정을 생각해보면,  p_rgb가 갖는 컬러값은 곧 p_ir 이 가져야 하는 컬러값으로 볼 수 있으며, 따라서 깊이 영상에 컬러값을 매핑하는 것도 가능하다. 

계산된 깊이 영상과 컬러 영상을 겹쳐놓은 것.

아래 그림은 다른 장면에 대해서 컬러 영상, 깊이 영상, 깊이 값을 컬러 영상으로 매핑한 것, 컬러 값을 깊이 영상에 매핑한 것, 그리고 두 가지를 오버레이한 것을 보여주는 그림이다. 앞서와 마찬가지로 깊이 값이나 컬러가 서로 다른 쪽으로 잘 매핑되는 것을 볼 수 있다. 

컬러-깊이 영상간 매핑 결과 (클릭하면 커짐.)


구현 동영상 


카메라 보정과 양쪽 영상간의 매핑을 통해 컬러 영상에 대한 깊이 값을 구할 수 있게 되었으니 이제 깊이 값을 컬러 영상과 연동하여 사용할 수 있게 되었다~.  





저작자 표시 비영리 변경 금지
프로그래밍 이야기 2011/02/24 14:42

Loop Unrolling 은 성능 향상에 도움이 된다.

CUDA로 알고리즘을 짠 후 성능 향상 방안을 생각하고 있었는데, 웹 사이트 들의 글들을 보면 Loop unrolling 이 성능 향상에 꽤나 도움이 된다는 이야기들이 많았다. 30%정도 성능 향상이 있었다는 글도 있었다. 

그래서 중첩된 for loop 가 있는 곳에서 #pragma unroll 매크로를 사용해보니 확실히 성능 향상이 있긴 있었다. 아래와 같이 내부 for 문에 대해서 적용해 주었더니, 

for(int y=0; y<h; y++) 
{
#pragma unroll
	for(int x=0; x<w x++) 
	{
		// Do something ... 
	}
} 

대략 34 ms 걸리던 작업이 25-26 ms 로 향상 되었다. 영상 처리할 때 써먹으면 괜찮을 듯 하다. % 로 보면 대략 20% 정도인데, 경우에 따라 다른 결과가 나오지 않을까 싶다. 내 경우에 있어서는 소요 시간을 따져보면 7-8 ms 정도가 그렇게 큰 성능 향상이 아니기 때문에 그다지 이득은 없었다고 볼 수 있다.  

저작자 표시 비영리 변경 금지
cuda, GPGPU, GPU
프로그래밍 이야기 2011/02/24 11:26

Nicolas Burrus 의 키넥트(Kinect) RGB Demo


Nicolas Burrus 라는 사람이 만들어 놓은 키넥트용 데모 프로그램.

간단하게 테스트 해 볼 수 있는 데모 프로그램이다. 맥용도 빌드된 것을 배포하고 있어서 내려받아 실행해 보았다. 
뭐.. 실행은 잘 된다 ^^. 프로그램 짜지 않고 간단하게 테스트 용도로는 괜찮은 것 같다. 


저작자 표시 비영리 변경 금지
프로그래밍 이야기 2011/02/24 11:16

Mac OS X 에서 키넥트(Kinect) 사용하기 ~

요즘 인기를 끌고 있는 마소의 키넥트(Kinect)는 Computer Vision 연구자들 사이에서도 주목 받는 카메라이다. 실시간으로 영상의 깊이 값을 얻어올 수 있다는 매력이 있기 때문인데, 깊이 정보를 이용하는 것은 대체로 로봇 비전 쪽에서 많이 사용되어 왔는데, 로봇 비전 쪽에서는 레이저 스캐너 같은 것들을 가지고 영상을 3차원으로 복원한 것을 많이 이용해 왔는데, 키넥트와 같이 하나의 패키지로 구성된 장치가 나왔으니 앞으로 로봇 쪽에서도 활용도가 높아질 것 같다. 

키넥트를 맥에서 사용하기 위해서 필요한 것은 라이브러리 두 개를 내려받아 빌드하고 설치해 주면 된다 (libFreenect 와 libusb-1.0-devel).  libFreenect는 키넥트를 사용하기 위한 API를 제공해주고, libusb-1.0은 USB 카메라를 통해 키넥트로부터 영상을 받아오는 드라이버 역할을 한다. 

OpenCV를 MacPorts 를 통해 설치한 경우, libusb-1.0이라는 라이브러리가 이미 설치되어 있으나, 키넥트에 필요한 libusb-1.0-devel 과는 다른 버전이므로 따로 설치해주어야 한다. MacPorts를 통해 설치하려 하면 기존의 libusb-1.0 과 충돌이 나서 설치가 되지 않으므로, 이 경우 그냥 소스를 다운받아 설치해 준다. 

두 라이브러리가 설치후 glview 예제를 실행하면 컬러와 깊이 영상을 렌더링해주는 것을 볼 수 있다. 아래 그림은 예제 코드를 XCODE 프로젝트로 만들어서 실행한 결과이다. 소스코드를 약간 고쳐서 컬러영상과 깊이 영상이 겹치도록 하였다. 

테스트는 맥북 프로 (15인치, Core2Duo 2.3Ghz, Geforce 9600M GT) 에서 했는데 팬이 돌아가는 것 빼곤 아주 잘 실행되었다.  

이제 이걸로 뭘 하지 ?? ^^





테스트용 XCODE 프로젝트 ZIP :  
저작자 표시 비영리 변경 금지
프로그래밍 이야기 2011/02/17 21:57

CUDA 와 OpenCL을 모두 사용해 본 느낌.

최근에 연구를 위해서 GPU 프로그래밍 언어를 공부하던 중, OpenCL과 CUDA를 사용해보고 느낀 점을 간략하게 적어본다. 

CUDA는 CUDA Runtime API와 Driver API의 두가지로 프로그래밍 할 수 있다. Runtime API는 사용의 편의를 위해 여러가지 설정이나 인자들을 자동으로  셋팅되도록 해 놓은 것으로 내부적으로는 Driver API로 구현되어 있다. Rumtime API를 사용하면 CUDA 커널 한수와 C++ 과의 연동 시 코드를 쉽고 직관적으로 짤 수 있다. 이에 반해 Driver API는 CUDA Device나 메모리 등을 일일이 프로그래머가 핸들링 해주어야 하므로 같은 작업이라도 Runtime API에 비해 보다 손이 많이 간다. 그렇지만 Driver API를 사용해야만 하는 작업들이 있으므로 나름 각각 장단점이 있다. OpenCL의 경우 CUDA의 Driver API에 가깝게 만들어져 있다. 그래서 장치 설정부터 context, command queue, memory object등 'cl_' 가 접두어로 붙는 변수 형들을 통해 관리해 주어야 할 것들이 많다. 

커널을  실행할 때 CUDA의 경우 Thread 블록의 갯수와 한 블록당 Thread의 갯수를 명시하도록 되어 있는 반면, OpenCL의 경우 실행할 전체 Thread의 갯수와 Thread 그룹 당 Thread의 갯수를 명시하도록 되어 있다. 결국 동일한 내용이긴 하지만, 실행해야 할 작업의 갯수 (총 Thread의 수)를 명시하는 OpenCL쪽이 좀 더 직관적인 것 같다. 

CUDA에서 kernel 실행을 위해서는 my_function<<<,>>>(...) 과 같이 함수 호출 형태로 사용할 수 있는 반면 (Runtime API 사용시), OpenCL에서는 command queue 라는 것에 enqueue 하는 형태로 사용하게 된다. 그래서 CUDA는 약간 더 직관적으로 코드를 이해할 수 있는 반면에, OpenCL은 CPU/GPU를 나누어 활용한다는 느낌이 강하다.  

Device에 메모리를 할당하고 관리하는 것은 CUDA가 보다 직관적이다. OpenCL은 cl_mem 이라는 메모리 오브젝트를 선언해서 사용하는 반면 CUDA에서는 C의 데이터 형을 그대로 사용한다. 대신 Device 메모리와 Host 메모리 사이에 혼동을 할 수도 있으니 naming convention을 잘 만들어 쓸 필요가 있다. 

커널 함수의 경우 어느 한쪽에서 작성된 코드를 다른 쪽으로 포팅하는 것은 (OpenCL <---> CUDA)  매우 쉬웠다. 단지 커널 Thread 의 인덱스만 알맞게 계산해 주는 부분만 잘 처리하면 그 외에는 거의 바꿀 것이 없었다.

OpenCL은 Thread의 Index를  get_global_id() 라는 함수를 통해 바로 제공하는 반면, CUDA에서는 Block Index 와 Thread의 Index 로부터 계산해야 한다. 개인적으로는 OpenCL의 방식이 더 마음에 들었다. CUDA에서도 자동으로 해줄 수 있을 것 같은데, 그렇지 않을 걸 보면, 뭔가 이유가 있을 듯 (?).

CUDA는 kernel 함수가 있는 .cu 파일을 프로그램 빌드시 함께 빌드해서 집어넣는 반면 OpenCL은 OpenGL의 쉐이더와 같이 프로그램 실행시에 커널 소스를 불러와 컴파일한다. 따라서 OpenCL은 실행시켜 보기 전에는 에러가 있는지 알기 어렵다.  그렇지만 NVCC 같은 커스텀 컴파일러가 필요한 게 아니라서 개발시 귀찮은 설정이 필요하지 않다는 장점도 있다. OpenCL을 플랫폼 독립적으로 만들기 위해서는 당연한 선택이었을 것 같다.  모든 기기에 대해 NVCC 같은 커스텀 컴파일러를 제공할 수는 없을테니 말이다. 

CUDA와 OpenCL 둘 중 어느쪽이 좋다고 말할 수는 없지만, 현재 상황에서 연구 목적이라면 여러가지 자료도 많고, 공개된 소스코드도 많이 공개되어 있는 CUDA가 낫지 않을까 싶다. 


저작자 표시 비영리 변경 금지
cuda, GPGPU, GPU, OpenCL
프로그래밍 이야기 2011/02/17 21:33

Dominant Orientation Template 코드 Mac OS X 에서 실행하기

Stefan Hinterstoisser 가 공개한 Dominant Orientation Template 은 텍스쳐가 별로 없는 객체들을 영상에서 Detection 하는데 유용한 알고리즘이다. 기본적으로 VS 용 프로젝트로 공개가 되어 있는데, 맥에서 실행하는 방법을 정리한 글이 여기에 있다.  


저작자 표시 비영리 변경 금지
프로그래밍 이야기 2011/02/14 16:50

Mac OS X 의 OpenCL 에서 Image Object

Apple의 하드웨어는 NVIDIA 와 ATI 의 제품 모두를 채택해서 출시되고 있다. 맥북프로의 경우 Geforce를, 아이맥의 경우 Radeon을 채택하고 있다. XCODE에서 제공하는 OpenCL은 동일한 인터페이스를 제공하므로 맥북프로와 아이맥에서 작업하는 것에 있어서 아무런 문제가 없을 것이라고 볼 수 있으나, 현실은 항상 기대와는 어긋나는 법이다. 

이전에 언급한 ATI 하드웨어 드라이버가 갖는 문제점 외에도, ATI OpenCL 드라이버는 Image object라는 것을 지원하지 않는다는 제약 사항이 있다는 것을 발견했다. ATI 하드웨어를 가진 아이맥에서 CL_DEVICE_IMAGE_SUPPORT 라는 쿼리를 날려보면 NO라는 대답이 자신있게 돌아오면 OpenCL kernel 빌드도 실패한다.  스노우 레퍼드 출시 당시에 ATI 에서 제공한 드라이버가 Image object를 지원하지 않았나보다. 그렇지만 Apple의 OpenCL 문서에는 Image object에 대해 떡하니 설명하고 있는데, 자사의 하드웨어에서 실행이 안되는 걸 보면 OpenCL 적용하는 것이 아직 성숙한 단계가 아닌 것 같기도 하다. 

이런 문제는 Apple 만의 문제가 아니라 OpenCL 자체가 갖는 문제점으로  지적되고 있는 점이다. 하드웨어에 독립적으로 코딩을 하는 것이 가능하지만, 하드웨어의 특성에 따라 최적화 하기 위해서는 어쩔 수 없이 kernel 함수의 수정이 불가피 하다는 것이다. 따라서 특정 하드웨어에서 높은 성능을 끌어내도록 디자인된 kernel 함수가 다른 벤더의 하드웨어에서도 동일한 성능을 보여준다는 보장이 없다는 것을 OpenCL  사용시 고려해야만 한다.  

OpenCL, 산 너머 산 ~ 

저작자 표시 비영리 변경 금지
프로그래밍 이야기 2011/02/14 16:40

Apple Mac OS X 의 OpenCL 구현에 존재하는 버그

많은 분들이 아시다시피, 애플 하드웨어용 그래픽카드 드라이버는 Apple 을 통해서만 업데이트가 가능하다. 그러니까 ATI나 NVIDIA에서 나오는 Geforce 나 Catalyst 드라이버를 다운받아 내 맘대로 설치해 사용할 수 없다는 얘기다. 그래서 CUDA 나 OpenCL을 위한 SDK 및 드라이버가 지속적으로 업데이트 되고 있지만 해당 기능을 맘대로 사용하지 못하는 경우가 있다 (NVIDIA 같은 경우 CUDA SDK 다운로드 시 개발자용 드라이버를 같이 제공하기는 한다). 현재 Mac OS X가 제공하는 OpenCL 버전은 1.0으로,  언제 다음 버전으로 업데이트 될지 알 수 없는 상황이다. 

OpenCL을 공부하면서 서로 다른 그래픽 카드를 가진Macbook Pro (NVDIA GeForce 9600m GT)와 iMac (ATI Radeon HD 5750)을 대상으로 테스트를 했는데, iMac에서 결과가 이상하게 나오는 경우가 있었다.  아래와 같이 단순히 이미지 복사를 위해 만든 kernel 함수가 NVIDIA GPU에서는 정상적인 결과를 내어주었지만, ATI 하드웨어를 가진 iMac에서는 잘못된 결과가 나왔다. 

__kernel void image_copy_simple
__global unsigned char *src, 
__global unsigned char *dst,
const int pitch,
const int chans
)
{
  const uint gid_x = get_global_id(0) ;
  const uint gid_y = get_global_id(1) ;
  uint index = gid_x *chans + gid_y * pitch ;

  if(chans == 1) 
    dst[index] = src[index] ;
  else if(chans == 2) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
  else if(chans == 3) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
  else if(chans == 4) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
}

어디에 문제가 있는지 모르겠지만 위 kernel 에서 dst 의 형을 float* 으로 바꾸면 올바로 동작하는 걸로 봐선 아마도 ATI 용 OpenCL 드라이버의 버그가 아닌가 싶다. Apple에서 드라이버 업데이트를 해주지 않는 한 이 버그는 계속 존재할테니, 문제 해결은 요원한 것 같다.
저작자 표시 비영리 변경 금지
Mac OS X, OpenCL
프로그래밍 이야기 2011/02/11 16:11

OpenCL Shared Memory 사용 예제.

OpenCL 이나 CUDA 구현을 통해 이득을 얻으려면 shared memory 를 잘 사용하는 것이 필수적이다. 행렬 곱셈을 구현하면서 어떻게 속도 향상을 이룰 수 있는지 단계별로 잘 설명하는 튜토리얼이 여기에 있다. 

처음에 global memory 만을  사용하여 구현하고 이것을 shared memory 를 사용해서 계산을 보다 빠르게 하는 과정이 단계별로 잘 설명되어 있으며, 그 내용은 아래와 같다. 



저작자 표시 비영리 변경 금지
OpenCL
프로그래밍 이야기 2010/06/16 19:21

iOS 4 에서 마커 트래킹 데모

알려진 바대로 iOS 4 에서는 Full camera access 를 지원한다. Developer 도큐먼트를 참고해서 예전에 만들었던 마커 트래킹 하는 프로그램을 iOS 4로 포팅하고 카메라 부분만 바꾸어주었다. 대충 잘 되는 듯 하지만 아직 정확히 얼마나 성능이 나오는지는 측정을 해보지 못했는데, iPhone 3G 라서 그런지 딜레이가 좀 있는 것으로 봐서 꽤나 느리게 동작하는 것 같다. 

아래는 발로 만든 데모 동영상 ~ 


 
저작자 표시 비영리 변경 금지
TOTAL 73,112 TODAY 4