Forum: D Programlama Dili RSS
Cuda gunlukleri
kerdemdemir #1
Üye Eyl 2013 tarihinden beri · 121 mesaj · Konum: Danimarka
Grup üyelikleri: Üyeler
Profili göster · Bu konuya bağlantı
Konu adı: Cuda gunlukleri
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/…

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…  

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
kerdemdemir #2
Üye Eyl 2013 tarihinden beri · 121 mesaj · Konum: Danimarka
Grup üyelikleri: Üyeler
Profili göster · Bu konuya bağlantı
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;
}
acehreli (Moderatör) #3
Kullanıcı başlığı: Ali Çehreli
Üye Haz 2009 tarihinden beri · 4511 mesaj
Grup üyelikleri: Genel Moderatörler, Üyeler
Profili göster · Bu konuya bağlantı
Hız farkı büyük! :)

Ali
kerdemdemir #4
Üye Eyl 2013 tarihinden beri · 121 mesaj · Konum: Danimarka
Grup üyelikleri: Üyeler
Profili göster · Bu konuya bağlantı
D tarafında CPU ile beamform işine başladım.

https://github.com/kerdemdemir/DSharpEar/blob/master/DShar…

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
acehreli (Moderatör) #5
Kullanıcı başlığı: Ali Çehreli
Üye Haz 2009 tarihinden beri · 4511 mesaj
Grup üyelikleri: Genel Moderatörler, Üyeler
Profili göster · Bu konuya bağlantı
kerdemdemir:
yazı hazırlamıştım fakat browser penceresini yanlışlıkla kapatınca hepsi uçtu

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

Ali
Doğrulama Kodu: VeriCode Lütfen resimde gördüğünüz doğrulama kodunu girin:
İfadeler: :-) ;-) :-D :-p :blush: :cool: :rolleyes: :huh: :-/ <_< :-( :'( :#: :scared: 8-( :nuts: :-O
Özel Karakterler:
Bağlı değilsiniz. · Şifremi unuttum · ÜYELİK
This board is powered by the Unclassified NewsBoard software, 20100516-dev, © 2003-10 by Yves Goergen
Şu an: 2017-10-18, 20:45:31 (UTC -07:00)