OpenCL kullanarak mizaçsal dizi toplamı OpenCL

Question

OpenCL kullanarak n boyutlu noktalar arasındaki Öklid mesafesini hesaplıyorum. İki n-boyutlu nokta listesi alıyorum ve ilk tablodaki her noktadan ikinci tablonun her noktasına kadar olan mesafeleri içeren bir dizi döndürmeliyim.

Benim yaklaşımım {Table2 {her noktasında ...} için} Tablo-1'de her noktası için (düzenli doble döngü yapmak ve sonra paralell noktaların her çifti için hesaplama yapmaktır.

Öklid mesafe daha sonra 3 bölüme ayrılır: 1. noktalarındaki her bir boyut arasındaki farkı alın. (ki her boyut için hala) arasındaki fark. 2. 4. elde edilen tüm değerleri toparlar. 3'te elde edilen değerin (bu adım bu örnekte atlanmıştır.)

Her şey bir çekicilik gibi çalışır Tüm farklılıkların toplamını (yani, yukarıda açıklanan prosedürün 3. adımını gerçekleştirerek, aşağıdaki kodun 49. satırını) toplamaya çalışın.

Test verileri olarak DescriptorLists kullanıyorum her biri 2 nokta ile: DescriptorList1: 001,002,003, ..., 127,128; (p1) 129,130,131, ..., 255,256; (p2)

TanımlayıcıList2: 000,001,002, ..., 126,127; (p1) 128,129,130, ..., 254,255; (p2)

Sonuçta elde edilen vektör şu değerlere sahip olmalıdır: 128, 2064512, 2130048, 128 Şu anda her koşula göre değişen rasgele sayılar alıyorum.

Neyi yanlış yaptığım konusunda herhangi bir yardım ya da ipucuna minnettarım. Umarım her şey çalışıyorum senaryo hakkında açıktır

#define BLOCK_SIZE 128 

typedef struct 
{ 
    //How large each point is 
    int length; 
    //How many points in every list 
    int num_elements; 
    //Pointer to the elements of the descriptor (stored as a raw array) 
    __global float *elements; 
} DescriptorList; 

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float As[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 

    int featA = get_local_id(0); 

    //temporary array to store the difference between each dimension of 2 points 
    float dif_acum[BLOCK_SIZE]; 

    //counter to track the iterations of the inner loop 
    int loop = 0; 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the i-th descriptor. Returns a DescriptorList with just the i-th 
     //descriptor in DescriptorList A 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory. 
     //returns one element of the only descriptor in DescriptorList tmpA 
     //and index featA 
     As[featA] = GetElement(tmpA, 0, featA); 
     //wait for all the threads to finish copying before continuing 
     barrier(CLK_LOCAL_MEM_FENCE); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current points 
      dif_acum[featA] = As[featA]-B.elements[k*BLOCK_SIZE + featA]; 
      //wait again 
      barrier(CLK_LOCAL_MEM_FENCE); 
      //square value of the difference in dif_acum and store in C 
      //which is where the results should be stored at the end. 
      C[loop] = 0; 
      C[loop] += dif_acum[featA]*dif_acum[featA]; 
      loop += 1; 
      barrier(CLK_LOCAL_MEM_FENCE); 
     } 
    } 
} 

Answer 1

Senin sorunun bu kod hatlarında yatıyor. Çalışma grubunuzdaki

C[loop] = 0; 
C[loop] += dif_acum[featA]*dif_acum[featA]; 

tüm ipler (Evet, aslında tüm ipler, ancak daha sonra gelelim) herhangi bir senkronizasyon olmadan aynı anda bu dizi pozisyonunu değiştirmeye çalışıyoruz. Birçok faktör bu gerçekten sorunlu hale:

1. çalışma grubu diğer konu zaten sonraki çizgisini infaz ettikten sonra bazı ipler C [döngü] için = 0 çağrılabilir yani paralel tamamen çalışacağı garanti edilmez
2. Paralel olarak yürütenler, aynı değeri C [döngü] 'den okurlar, artıları ile değiştirirler ve aynı adrese tekrar yazmaya çalışırlar. Bu yazarın sonucunun ne olduğundan tam olarak emin değilim (sanırım bir tanesi geri yazarken başarılı oluyor, diğerleri başarısız, ama tam olarak emin değilim), ama her iki şekilde de yanlış. Of

   local float* accum; 
   ... 
   accum[featA] = dif_acum[featA]*dif_acum[featA]; 
   barrier(CLK_LOCAL_MEM_FENCE); 
   for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
   { 
       if ((featA % (2*i)) == 0) 
        accum[featA] += accum[featA + i]; 
       barrier(CLK_LOCAL_MEM_FENCE); 
   } 
   if(featA == 0) 
       C[loop] = accum[0]; 
   

   : bu nedenle yerel bellekte birikmesine olanak sağlar, bu, hızlı olmayacak atomics kullanarak genel bellek üzerinde çalışmaya almak mümkün olabilir iken:

Şimdi Bunu düzeltmek sağlar Tabii ki bunun için diğer yerel tamponları yeniden kullanabilirsiniz, ama bence bu nokta açıktır (btw: Bu dif_acum'un yerel bellekte oluşturulacağından emin misiniz, çünkü bunun yerel bellekte bulunmayacağını bir yerde okuduğumu düşünüyorum. yerel bellek türünü önyüklemeli hale getirir.

bu kod hakkında bazı diğer noktalar:

1. Kodunuz için, (üzerinde çalışmak hangi öğeleri görmek için groupid ne de küresel kimliğini ya kullanmıyorsunuz) sadece çalışma grubu üzerinde kullanmaya dişli görünüyor Optimal performans daha sonra kullanmak isteyebilirsiniz.
2. kişisel tercih olabilirdi ama benim için
3. (sizin OpenCL kodunu değiştirildi gerekirdi sizi farkında olmadan ev sahibi kodunda değişebilir beri) bir tanımla kullanmak daha workgroupsize için get_local_size(0) kullanmak iyi görünüyor Kodunuzdaki engeller gereksizdir, çünkü hiçbir iş parçacığı, başka bir iş parçacığı tarafından yazılan yerel bellekte bir öğeye erişemez. Bu nedenle bunun için yerel belleği kullanmanız gerekmez. yapabildin son mermiyi düşünüldüğünde

basitçe yapın:

float As = GetElement(tmpA, 0, featA); 
... 
float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

Bu (ilk iki mermi düşünmediğini) kodu yapacak: Grizzly

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE]) 
{ 
    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 
    int loop = 0; 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 
     DescriptorList tmpA = GetDescriptor(A, i); 
     float As = GetElement(tmpA, 0, featA); 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
      for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
      { 
       if ((featA % (2*i)) == 0) 
       accum[featA] += accum[featA + i]; 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 
      if(featA == 0) 
       C[loop] = accum[0]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      loop += 1; 
     } 
    } 
} 

Answer 2

teşekkürler, şimdi var çalışan bir çekirdek. Grizzly:

'un cevabına göre değişiklik yapmak için ihtiyaç duyduğum bazı şeyler Kullanmakta olduğum dizilerde geçerli bir konuma başvurmayacak tüm konuları atmak için rutin başında bir IF ifadesi ekledim.

if(featA > BLOCK_SIZE){return;} 

(Bs i.g.) yerel (paylaşılan) belleğe ilk tanımlayıcısı kopyalarken, indeks fonksiyonu GetElement arama başına sadece bir elemanını döndürür beri belirtilmesi gerekir (Sorumun o atlanır). tampon her yineleme sonrası üzerine olduğu için bu

Bs[featA] = GetElement(tmpA, 0, featA); 

Daha sonra, TARAMA döngü küçük verdiği gerekli ve bir veri birinci iplik erişimi kontrol edemez. Bu yüzden kısmi sonuçları saklamak için dif_acum arabelleğini 'geri dönüştürüyorum' ve bu şekilde bu döngü boyunca tutarsızlıkları önlerim.

dif_acum[featA] = accum[featA]; 

Bir araya getirilecek terimleri güvenilir şekilde belirlemek için SCAN döngüsünde bazı sınır denetimi de vardır.

if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 

Son, ben deyim böylece yalnızca bir iplik onu değiştirir EĞER son dahilinde döngü değişkeni artışı dahil etmek mantıklı düşündüm.

İşte bu kadar. Hala BLOCK_SIZE tanımını ortadan kaldırmak için group_size'yi nasıl kullanabilirim ve eğer thread kullanımı ile ilgili daha iyi politikalar varsa bunu kullanabilirim.

Yani kod şöyle nihayet görünür:

Ben Grizzly'nın cevap çok müteşekkir olduğumu, her şeyden önce, söylemek zorunda

__kernel void CompareDescriptors(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE], __local float Bs[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 

    //global counter to store final differences 
    int loop = 0; 

    //auxiliary buffer to store temporary data 
    local float dif_acum[BLOCK_SIZE]; 

    //discard the threads that are not going to be used. 
    if(featA > BLOCK_SIZE){ 
     return; 
    } 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the gpidA-th descriptor 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory 
     Bs[featA] = GetElement(tmpA, 0, featA); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current descriptors 
      dif_acum[featA] = Bs[featA]-B.elements[k*BLOCK_SIZE + featA]; 

      //square the values in dif_acum 
      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      //copy the values of accum to keep consistency once the scan procedure starts. Mostly important for the first element. Two buffers are necesarry because the scan procedure would override values that are then further read if one buffer is being used instead. 
      dif_acum[featA] = accum[featA]; 

      //Compute the accumulated sum (a.k.a. scan) 
      for(int j = 1; j < BLOCK_SIZE; j *= 2){ 
       int next_addend = featA-(j/2); 
       if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 
        dif_acum[featA] = accum[featA] + accum[next_addend]; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 

       //copy As to accum 
       accum[featA] = GetElementArray(dif_acum, BLOCK_SIZE, featA); 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 

      //tell one of the threads to write the result of the scan in the array containing the results. 
      if(featA == 0){ 
       C[loop] = accum[BLOCK_SIZE-1]; 
       loop += 1; 
      } 
      barrier(CLK_LOCAL_MEM_FENCE); 

     } 
    } 
}