CUDA ile OpenCV kullanarak Webcam Görüntü İşleme

For English version: http://www.codeproject.com/KB/GPU-Programming/opencv-cuda-filters.aspx

Biraz uzun soluklu olan bu yazımıza OpenCV ve CUDA hakkında kısa maddeler ile başlayalım.

NOT: Bu yazıdaki kod örnekleri kısaltılarak verilmiştir. Kod belgeleri için Doxygen ile yaratılmış olan mikrositeyi inceleyebilirsiniz.

NOT: Her tür eleştirinizi yorumlara yazabilirsiniz, teşekkür ederiz.

İçindekiler

  • Gerekenler
  • Görüntü İşleme ve Webcam
  • ISingleImageFilter Kullanımı
  • En basit filtre: IdentityFilter
  • CPU Üzerinde Negatif Görüntü Filtresi
  • Yeniden Hoşgeldin CUDA!
  • Görüntü Negatifi – CUDA
  • Texturelar
  • Görüntü Negatifi – CUDA Texture
  • Filtre Zinciri: SingleImageFilterChain
  • UML Şeması
  • Filtre Örnekleri

Gerekenler

  • VS 2010 ve C/C++ Bilgisi
  • OpenCV ile görüntü yakalama aşinalığı. Windows için OpenCV 2.2’yi patchlemeyi unutmayın! Önceki yazımızı inceleyebilirsiniz.
  • VS2010 ile CUDA projesi yaratabilmek. Önceki yazımızı inceleyebilirsiniz.
  • CUDA içeren kodları çalıştırabilmek için CUDA destekli bir ekran kartı.
  • Derleyici includeları, libraryleri ve hataları ile uğraşmak için bolca sabır.

OpenCV Nedir?

OpenCV, açık kaynaklı bir bilgisayarlı görselleştirme kütüphanesidir. Amacı görüntü işleme ve görselleştirmede sık kullanılan metodların kolaylıkla erişilebilir olmasını sağlamaktadır.

CUDA Nedir?

CUDA, Compute-Unified Device Architecture, hesaplamaların grafik işlemcisi üzerinde gerçekleştirilmesi amacıyla CPU’dan zaman bağımsız olarak ve yüzlerce çekirdek üzerinde binlerce thread kullanarak paralel hesaplamaya olanak veren nVidia’nın bir kaç senedir satışa sunulan ekran kartlarında uyguladığı bir mimaridir. Bu muazzam paralellik(massively parallel) bize doğru kullanıldığında hesaplamada hız olarak dönecektir.

Görüntü İşleme ve Webcam

Görüntü işleme, görüntüye uygulanan matematiksel bir işlemi ifade eder. Örneğin görüntünün negatifinin alınması, görüntünün 90 derece çevrilmesi, görüntünün bir kurala göre bulandırılması vb. basit görüntü işleme metodlarıdır. Bir önceki yazımızda bahsettiğimiz OpenCV ile webcam görüntüsü yakalamayı bu yazımızda bir adım ileri götürerek görüntü işleme ile genelleştirilmiş olarak birleştireceğiz.

Görüntünün webcam üzerinden, bir avi dosyasından veya başka bir kaynaktan geliyor olması OpenCV kullandığımız için fark etmemektedir.

Özetle OpenCV Görüntü Yakalama

  1. OpenCV ile kamerayı açıyoruz.
  2. İlk kareyi yakalıyoruz.
  3. İlk kareyi hatasız yakaladıysak devam ediyoruz.
  4. Kullanıcının ‘q’ tuşuna basması ile çıkacağı bir döngüye giriyoruz.
    1. Kareyi yakalıyoruz.
    2. Kareyi istediğimiz boyuta getiriyoruz.(Örnekte 640 x 480)
    3. Görüntü işleme metodumuzu çağırıyoruz.
    4. Kareyi ekranda gösteriyoruz.
  5. Programdan çıkıyoruz.

Arayüzler, Sınıflar, Kalıtım ve dahası

Burada geliştireceğimiz nokta, görüntü işlemenin main.cpp içerisindeki bir global metod olmaktan çıkarılmasıdır. Görüntü işleme için kullanabileceğimiz en basit arayüzü düşündüğümüzde, bu arayüzde görüntünün boyutları ve Filtrele metodundan başka bir içeriğe ihtiyaç olmadığını görüyoruz.

Bellek ile ilgili ayırma ve serbest bırakma özelliğinin de görüntü işleme sınıfında olması için InitFilter ve ReleaseFilter metodları kullanılmıştır.

Bu basit arabirimi aşağıdaki UML şemasında inceleyebiliriz.

ISingleImageFilter Arayüzü

main.cpp dosyasında ISingleImageFilter Kullanımı

  • Filtreyi yaratmak,
  • InitFilter’ı looptan önce çağırmak,
  • ReleaseFilter’ı looptan çıkınca çağırmaktır.
// ... OpenCV kamera açılır, kare yakalanır.

// Kullanacağımız filtre yaratılır.
ISingleImageFilter* myFilter = new IdentityFilter();

// Filtreye görüntünün boyutu bildirilir.
myFilter->InitFilter( 640, 480 );

// OpenCV görüntü yakalama döngüsü.
while( key != 'q' )
{
//	... kare yakalama.

	// Yakalanan görüntü işlenir.
	myFilter->FilterImage( imageData );

//	... kare gösterim.
}

// Filtre kaynaklarının serbest bırakılması istenir.
myFilter->ReleaseFilter();

// ... program çıkışı.

ISingleImageFilter Kullanımı

ISingleImageFilter‘da tüm filtrelerin kullanacağı en az iki değişken vardır. Görüntünün genişliği ve yüksekliği. Bu değişkenleri üye değişkenlerinde tutan ve tüm filtrelere ortak özellik eklenmesi gerektiğinde geliştirilebilecek bir sınıf olarak SingleImageFilter sınıfı yazılmıştır.

SingleImageFilter, ISingleImageFilter sınıfından kalıtlanmıştır ve InitFilter metodu aşağıda verilmiştir.

	virtual void InitFilter(int width, int height)
	{
		this->width = width;
		this->height = height;
		this->isInited = true;
	}

En basit filtre: IdentityFilter

ISingleImageFilter arayüzünün kullanımını en basit şekliyle göstermek amacıyla yazılmış bir filtre sınıfıdır. Kalıtladığı SingleImageFilter sınıfının FilterImage metodunu override eder.

class IdentityFilter : public SingleImageFilter
{
public:
	IdentityFilter()
	{
	}

	/** Görüntüde değişiklik yapmadan çıkar. */
	virtual void FilterImage(char* imageData)
	{
		return; // imajı değiştirmeden dön.
	}
};

CPU Üzerinde Negatif Görüntü Filtresi

CPU üzerinde çalışan filtrelere örnek olarak negatif görüntü filtresi verilmiştir. Bu sınıfın sadece FilterImage metodu aşağıdadır.

	/** Görüntünün RGB kanallarının tersini alır. */
	virtual void FilterImage(char* imageData)
	{
		for(int i=0; i<3*width*height; i++)
		{
			*( imageData + i ) = ( unsigned char ) ( 255 - *( imageData + i ) ); // her pikselin her kanalının negatifini al.
		}
	}

Buraya kadarki filtrelerin UML şemasını gözden geçirecek olursak:

CPU filters

Yeniden Hoşgeldin CUDA!

CUDA kodlarının çalışabilmesi için CUDA SDK 3.2 yazılımının bilgisayarınızda yüklü olması gerekmektedir. Kurulum için: VS2010 ile CUDA 3.2 Projesi Yaratmak

Görüntü işleme, CUDA kullanarak paralelleştirilmeye çok yatkın bir alandır. Her ne kadar bu yazıda anlatılan görüntü işleme metodları muazzam bir işlem gücü gerektirmese de metodların GPU’da CPU’dan daha hızlı çalıştığını farkedebilirsiniz.

CUDA üzerinde OpenCV’de yakaladığımız görüntüyü işleyebilmek için öncelikle bu görüntünün verisini ekran kartının belleğine yüklemeliyiz. Hesaplamayı GPU üzerinde yaptıktan sonra sonucu da CPU belleğine geri almalıyız.

SingleCudaFilter::FilterImage metodu

  • Kareyi yakala
  • SingleCudaFilter::FilterImage
    • Kareyi ekran kartı belleğine yükle
    • Filtreyi GPU üzerinde uygula
    • Sonucu CPU belleğine geri al
  • Sonraki kareye geç

Ekran kartı belleğinin ayrılması InitFilter metodunda, serbest bırakılması ReleaseFilter metodunda, verilerin ekran kartına kopyalanması ise FilterImage metodunda yapılır.

CUDA kullanmak için bir de kernele ihtiyacımız var. SingleCudaFilter sınıfında bu kernel parametrik olarak constructorda alınmaktadır. CUDA kernellerini derlemek için nvcc adlı nVidia derleyici sürücüsünün çalışması gerekmektedir ve bu kernellere olan referanslar MSVC derleyicisi tarafından algılanmamaktadır, dolayısıyla kernellerin metod işaretçilerini doğrudan kullanamıyoruz.

Kernelleri çağırmak için gereken griddeki blok ve bloktaki thread sayısını kullanarak kerneli çağıran metodlara proje kapsamında “kernel launcher” adı verilmiştir. Kernel launcherlar, aynı zamanda ptKernelLauncher tipindeki metod işaretçisi imzasını taşırlar.

/**
	Yaratıcıda alınan kerneli çağırır.

	\param imageData Görüntünün BGR kanal sıralı bellekteki adresi.

	Görüntüyü normalize ederek kernelLauncher işaretçisinin gösterdiği metodu çağırır ardından görüntüyü denormalize eder( [0, 255] aralığına ).
	Kernelde işlenen görüntüden sonuç olarak [0, 1] aralığı dışında bir değer dönerse o kanalın değeri [0, 255] aralığından dışarıda olabilir. Bu durumda değer yakın olduğu sınıra indirgenir.
*/
virtual void FilterImage(char* imageData)
{
	// imageData değişkenindeki görüntü verisi normalize edilerek h_Image değişkenine aktarılır.
	for(int i=0; i<3*width*height; i++)
	{
		*(h_Image + i) = (unsigned char)*(imageData + i) / 255.0f; // normalize and copy image
	}

	/*
		Görüntü GPU belleğine kopyalanır.
	*/

	cudaMemcpy( d_Image, h_Image, 3 * sizeof(float) * width * height, cudaMemcpyHostToDevice );
	checkCUDAError("FilterImage: memcpy");

	/*
		Constructorda verilen kernel çalıştırılır.
	*/
	kernelLauncher( d_Image, width, height );

	/*
		Sonuçlar CPU belleğine kopyalanır.
	*/
	cudaMemcpy( h_Image, d_Image, 3 * sizeof(float) * width * height, cudaMemcpyDeviceToHost);
	checkCUDAError("FilterImage: memcpy2");

	/*
		h_Image değişkenindeki normalize edilmiş görüntü verisi [0, 255] aralığına çekilir.
	*/
	for(int i=0; i<3*width*height; i++)
	{
		*(imageData + i) = satchar(*(h_Image + i) * 255);
	}
}

Projedeki kernelleri çağırmak için kullanılan metod işaretçisi ve kullanımı:

typedef void (*ptKernelLauncher)(float*, int, int);

ptKernelLauncher kernelLauncher; // kernelLauncher değişkeni tanımlanır.

kernelLauncher = deviceInvertLaunch; // deviceInvertLaunch metodunun işaretçisi tanımlanan değişkene atanır.

kernelLauncher( d_Image, width, height ); // metod, işaretçisi üzerinden çağrılır. Çağrılan metod, GPU üzerinde çalışacak olan kerneli configuration parametreleri ile çağırır.

Kerneli parametreleri ile çağırmak:

    dim3 dimBlock( BLOCK_SIZE, BLOCK_SIZE );
    dim3 dimGrid( height / dimBlock.x, width / dimBlock.y );

    gpuInvert<<< dimGrid, dimBlock >>>( d_Image, width, height);

Görüntü Negatifi – CUDA

CUDA kerneli kullanarak görüntünün negatifini almak için kernelimizde görüntünün her bir pikselinin tersini almalıyız. Bu işlem için kernel aşağıdaki gibidir.

/**
	Görüntünün tersini alan kernel.

	\param image [0, 1] aralığına normalize edilmiş, BGR kanal sıralı görüntünün GPU belleğindeki adresi.
	\param width Görüntünün piksel olarak genişliği
	\param height Görüntünün piksel olarak yüksekliği

	Metod GPU üzerinde çalışır, çıktısını image parametresinin üzerine yazar.

	*/
__global__
void gpuInvert(
	float* image,
	int width,
	int height
	)
{
	int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
	int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;

	int cIdx = ( row * width + col ) * 3; // 3 ile çarpım RGB için, linearIndex.

	// normalize edilmiş pikselleri 1'den çıkarttığımızda görüntünün negatifini almış oluruz.
	*( image + cIdx     ) = 1 - *( image + cIdx     ); // Blue kanalı
	*( image + cIdx + 1 ) = 1 - *( image + cIdx + 1 ); // Green kanalı
	*( image + cIdx + 2 ) = 1 - *( image + cIdx + 2 ); // Red kanalı
}

Bu kernelin kullanımı iki şekildedir. Birincisi SingleCudaFilter sınıfını doğrudan main.cpp içinde kullanarak:

// deviceInvertLaunch kernel launcher metodu SingleCudaFilter sınıfına parametre olarak verilir.
ISingleImageFilter* myFilter = new SingleCudaFilter(deviceInvertLaunch);

İkincisi ise bu parametreyi kendi constructorında SingleCudaFilter sınıfının constructoruna veren bir sınıf yazmaktır. Tek bir filtre için tek bir sınıf yazmanın getirisi, hangi kernelin kullanılacağını bilmeye gerek olmamasıdır, ptKernelLauncher, sınıfın içinde bir kere yazılır.

/**
	CUDA kullanarak GPU üzerinde görüntünün negatifini alan filtre sınıfı.
*/
class CudaInvertFilter : public SingleCudaFilter
{
public:
	/**
		\ref deviceInvertLaunch metod işaretçisi parametresi ile SingleCudaFilter yaratıcısını çağırır.
	*/
	CudaInvertFilter()
		: SingleCudaFilter(deviceInvertLaunch)
	{
	}
};

Texturelar

Buraya kadar bahsedilen CUDA filtreleri, tamamen ekran kartı global belleği üzerinden çalışmaktadır. Okumalar texture belleğinden yapılarak zaman kazanılabilir ve texture önbelleğinden otomatik olarak faydalanılabilir. Texture önbelleği, erişimde yerellik olduğu zaman örneğin bir piksele eriştiğinizden hemen ardından o pikselin 2 boyutlu düzlemdeki komşularına eriştiğinizde devreye girerek bellek okuma işlemlerinin bir miktar hızlanmasını sağlayabilir.

Textureları CUDA ile kullanabilmek için cuArray tipinde tanımlanan değişken üzerinden işlemler yürütülür. Texture üzerine hiçbir zaman doğrudan veri kopyalanmaz çünkü texturelar salt-okunurdur. Bir CUDA textureına cuArray atanarak yazma işlemleri cuArray‘e yapıldığında, texture otomatik olarak yazılan veriyi üzerine alacaktır.

Texturelar ile ilgili karşılaşılan sorun, kullanılacak textureın, kernelin bulunduğu .cu dosyasında olması zorunluluğudur. Ortak bir başlık dosyasında texture tanımlanıp hem kernelin bulunduğu dosyadan, hem de bu textureı yönetmek istediğimiz sınıfın dosyasından içerildiğinde, texturelar nvcc tarafından farklı olarak algılanmaktadır.

Texture değişkeninin scopeu:

texture degiskeninin scopeu

Kernelin bulunduğu dosya dışındaki bir sınıftan texturea veri yüklemesi ve diğer texture ayarlarının yapılması için nVidia’nın CUDA Driver API’si kullanılabilir. Adı sizi korkutmasın, tek satırda yapılan işlemi başka bir dosyadan yapmak istediğimizde bir kaç satırda yapıyoruz, tek fark bu.

SingleCudaTexFilter sınıfında texture yönetimi ile ilgili bir diğer ince nokta da CUDA kodlarında cudaGetTextureReference ve benzeri metodların parametre olarak const bir referans beklemesidir. Bizim niyetimiz textureı değiştirmek olduğundan CUDA Driver API’si ile anlaşabilmek için ara bir değişken kullanıyoruz.

cuda_runtime_api.h içinde cudaGetTextureReference tanımı:

extern __host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol);

SingleCudaTexFilter içerisinde textureın const‘unun kaldırılarak ara değişkene atanması ve kullanımı:

	const textureReference* constTexRefPtr;
	textureReference* texRefPtr;

	...

	// cudaGetTextureReference, istediği gibi const texture referansı ile çağrılır.
	cudaGetTextureReference(&constTexRefPtr, textureSymbolName);
	checkCUDAError("get texture reference");

	// const olan referans, const olmayan referansa dönüştürülerek saklanır.
	texRefPtr = const_cast<textureReference*>( constTexRefPtr );
	channelDesc = cudaCreateChannelDesc<float4>();

	// const olmayan texture referansının kullanımı.
	cudaMallocArray( &cu_array, &texRefPtr->channelDesc, width, height );
	checkCUDAError("malloc device image");

	...

	// cu_array üzerine kopyalanan veriler texture atanacak.
	cudaMemcpyToArray( cu_array, 0, 0, h_Image, sizeof(float4) * width * height, cudaMemcpyHostToDevice);
	checkCUDAError("FilterImage: memcpy");

	// cu_array, texturea atanır.
	cudaBindTextureToArray( texRefPtr, cu_array, &texRefPtr->channelDesc );

Texture kullanımı ile ilgili olarak SingleCudaTexFilter.cu ve SingleCudaTexFilter.h dosyalarını inceleyebilirsiniz.

Görüntü Negatifi – CUDA Texture

CUDA ile texture belleğini kullanarak görüntünün negatifini almak için kernelimizde görüntünün her bir pikselini textureın her bir pikselinin tersinden oluşturmalıyız. Bu işlem için kernel aşağıdaki gibidir.

/**
	Texture kullanarak görüntünün negatifini alan kernel.

	\param image [0, 1] aralığına normalize edilmiş, BGR kanal sıralı görüntünün GPU belleğindeki adresi.
	\param width Görüntünün piksel olarak genişliği
	\param height Görüntünün piksel olarak yüksekliği

	Metod GPU üzerinde çalışır, çıktısını image parametresinin üzerine yazar.
*/
__global__
void gpuTexInvert(
	float* image,
	int width,
	int height
	)
{
	int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
	int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;

	int cIdx = ( row * width + col ) * 3; // 3 ile çarpım RGB için, linearIndex.

	// threade ait kordinatın texture uzayındaki kordinatı bulunur.
	float tu = (float)col / width;
	float tv = (float)row / height;

	// Texture üzerinden görüntü verisi okunur.
	float4 texVal = tex2D( texInvert1, tu, tv );

	// Texture değerleri 1'den çıkartılarak global belleğe yazılır.
	*( image + cIdx )     = 1 - texVal.x;
	*( image + cIdx + 1 ) = 1 - texVal.y;
	*( image + cIdx + 2 ) = 1 - texVal.z;
}

Texture kullanan kernel de texture kullanmayan CUDA negatif kerneli gibi iki şekilde kullanılabilir. Detaylar için kaynak kodu inceleyebilirsiniz.

Filtre Zinciri: SingleImageFilterChain

Filtre zinciri, birbiri ardına uygulanacak filtreleri barındıran bir sınıftır. Bu sınıf SingleImageFilter sınıfından kalıtlanmıştır. AppendFilter metodu ile sınıfa işlemesi için filtreleri sıralı olarak verebilirsiniz. FilterImage metodu çağrıldığında zincirdeki tüm filtreler için FilterImage metodu çağrılacaktır.

ISingleImageFilter* myFilter1 = new SingleCudaTexFilter(deviceTexAbsDiffLaunch, "texAbsDiff1");
ISingleImageFilter* myFilter2 = new CpuInvertFilter();

SingleImageFilterChain* myFilter = new SingleImageFilterChain();
myFilter->AppendFilter( myFilter1 );
myFilter->AppendFilter( myFilter2 );

// myFilter'a atanmış olan filtre zinciri, ISingleImageFilter arayüzünü gerçeklediği için diğer filtreler gibi kullanılabilir.

UML Şeması

UML Şeması basitleştirilmiş olarak verilmiştir:

UML Şeması

Filtre Örnekleri

IdentityFilter:
IdentityFilter

CudaSepiaFilter:
CudaSepiaFilter

CudaTexBoxBlurFilter:
CudaTexBoxBlurFilter

SingleImageFilterChain: deviceTexAbsDiffLaunch kerneli ardından CpuInvertFilter:
SingleImageFilterChain

İndirmeler

Dosyalar 7-zip ile sıkıştırılmıştır.

NOT: Kerneller optimize edilmemiştir.