Merhabalar,

Yuksek lisans tezimde CPU ve C++ ile yaptigim mikrofon dizisi projesini GPU ve D gecirerek gelistirmeye baslicagimizi konusmustuk Ali Abi ile.

Ben DCompute'u derledim. LDC 1.4 u indirip bir kac degisiklik yapmak yetti. Eger LDC 1.4 cikmis olmasa idi LDC yi belirli LLVM opsiyonlari ile derlemem gerekiyordu fakat gerek kalmadi.

DCompute'u derledim fakat CPU ve C++ dan direk GPU ve D gecmektense bir ara adim olarak GPU ve C++ ile basladim. Cunku ben GPU cok kodlamadim. Benim deneyimsizligim ile DCompute'un deneyimsizligi ayni anda yasamaktansa olgunlasmis C++ kullanmak ilk adim icin daha iyi oldugu dusunuyorum.

Proje surda gorulebilir : https://github.com/kerdemdemir/CUDABeamformer/blob/master/CudaBeamformer/kernel.cu

C++ kismini hizli gecmeye calistigimdan cok ozenmedim oldukca hizli yapmaya calisiyorum.

CUDA kernel kodu su ornege biraz benziyor:
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory

Su anki durumum kernel derleniyor fakat calisirken hata veriyor bende GPU tarafinda calisan kodu nasil debug edicegimi bilmedigimden bulamadim simdilik. DCompute veya Stackoverflow forumlarinda soracagim.

Bunu calistirdiktan sonra GPU icinde L1 cache var shared memory diyorlar onu kullanip bir daha benchmark yapacagim. Sonra D ye gecis baslayacak. CPU -GPU - GPU L1 cache zamanlama karsilastirmalarini yapcagim. Suan da bu is CPU ile 88100 ornekden olusan bir ses dosyasi icin 900ms suruyor.

Durumlar bunlar calismaya devam.
Saygilar
Erdem

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

C++ ile GPU beamforming işini tamamladım. Sonuçlar gerçekten çok iyi

CPU ~ 1.1 saniye GPU ~ 0.018 saniye. Daha L1 cache'ini kullanmadım GPU'nun yinede sonuçlar şimdiden çok iyi. Biraz heyecan yarattı iyi sonuçlar L1 adımını geçip direk D 'ye geçirme işine başlıyacağım.

GPU kerneli:

struct MicParamsGpu
{
	float *outputData;
	float *rawData;
	int    packetSize;
	float *leapData;
	int    leapStride;
	int    *delays;
	int    arraySize;
	int    stride;
};

__device__ float GetElement(MicParamsGpu params, int curMic, int index ) {
	if (params.packetSize > index)
	{
		return params.rawData[curMic*params.packetSize + index];
	}
	else
	{
		int leapIndex = (index - params.packetSize) + curMic*params.leapStride;
		return params.leapData[leapIndex];
	}
}

__global__ void beamformKernel2(MicParamsGpu params)
{
	int xIndex = threadIdx.x;
	int currentStartIndex = xIndex * params.stride;

	for (int k = 0; k < params.stride; k++)
	{
		float curVal = 0;
		for (int i = 0; i < params.arraySize; i++)
		{
			curVal = GetElement(params, i, currentStartIndex + k + params.delays[i]);
		}
		params.outputData[currentStartIndex + k] = curVal;
	}
}

Cuda Memory Allocation:

// Helper function for using CUDA to add vectors in parallel.
cudaError_t beamformWithCudaHelper(MicrophoneArray& array, SharpVector& outputData)
{
	cudaError_t cudaStatus;
	MicParamsGpu params;
	params.arraySize = array.micropshoneList.size();
	params.packetSize = Config::getInstance().packetSize;
	params.leapStride = Config::getInstance().getMicMaxDelay()*2;
	cudaMalloc(&params.rawData, array.micropshoneList.size() * sizeof(float) * params.packetSize);
	cudaMalloc(&params.leapData, array.micropshoneList.size() * sizeof(float) * params.leapStride);
	cudaMalloc(&params.delays, array.micropshoneList.size() * sizeof(int));


	cudaMalloc(&params.outputData, Config::getInstance().packetSize * sizeof(float) );
	std::vector<int> delayVec;
	params.stride = params.packetSize / 1000;
	cudaStatus = cudaGetLastError();
	for (int i = 0; i < params.arraySize; i++)
	{
		cudaMemcpy( params.rawData + i * params.packetSize, array.micropshoneList[i].getData().data(),
						params.packetSize* sizeof(float), cudaMemcpyHostToDevice);
		cudaStatus = cudaGetLastError();
		cudaMemcpy(params.leapData + i * params.leapStride, array.micropshoneList[i].getLeapData().data(),
						params.leapStride* sizeof(float), cudaMemcpyHostToDevice);
		cudaStatus = cudaGetLastError();
		delayVec.push_back(array.micropshoneList[i].getDelay(1000, 45) + Config::getInstance().getMicMaxDelay());

	}
	cudaStatus = cudaGetLastError();
	cudaMemcpy(params.delays, delayVec.data(), delayVec.size() * sizeof(int), cudaMemcpyHostToDevice);

	float startTime = get_time();
	// Launch a kernel on the GPU with one thread for each element.
	beamformKernel2 << <1, 1000 >> >(params);

	// Check for any errors launching the kernel
	cudaStatus = cudaGetLastError();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
	}

	// cudaDeviceSynchronize waits for the kernel to finish, and returns
	// any errors encountered during the launch.


	cudaMemcpy(outputData.data(), params.outputData, Config::getInstance().packetSize * sizeof(float), cudaMemcpyDeviceToHost);
	float endTime = get_time();
	std::cout << "GPU Time spent: " << endTime - startTime;
	return cudaStatus;
}

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

Hız farkı büyük! :)

Ali

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

D tarafında CPU ile beamform işine başladım.

https://github.com/kerdemdemir/DSharpEar/blob/master/DSharpEar/main.d

Beamformun ne olduğunu anladan uzun bir yazı hazırlamıştım fakat browser penceresini yanlışlıkla kapatınca hepsi uçtu. Artık yarın yazarım bir daha.

Şimdi DCompute kullanarak GPU işlerine geldi sıra. DCompute .lib 'ini projeme nasıl alıcam onun derdindeyim.
DCompute'u sadece dcompute kaynak kodundaki testlerde derleyebiliyorum.

Eğer projede güzelleştirceğim yerler varsa yorumlarınızı bekliyorum.

Bu GPU işi bittikten sonra bunun arayüzünü yazıcam vibe.d kullanıcam three.js ile opengl ile browser'dan çalışacak umarım.

Bakalım bakalım.
Erdemdem

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

Alıntı (kerdemdemir):

Ben herşeyi bir Emacs penceresinde yazıyorum ve tarayıcıya oradan kopyalıyorum. (Çok kısa istisnalar hariç.)

Ali

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

C++ GPU kısmını DCompute'e geçirme kısmında tıkandım. DCompute'u bir türlü derleyemedim Windows'da ayrı uğraştım Linux'dede ayrı uğraştım(sanıyorsam 1-2 haftadır uğraşıyorum birazda soğudum). D Learn forum'da bir başlık açmıştım fakat çok fazla soru sorarak insanları sıkmak istemiyorum. Belki bir süre sonra geri dönerim bu konuya.

DCompute'un kaynak kodu derlenmesi rağmen, dependency olarak eklendiğin derlenmiyor hem Windows hem Linux'de sonuç bu .

Saygılar
Erdemdem.

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]

Alıntı (kerdemdemir):

Yanlış düşünüyorsun çünkü orası tam da sorunları olan kişilerin kullandıkları bir ortam.

Alıntı:

Bir yerde çok basit bir yanlış anlama olmalı.

Ali

--
[ Bu gönderi, http://ddili.org/forum'dan dönüştürülmüştür. ]