CUDA ile Matris Çarpımı – 2

Matris Çarpımı CUDA – 2

Bir önceki yazımızda CUDA ile matris çarpımının nasıl yapılabileceğini en basit hali ile görmüştük. Burada geçen sefer kaldığımız yer olan hafıza erişim optimizasyonuna devam ediyoruz.
En son matMul2 isimli kernel içerisinde toplam alınan satırda hafızaya yapılan yazmaları azaltmıştık.

matMul1 (n adet global store):

d_C[ cIdx ] = 0;

// C matrisinin her bir hücresi için
for(int j=0; j
matMul2 (1 adet global store):

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

// C matrisinin her bir hücresi için
for(int j=0; j

Paylaşık Hafıza (Shared Memory)

Paylaşılmış olmayan hafıza mı vardır dediğinizi duyar gibiyim. Global Hafıza da paylaşılmıştır elbet, herkes yazabilir herkes okuyabilir. Fakat burada bir nVidia terimi olarak geçen "paylaşık hafıza" bir blok içerisindeki threadler arasında paylaşılmıştır. Yani ancak aynı blok içerisindeki threadler aynı yeri okuyabilir ve aynı yere yazabilir.

Bir bloktaki threadler paylaşık hafızaya yardımlaşarak yükledikleri verileri kullanabilirler.

...

Ais[threadIdx.y][threadIdx.x] = *loc( d_A, n, i, j ); // A(i,j) elemanı Ais paylaşık hafızasına yüklenir.
Bis[threadIdx.y][threadIdx.x] = *loc( d_A, n, j, k ); // A(j,k) elemanı Bis paylaşık hafızasına yüklenir.

// kullanılacak elemanlar paylaşık hafızaya yüklenir.
__syncthreads(); // syncthreads, bloktaki tüm threadlerin bu noktaya erişmesini bekletir.

for(int k=0; k

2. Optimizasyon (Paylaşık Hafıza)

Yapılan optimizasyonda paylaşık hafıza kullanılarak global hafızadan okuma sayısı düşürülmüştür.

__global__
void matMul3(
	float* d_C,
	float* d_A,
	float* d_B,
	int m,
	int n,
	int r)
{

	__shared__ float Ais[BLOCK_SIZE][BLOCK_SIZE];
	__shared__ float Bis[BLOCK_SIZE][BLOCK_SIZE];

	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

	for(int j=0; j < n / BLOCK_SIZE; j++)
	{
		Ais[threadIdx.y][threadIdx.x] = *loc( d_A, n, i, j );
		Bis[threadIdx.y][threadIdx.x] = *loc( d_B, n, j, k );

		// kullanılacak elemanlar paylaşık hafızaya yüklenir.
		__syncthreads();

		for(int k=0; k

Zamanlama

512 x 512 boyutlarında iki matrisin çarpım sürelerini karşılaştırabiliriz.

Versiyon Hız (ms) Hızlanma
Host 9619 1x
MatMul1 509 19x
MatMul2 183 53x
MatMul3 72 135x

Referanslar

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