CPU üzerinde Connected Component Labelling

Önceki CUDA ile OpenCV kullanarak Webcam Görüntü İşleme 
başlıklı yazımızda bahsedilen projeye ekleme bir filtre olarak gerçekleştirilen Connected Component Labeling (Birleşik Eleman Etiketleme) filtresini OpenCV ile webcam görüntü işlemede nasıl kullandığımıza bakalım.

NOT: Bu yazıdaki kod örnekleri kısaltılarak verilmiştir. Kod belgeleri için Doxygen ile yaratılmış olan mikrositeyi inceleyebilirsiniz (yazıdaki eklemeleri içerecek şekilde güncellenmiştir).

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

İçindekiler

  • Gerekenler
  • CCL Nedir?
  • CCL için Önişleme
  • CPU CCL Filtresi: CpuCCLFilter
  • Filtre Örnekleri
  • İndirmeler

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.
  • Önceki yazımızda detayları bulabilirsiniz: CUDA ile OpenCV kullanarak Webcam Görüntü İşleme

CCL Nedir?

Connected Components Labeling, görüntüde birleşik olan nesneleri etiketleyerek birbirinden ayırdedebilecek hale getiren algoritmaya verilen addır. Burada en önemli tanım, “görüntüde birleşik” tamlamasıdır. Bu filtredeki kodlar, wikipediadaki CCL açıklamasından yola çıkılarak yazılmıştır.

Görüntünün birleşik olması, komşu piksellerin renklerinin aynı veya belli bir derecede yakın olması ile ölçülür. Komşulukları gözönüne alırken 4-komşu ve 8-komşu modelleri kullanılabilir. 4-komşu modeli, ana yönler olan kuzey, güney, doğu ve batı komşularıdır. 8-komşu modelinde ara yönlerdeki komşular da hesaba katılır.

4-komşuluk

8-komşuluk

CCL için Önişleme

Renk Benzerliği

Renk benzerliği, CCL filtremizde birbiri ile aynı olarak nitelendirilecek renkleri ayırdetmemize yarar. Basit olarak gri skala uzayında uzaklığı 10’dan küçük olan renkler aynı sayıldı. Standart bir webcamin yakaladığı görüntülerde (en azından benimkinde) RGB noise denilen her renkten gürültü vardır. Webcam sabit bir yere bakarken ve görüntüde hiç hareket yokken bile yakalanan resmin pikselleri kısıtlı bir aralıkta sürekli değişir. Renklerin aynı olmasını katı bir şekilde gerektirmeyerek saptanan etiket sayısını düşürüyoruz, bir bakıma bu gürültüye bağışıklık kazanıyoruz.

Thresholding

Thresholding, belli bir aralığa sahip değeri, iki adet değere indirgemenin bir yoludur. Örneğin değerimiz 0 ile 255 arasında ise, threshold değerimiz 90 ise, 90’dan küçük değerler 0, büyük ve eşit değerler 1 grubuna denk gelir. CCL, threshold işlemi görmüş görüntü üzerinde çalışır. CCL algoritması için 0 arkaplan, 1 nesne demektir.

CPU CCL Filtresi: CpuCCLFilter

ISingleImageFilter arayüzünü kullanarak yazılan filtrede en önemli metod olan findConnectedComponents kısaca anlatılarak gösterilmiştir.

/**
	\ref ptKernelLauncher tipinde metod.

	\param width Görüntünün piksel olarak genişliği
	\param height Görüntünün piksel olarak yüksekliği
	\param rowStride Görüntünün hafızadaki byte olarak genişliği.
	\param imageData Görüntü verisi.
	\param outImageData Etiket görüntü verisi.
	\param outLabelBitmap Etiket integer bitmap.

	Görüntüdeki birleşmiş birimleri bulur. Connected Component Labeler.
*/
template<bool eightConnected>
void CpuCCLFilter::findConnectedComponents(int width, int height, int rowStride, char* imageData, char* outImageData, int* outLabelBitmap)
{
	// label bitmap	

	// disjoint set için tip ve map tanımları.
	typedef std::map<short, std::size_t> rank_t; // element order.
	typedef std::map<short, short> parent_t;

	rank_t rank_map;
	parent_t parent_map;

	boost::associative_property_map<rank_t>   rank_pmap( rank_map );
	boost::associative_property_map<parent_t> parent_pmap( parent_map );

	// disjoint sets yaratılır.
	boost::disjoint_sets<boost::associative_property_map<rank_t>, boost::associative_property_map<parent_t>> 
		ds( rank_pmap, parent_pmap );

	boost::unordered_set<short> elements;

	short labelNumber = 0;

	for(int i = 0; i < height; i++)
	{
		for(int j = 0; j < width; j++)
		{
			outLabelBitmap[i*rowStride + j] = 0;
		}
	}

	// tüm pikseller gezilir
	for(int i = 0; i < height; i++)
	{
		for(int j = 0; j < width; j++)
		{
			// if pixel is not background.
			if( imageData[i*rowStride + j] != 0 )
			{
				if(j > 0 & i > 0 &
					isSameColor( imageData[i*rowStride + j - 1], imageData[(i - 1)*rowStride + j]) // same value
					& isSameColor( imageData[i*rowStride + j - 1], imageData[i*rowStride + j] )
				)
				{
					// west ve north aynı
					ds.union_set( outLabelBitmap[i*rowStride + j - 1], outLabelBitmap[(i - 1)*rowStride + j] );
					outLabelBitmap[i*rowStride + j] = outLabelBitmap[i*rowStride + j - 1];
				}

				if(j==0 & i > 0 & isSameColor( imageData[i*rowStride + j], imageData[(i - 1)*rowStride + j] ))
				{
					outLabelBitmap[i*rowStride + j] = outLabelBitmap[(i-1)*rowStride + j];
				}

				// West pixel
				if(j > 0 & isSameColor( imageData[i*rowStride + j - 1], imageData[i*rowStride + j] ))
				{
					// west pixel ile aynı label.
					outLabelBitmap[i*rowStride + j] = outLabelBitmap[i*rowStride + j - 1];
				}

				// west farklı north aynı.
				if(
					(j > 0 & !isSameColor( imageData[i*rowStride + j - 1], imageData[i*rowStride + j] )) // west different value
					& (i > 0 & isSameColor( imageData[(i-1)*rowStride + j], imageData[i*rowStride + j] )) // north same value					
				)
				{
					// north ile aynı
					outLabelBitmap[i*rowStride + j] = outLabelBitmap[(i-1)*rowStride + j];					
				}

				// west ve north farklı
				if(   ((j > 0 &;amp;& !isSameColor( imageData[i*rowStride + j - 1], imageData[i*rowStride + j] )) || j == 0) // west different value
					& ((i > 0 & !isSameColor( imageData[(i-1)*rowStride + j], imageData[i*rowStride + j] )) || i == 0) // north different value					
				)
				{
					labelNumber++;
					ds.make_set(labelNumber);
					elements.insert(labelNumber);
					outLabelBitmap[i*rowStride + j] = labelNumber;
				}

				// northeast ve northwest kontrol edilir. (çaprazlar).
				if(eightConnected)
				{
					// northwest
					if(j > 0 & i > 0
						& isSameColor( imageData[(i-1)*rowStride + j - 1], imageData[i*rowStride + j] )
					)
					{
						// northwest ile aynı
						ds.union_set( outLabelBitmap[(i-1)*rowStride + j - 1], outLabelBitmap[i*rowStride + j] );
					}

					// northeast
					if(j+1 < width & i > 0
						& isSameColor( imageData[(i-1)*rowStride + j + 1], imageData[i*rowStride + j] )
					)
					{
						// northeast ile aynı
						ds.union_set( outLabelBitmap[(i-1)*rowStride + j + 1], outLabelBitmap[i*rowStride + j] );
					}
				}
			}
		}
	}

	int cnt = ds.count_sets(elements.begin(), elements.end());

	printf("Component count: %i\n", cnt);

	// second pass - label output image and colorize.
	for(int i = 0; i < height; i++)
	{
		for(int j = 0; j < width; j++)
		{
			int labelNo = ds.find_set( outLabelBitmap[i*rowStride + j] ); // pikselin etiketi bulunur.
			int R =0, G=0, B=0;

			// etiketler renklendirilir.
			hsl_to_rgb( 1.0f * labelNo / labelNumber, .8f + .2f * labelNo / labelNumber, .75f, &R, &G, &B );

			int idx = i * rowStride + j;

			outImageData[idx * 3 + 0] = (char)(B);
			outImageData[idx * 3 + 1] = (char)(G);
			outImageData[idx * 3 + 2] = (char)(R);
		}
	}	
}

 

Filtre Örnekleri

CpuCCLFilter:

İndirmeler

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

CUDA ile Matris Çarpımı – 1

Matris Çarpımı CUDA

Matrislere iki boyutlu vektörler dersek vektörleri tek boyutlu değişkenler gibi görmek, değişkenlere de “boyutsuz” benzetmesini yapmak yerinde olacaktır. Matris nedir sorusunun cevabını aramayacağımız bu yazımızda matris çarpımını C dilinde CPU üzerinde ve GPU üzerinde en basit haliyle nasıl yapabileceğimizi göreceğiz.
Örnek kodlarda kodun basitliği açısından kare matrisler kullanılmıştır.
A: M x N ( M satır, N sütun),
B: N x R
C: M x R boyutlarında seçilmiştir.
Örnek kodlarda matrislerin boyutları #define ile tanımlanacaktır fakat bu yazımızdaki örneklerde M = N = R olarak seçilmelidir.

Tanımlar:

A_{i,j} A matrisinin i. satırının j. sütunundaki eleman
C = AB
C_{i,k}=\sum^N_{j=1}A_{i,j} B_{j,k}
Bu yazımızdaki örnekler:
  • C ile yazılmış naif matris çarpımı
  • Optimize edilmemiş GPU kerneli
  • 1. Optimizasyon
Optimizasyon çok basit olarak yapılmıştır öyleki optimizasyon demeye dilim varmıyor. İleriki yazılarda kernel adım adım optimize edilerek GPU’nun suyu sıkılacaktır.

Naif CPU Matris Çarpımı

C matrisinin her elemanı için A matrisinin N sütunu ve B matrisinin N satırı çarpılarak toplanır. M x N x R kere loop döner ( O(n^3), n kare matrisimizin bir boyutu )

inline
float* h_loc(float* h_B, int width, int i, int j)
{
	return h_B + i*width + j;
}

void hostMatrixMultiply(
	float* h_C,
	float* h_A,
	float* h_B,
	int m,
	int n,
	int r)
{

	clock_t start = clock(), diff;

	for(int i=0; i<m; i++)
	{
		for(int k=0; k<r; k++)
		{
			float sum = 0;
			for(int j=0; j<n; j++)
			{
				sum += *h_loc(h_A, n, i, j) * *h_loc(h_B, r, j, k);
			}

			*h_loc(h_C, r, i, k) = sum;
		}
	}

	diff = clock() - start;

	int msec = diff * 1000 / CLOCKS_PER_SEC;
	printf("hostMatrixMultiply: %i ms\n", msec);

}

Bu algoritma, matris çarpımının “text-book” çözümüdür.

Optimize Edilmemiş GPU Kerneli

deviceMultiply1 fonksiyonunun çağırdığı matMul1 kerneli, performans hakkında en ufak tedirginlik hissedilmeden yazılmış bir matris çarpım kernelidir. deviceMultiply1 girizgah fonksiyonunda GPU üzerinde hafıza ayrılmakta, matrisler GPU hafızasına kopyalanmakta ve hesaplanan sonuç CPU hafızasına alındıktan sonra GPU üzerinde ayrılan hafızalar bırakılmaktadır. Sonucun hesaplanması matMul1 kernelinde gerçekleşir.

deviceMultiply1 fonksiyonunda matMul1 kerneli çağrılmadan önce grid ve block boyutları hesaplanır. GPU üzerinde çalışacak thread sayısının C matrisindeki eleman sayısı kadar olmasını sağlayacak değerler hesaplanır. Daha sonra matMul1<<< dimGrid, dimBlock >>>( d_C, d_A, d_B, m, n, r); satırıyla kernel çağrılır.

__global__
void matMul1(
	float* d_C,
	float* d_A,
	float* d_B,
	int m,
	int n,
	int r)
{
	int i = blockIdx.x * BLOCK_SIZE + threadIdx.x; // GPU threadinin hesaplayacağı C matrisinin satırı.
	int k = blockIdx.y * BLOCK_SIZE + threadIdx.y; // GPU threadinin hesaplayacağı C matrisinin sütunu

	int cIdx = i*m + k; // Hesaplanacak elemanın C indisi.

	d_C[ cIdx ] = 0;

	// C matrisinin her bir hücresi için
	for(int j=0; j<n; j++)
	{
		// her bir eleman için C matrisinin bir elemanı okunur ve yazılır.
		d_C[ cIdx ] += *loc( d_A, n, i, j ) * *loc( d_B, r, j, k );
	}
}

Memory access sayısının düşürülmesi, memory coalescing, tiling gibi hiçbir teknik kullanılmadan yazılmış olan naif CUDA matris çarpım kernelini birlikte inceleyelim.

1. Optimizasyon (Hafıza Erişimi)

matMul1 kernelinde yapılacak çok ufak bir değişiklik ile performans arttırılabilir. Bu değişikliğin getirdiği hız kazancı, GPU’da hafıza erişim işlemlerinin CPU hafızasına erişimden çok daha yavaş olması ile gün ışığına çıkmaktadır.
matMul1 kernelinde d_C, C matrisinin GPU hafızasındaki yerine işaretçidir(pointer). Bu işaretçinin kullanıldığı her yer GPU hafızasına erişmek demektir. İşaretçi başta sıfırlanmıştır ve toplama işlemi işaretçi üzerinde yapılmıştır. += operatörü, d_C[ Idx ] adresinin önce okunması, toplam işleminden sonra da tekrar yazılması demektir. Yani d_A ve d_B’yi saymazsak for döngüsünde toplam 2*n kere hafıza erişimi yapılmaktadır. Halbuki C matrisinin değerini hesaplayıp tek seferde hafızaya yazabiliriz.
matMul2 kernelinde yaratılan değişken ile global GPU RAM erişimi tek thread başına 1 adet yazmaya (d_A ve d_B hariç) düşmüştür.
__global__
void matMul2(
	float* d_C,
	float* d_A,
	float* d_B,
	int m,
	int n,
	int r)
{
	int i = blockIdx.x * BLOCK_SIZE + threadIdx.x;
	int k = blockIdx.y * BLOCK_SIZE + threadIdx.y;

	int cIdx = i*m + k;

	float val = 0; // ara toplam değişkeni

	// C matrisinin her bir hücresi için
	for(int j=0; j<n; j++)
	{
		val += *loc( d_A, n, i, j ) * *loc( d_B, r, j, k );
	}

	// C matrisine bir kere yazılır.
	d_C[ cIdx ] = val;
}

Zamanlama

Yapılan bu ufak değişiklikle kazanılan 128 x 128 boyutlarında iki matrisin çarpım sürelerini karşılaştırabiliriz.
matMul1 ve matMul2 zamanlaması

Proje dosyasını açabilmek için farklı kaydederek uzantısını 7z yapınız.
Kaynakları indir: CUDA matris carpimi 1

Visual Studio 2010 ile CUDA SDK 3.2 Projesi Yaratmak

CUDA Projelerine Giriş

For English: CUDA 3.2 on VS2010 in 9 steps

Bu kısa yazımızda CUDA ile programlamaya giriş yapmanın zevkini yaşayacağız. CUDA ile ekran kartındaki GPU üzerinde çalışacak programlar yazabilirsiniz. Bu programlar GPU’daki özel hardware threadler üzerinde paralel olarak çalışırlar. Detaylı bilgi için tıklayınız.

4. adımdaki dosyalar yok ise indiriniz: CUDA 3.2 Build Rules

Aşağıdaki adımları izlemek için gerekenler:

  • Visual Studio 2008 kurulumu
  • Visual Studio 2010 kurulumu
  • nVidia CUDA 3.2 SDK kurulumu
  • Genel Visual Studio menüleri ve kullanımı

9 Adımda CUDA Derleme

    1. Yeni bir Win32 Console Application yaratın.

    1. Projeyi yaratırken empty project seçeneğini işaretleyin. Boş bir projeye kaynak dosyasını kendimiz ekleyeceğiz.

    1. cu uzantılı kaynak dosyası ekleyin. Bu dosya nvidianın compiler driverı olan nvcc.exe dosyası tarafından işlenerek vs2008 ile gelen versiyon 9 C derleyicisine gönderilecektir.

    1. CUDA 3.2 SDK’nın kuruluşuile gelen Build Customizationların yerinde olduğu kontrol edilir.

    1. Platform Toolset v90 olarak değiştirilir. Visual Studio 2010, CUDA dosyalarını derlemeyi destekler fakat altyapı olarak bir önceki versiyon C derleyicisini kullanır. Dolayısıyla Visual Studio 2008’in makinada kurulu olmasına ihtiyaç vardır.

    1. Eklediğimiz CU dosyasının özelliklerinden Item Type olarak CUDA C/C++ seçilir.

    1. Project Build Customizations olarak CUDA 3.2 seçilir. Bende CUDA 3.1 SDK’sı da yüklü olduğu için o da seçilebiliyor.

    1. CUDA’nın kullandığı kütüphaneler eklenir( cuda.lib ve cudart.lib ).

    1. Yaratılan CU dosyasından CUDA fonksiyonlarının çağrılabildiği ve derlenebildiği görülür.