Einführung in Paralleles Programmieren 2.Teil

Zeitmessungen

Ich werde diese Unterprogramme dann in den folgenden Beispielen verwenden.

Zweidimensionales Beispiel

Zuerst die CPU-Variante (mit Threads):

void grauwerte(int& R[N][M],int& G[N][M],int& B[N][M],float& H[N][M],
               int i0,int nt)
{
 for(int i=i0;i<N;i+=nt)
  for(int j=0;j<M;j++)
   {
    H[i][j] = (0.299*R[i][j] + 0.587*G[i][j] + 0.114*B[i][j])/255;
   }
}

Jetzt gibts aber schon ein kleines Problem: in c sind mehrdimensionale Felder nur schlecht benutzbar. Insbesondere müssen N und M unbedingt Konstanten sein. Für unterschiedliche Bildgrössen also nicht verwendbar.
Wir müssen die zweidimensionalen Felder also wieder durch gewöhnliche eindimensionale ersetzen.
Also statt feld[N][M] in der Deklaration schreiben wir feld[N*M], oder als Parameter int*feld (oder float*feld).
Dann um auf ein Element zuzugreifen statt feld[i][j] schreiben wir feld[j*N+i].
(Einige Leute machen das anders herum (z.B. CUBLAS) und verwenden feld[i*M+j])
Man könnte sich auch ein Makro schreiben um dann feld[IDX(i,j)] zu benutzen.

Hier also das angepasste Beispiel ohne Makros:

void grauwerte(const int*R,const int*G,const int*B,float*H,
               int i0,int nt)
{
 for(int i=i0;i<N;i+=nt)
  for(int j=0;j<M;j++)
   {
    H[j*N+i] = (0.299*R[j*N+i] + 0.587*G[j*N+i] + 0.114*B[j*N+i])/255;
   }
}

Und jetzt der OpenCL-Kernel:

__kernel
void calc(global int*R,global int*G,global int*B,global float*H,int N,int M)
{
 int i=get_global_id(0);
 int j=get_global_id(1);
 H[j*N+i] = (0.299*R[j*N+i] + 0.587*G[j*N+i] + 0.114*B[j*N+i])/255;
}

  kernel_calc.setArg(0,d_R);
  kernel_calc.setArg(1,d_G);
  kernel_calc.setArg(2,d_B);
  kernel_calc.setArg(3,d_H);
  kernel_calc.setArg(4,N);
  kernel_calc.setArg(5,M);
  queue.enqueueNDRangeKernel(kernel_quadr, cl::NullRange,
                             cl::NDRange(N,M), cl::NDRange(32,32));
  queue.finish();

#define rup(a,b) ((a+b-1)/(b)*(b)) //aufrunden, damit durch b teilbar

  queue.enqueueNDRangeKernel(kernel_quadr, cl::NullRange,
  cl::NDRange(rup(N,32),rup(M,32)), cl::NDRange(32,32));

__kernel
void calc(global const int*R,global const int*G,global const int*B,
          global float*H,const int N,const int M)
{
 int i=get_global_id(0);
 int j=get_global_id(1);
 if(i<N && j<M)
   H[j*N+i] = (0.299*R[j*N+i] + 0.587*G[j*N+i] + 0.114*B[j*N+i])/255;
}

Um die Zeiten zu messen habe ich die oben erwähnten Unterprogramme so eingebaut:

 stoppuhr_reset();
 calc_opencl(Rot,Gruen,Blau,Hwerte,N,M);
 int usec=stoppuhr_read();
 copy_from_device(Hwerte,N*M);
 int usec2=stoppuhr_read();
 printf("Rechenzeit OpenCL: ");
 zeit_print(usec,FLOP);
 printf("   zurueckkopiert: ");
 zeit_print(usec2);

Optimierungen

 for(int i=0,k=j*N;i<N;i++,k++)
   H[k] = 0.299/255*R[k] + 0.587/255*G[k] + 0.114/255*B[k];

Download und Testlauf

Ausgabe vom Testlauf auf meinem Laptop:

2073600 Zufallszahlen erstellt.
 103 113 6 105 ... 41 6 251 10
Seriell auf CPU gerechnet:
Rechenzeit CPU: 37.965 ms  (0.273 GFLOPS)
Mit 4 CPU-Threads gerechnet:
Rechenzeit CPU (optimiert): 3.525 ms  (2.9 GFLOPS)
Auf Grafikkarte mit OpenCL gerechnet:
Using platform: NVIDIA CUDA
Using device: GeForce GTX 850M
Rechenzeit OpenCL: 1.277 ms  (8.1 GFLOPS)
   zurueckkopiert: 6.585 ms
Auf Grafikkarte mit CUDA gerechnet:
 Rechenzeit  CUDA: 2.357 ms  (4.4 GFLOPS)
   zurueckkopiert: 7.624 ms
 0.623502 0.516267 0.164071 0.300651 ... 0.292592 0.101984 0.748243 0.066392
Ergebnisse erfolgreich ueberprueft.

Und Testlauf auf meinem Gaming-PC:

2073600 Zufallszahlen erstellt.
 103 113 6 105 ... 41 6 251 10
auf CPU (12 Threads) gerechnet:
Rechenzeit CPU: 3.966 ms  (2.6 GFLOPS)
auf CPU (12 Threads) leicht optimiert:
Rechenzeit CPU (optimiert): 1.679 ms  (6.2 GFLOPS)
Auf Grafikkarte mit OpenCL gerechnet:
Using platform: NVIDIA CUDA
Using device: GeForce RTX 2080
Rechenzeit OpenCL: 0.183 ms  (56.7 GFLOPS)
   zurueckkopiert: 1.626 ms
Auf Grafikkarte mit CUDA gerechnet:
 Rechenzeit  CUDA: 0.233 ms  (44.5 GFLOPS)
   zurueckkopiert: 1.236 ms
 0.623502 0.516267 0.164071 0.300651 ... 0.292592 0.101984 0.748243 0.066392
Ergebnisse erfolgreich ueberprueft.

weitere Optimierungen

 H[k] = (float)(0.299/255)*R[k] + (float)(0.587/255)*G[k] + (float)(0.114/255)*B[k];

Im nächsten Teil werde ich dann ein Beispiel verwenden, bei dem mehr gerechnet werden muss, so dass Optimierungen sinnvoll werden.

Fortsetzungen:

Vergleich von OpenCL und CUDA:

Teilweise hier gefunden: Porting CUDA to OpenCL