OpenCL finden max im Array

Question

Ich versuche, Max-Wert in einem 1-D-Array durch Reduktionsoperatoren zu finden. Ich hatte das Verfahren in refered: OpenCL™ Optimization Case Study: Simple Reductions

Nach meinem Code ist:

__kernel void Normallize(__global float* input, __global float* output,__global float* cmax, int rows, int cols){ 

    int g_idx = get_global_id(0); 

    for(int i=0 ; i< get_global_size(0) ; i++) cmax[i] = 0; 

    barrier(CLK_GLOBAL_MEM_FENCE); 

    for(int offset = get_global_size(0)/2 ; offset >0 ; offset--){ 
     if(g_idx < offset){ 
      float pre = input[g_idx]; 
      float next = input[g_idx + offset]; 
      cmax[g_idx] = (pre > next) ? pre:next; 
     } 
     barrier(CLK_GLOBAL_MEM_FENCE); 
    } 

    output[g_idx] = cmax[0]; 
} 

Nach etwas Forschung zu tun, ich kann immer noch nicht das Problem in meinem Code herauszufinden.

Answer 1

Meinen Sie diese (60% VALU Nutzung für einen AMD GPU) ?:

__kernel void maxping(__global __read_only float * a, __global __write_only float *b){ 
         int threadId=get_global_id(0); 
         int localThreadId=get_local_id(0); 
         int localSize=get_local_size(0); 
         __local float fastMem[256]; 
         fastMem[localThreadId]=a[threadId]; 
         barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE); 

         for(int i=localSize/2;i>=1;i/=2) 
         { 
          if(localThreadId<i) 
          { 
           if(fastMem[localThreadId]<fastMem[localThreadId+i]) 
            fastMem[localThreadId]=fastMem[localThreadId+i]; 
          } 
          barrier(CLK_LOCAL_MEM_FENCE); 
         } 
         if(localThreadId==0) 
          b[threadId]=fastMem[localThreadId]; 
} 

wobei jede Gruppe (256 Fäden) in einem lokalen Speicher verringern und Einstellen jeden First-In-Gruppenwertes max seiner Gruppe. Dieses Beispiel hat 4096 Elemente von 0 bis 4095.

Für oberen Kern ist VALU Verwendung so etwas wie:

x: idle thread 
o: thread in process, m: thread in memory operation 

** :  m m m m m m m m m m m m m m m m 
i=0 :  o o o o o o o o x x x x x x x x 
i=1 :  o o o o x x x x x x x x x x x x 
i=2 :  o o x x x x x x x x x x x x x x 
i=3 :  o x x x x x x x x x x x x x x x 
** :  m m m m m m m m m m m m m m m m 

aber ich zählt Schritte und jede Reihe erstreckt sich über 250 Einheiten.

__kernel void maxpong(__global __write_only float * a, __global __read_only float *b){ 
         int threadId=get_global_id(0); 
         int localSize=get_local_size(0); 
         int maxGroups=4096/localSize; 
         if(threadId==0) 
         { 
          float maxv=FLT_MIN; 
          for(int i=0;i<maxGroups;i++) 
          { 
           if(maxv<b[i*localSize]) 
            maxv=b[i*localSize]; 
          } 
          a[0]=maxv; 

         } 
} 

wo nur erster Thread (am besten in CPU) nicht ein einfaches max (0,1,2, ..., M) und stellt erstes Element einer zu max (a).

Der erste Kernel macht 255/256 des gesamten Computing. Aber es lässt die Hälfte der Kerne jeder Recheneinheit unberührt. Sie können also eine andere Sache in dieser anderen Hälfte der Kerne sortieren. Das könnte ein anderes Array sein, das max() 'ed ist oder das gleiche Array' s min() 'oder das gleiche Maximum des gleichen Arrays, aber mit der Hälfte davon arbeitet, während andere Kerne auf der anderen Hälfte arbeiten.

% 73 VALU Nutzung für max (a) mit einem anderen Anfang kernel:

  __kernel void maxping(__global __read_only float * a, __global __write_only float *b){ 
       int threadId=get_global_id(0); 
       int localThreadId=get_local_id(0); 
       int localSize=get_local_size(0); 
       __local float fastMem[256]; 
       __local float fastMem2[256]; 
       fastMem[localThreadId]=a[threadId]; 
       fastMem2[localThreadId]=a[threadId+2048]; 

       barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE); 

       for(int i=localSize/2;i>=1;i/=2) 
       { 
        if(localThreadId<i) 
        { 
         // sorting first part 
         if(fastMem[localThreadId]<fastMem[localThreadId+i]) 
          fastMem[localThreadId]=fastMem[localThreadId+i]; 
        } 
        else if(localThreadId>localSize-i) 
        { 
         // sorting second part 
         if(fastMem2[localThreadId]<fastMem2[localThreadId-i]) 
          fastMem2[localThreadId]=fastMem2[localThreadId-i]; 
        } 
        else 
        { 
         // idle thread. Free compute slot. 
         // can squeeze some geometry computing 
         // or up-sweep scan of another reduction type 
        } 
        barrier(CLK_LOCAL_MEM_FENCE); 
       } 
       if(localThreadId==0) 
        b[threadId]=(fastMem[localThreadId]>fastMem2[255]?fastMem[localThreadId]:fastMem2[255]); 
      } 

Dies nutzt 2048 Threads für 4096-Element-Array. Stellt 0., 256., 512., .. Elemente auf ihre jeweiligen Gruppenmaximumwerte ein, dann können Sie einfach überprüfen, welches auf der Hostseite größer ist.

Es gibt immer noch unbenutzte Kerne.

Für oberen Kern ist VALU Verwendung so etwas wie:

x: idle thread 
o: thread in process, m: thread doing memory operation 

** :  m m m m m m m m m m m m m m m m 
i=0 :  o o o o o o o o o o o o o o o o 
i=1 :  o o o o x x x x x x x x o o o o 
i=2 :  o o x x x x x x x x x x x x o o 
i=3 :  o x x x x x x x x x x x x x x o 
** :  m m m m m m m m m m m m m m m m 

aber ich Schritte für log2 (256) mal so gibt es mehr „i“ Schritte und AMD-Hardware verfügt über 64 Kerne, die selbst voll dienen, wenn es sind 64 Threads in einem Schritt. Wenn wir die gesamte Thread-Verwendung für diese Schleife summieren, gibt es nicht% 73, aber wenn andere "Warps" (40 davon) zu derselben Recheneinheit strömen, werden mehr Löcher gefüllt, daher werden mehr Vektor-Arithmetik-Logik-Einheiten häufiger verwendet. Selbst der lokale Speicherzuweisungsteil ist wichtig, da die Speicheroperationseinheiten aller Kerne beschäftigt sind (global zu lokal, lokal zu global), während andere Warps Vergleichseinheiten beschäftigt halten.

Bearbeiten: Wenn Sie nicht ein Vielfaches von 256 der globalen Größe benötigen, können Sie eine globale ID-Prüfung nach dem lokalen Speichervorgang hinzufügen, so dass es nicht undefiniert Verhalten. Vielleicht können Sie Array stattdessen mit zusätzlichen FLT_MIN-Werten auffüllen.