Gleichzeitiges Programmieren mit CUDA. Teil 2: GPU-Hardware und parallele Kommunikationsmuster

  • Tutorial

Inhalt


Teil 1: Einführung.
Teil 2: GPU-Hardware und parallele Kommunikationsmuster.
Teil 3: Grundlegende GPU-Algorithmen: Verkleinern, Scannen und Histogramm.
Teil 4: Grundlegende GPU-Algorithmen: Kompaktes, segmentiertes Scannen, Sortieren. Praktische Anwendung einiger Algorithmen.
Teil 5: Optimierung von GPU-Programmen.
Teil 6: Beispiele für die Parallelisierung von sequentiellen Algorithmen.
Teil 7: Zusätzliche Themen der parallelen Programmierung, dynamische Parallelität.

Vorlagen für die parallele Kommunikation



Was ist Parallel Computing? Nichts als eine Vielzahl von Threads, die ein bestimmtes Problem gemeinsam lösen . Das Schlüsselwort lautet hier „kooperativ“ - um eine Kooperation zu erreichen, müssen bestimmte Kommunikationsmechanismen zwischen den Flüssen angewendet werden. Bei Verwendung von CUDA erfolgt die Kommunikation über den Speicher: Streams können Eingabedaten lesen, Ausgabedaten ändern oder Zwischenergebnisse austauschen.
Abhängig davon, wie die Flüsse durch den Speicher kommunizieren, werden verschiedene Muster der parallelen Kommunikation unterschieden .
Im vorherigen Teil wurde die Konvertierung eines Farbbildes in Graustufen als einfaches Beispiel für die Verwendung von CUDA betrachtet. Hierzu wurde die Intensität jedes Pixels des Ausgabebildes in Graustufen nach der Formel I = A * pix.R + B * pix.G + C * pix.B berechnet , wobei A, B, C Konstanten sind, pix das entsprechende Pixel des Originalbildes ist. Grafisch ist dieser Prozess wie folgt:

Wenn wir die Methode zur Berechnung des Ausgabewerts ignorieren, erhalten wir die erste parallele Kommunikationsvorlage - Zuordnung : Die gleiche Funktion wird für jedes Eingabedatenelement mit dem Index i ausgeführt und das Ergebnis im Ausgabedatenfeld unter gespeichert gleicher Index i. Шаблон map очень эффективен на GPU, и к тому же просто выражается в рамках CUDA — достаточно запустить по одному потоку на каждый входной элемент (как и было сделано в предыдущей части для задачи конвертации изображения). Однако, лишь малую часть задач можно решить используя только этот шаблон.
Если же для вычисления выходного значения с индексом i используются несколько входных элементов, то такой шаблон называют gather, и выглядеть он может так:

Или так:

Эффективность реализации данного шаблона на CUDA зависит от того, какие именно входные значения используются при расчете выходного и их количества — лучше всего когда используется небольшое количество идущих подряд элементов.
Обратный шаблон, ein Scatter - jedes Eingangselement wirkt sich ein paar (oder einer) der Ausgangselemente, grafisch sowie sieht die sammeln , sondern die Änderung der Bedeutung: wir jetzt haben „abgestoßen“ nicht auf den Ausgangselemente, die ihren Wert berechnet und die Eingabe , die Einfluss Werte bestimmter Ausgabeelemente. Ziemlich oft ein und dasselbe Problem kann sowohl im Rahmen des Musters gelöst werden das sammeln , und eine Streuung . Wenn wir beispielsweise drei benachbarte Eingabeelemente mitteln und in das Ausgabearray schreiben möchten, können wir:
  • Führen Sie für jedes Ausgabeelement einen Stream aus, wobei jeder Stream die Werte von 3 benachbarten Eingabeelementen mittelt .
  • Oder führen Sie einen Stream für jedes Eingabeelement aus, wobei jeder Stream 1/3 des Werts seines Eingabeelements zum Wert des entsprechenden Ausgabeelements addiert - Scatter .

Как вы скорее всего догадались, при использовании подхода scatter встает проблема синхронизации, так как несколько потоков могут пытаться модифицировать один и тот же выходной элемент одновременно.
Также стоит выделить подвид шаблона gatherstencil: в данном шаблоне накладывается ограничение на входные элементы, которые участвуют в вычислении выходного элемента — а именно, это могут быть только соседние элементы. В случае с 2D/3D изображениями, могут использоваться различные виды данного шаблона, например двумерный трафарет фон Неймана:

или двумерный трафарет Мура:

В связи с этим ограничением, шаблон stencilEs wird normalerweise recht effizient im Rahmen von CUDA implementiert: Es reicht aus, einen Stream pro Ausgabeelement auszuführen, während der Stream die benötigten Eingabeelemente liest. Bei dieser Organisation der Berechnungen wird die Effizienz durch zwei Faktoren sichergestellt:
  1. Alle Daten, die für einen Stream benötigt werden, sind im Speicher gruppiert (im Fall eines eindimensionalen Arrays handelt es sich um ein festes "Stück" Speicher, im 2D-Fall um mehrere Speicherstücke, die sich im gleichen Abstand voneinander befinden).
  2. Der Wert einiger Eingabeelemente wird mehrmals aus benachbarten Streams gelesen (die genaue Anzahl der Lesungen hängt von der ausgewählten Maske ab) - es wird möglich, die von CUDA bereitgestellten Daten "wiederzuverwenden" - dies wird später in diesem Artikel erläutert.

GPU-Hardware


Betrachten Sie die allgemeine Struktur der GPU-Hardware:

  • Eine CUDA-kompatible GPU besteht aus mehreren (normalerweise Dutzenden) Streaming-Multiprozessoren (Streaming-Multiprozessoren), im Folgenden als SM bezeichnet .
  • Jede SM besteht wiederum aus mehreren Dutzend Simple / Streaming-Prozessoren (SP) (Regular / Stream-Prozessoren) oder genauer gesagt CUDA-Kernen ( CUDA-Kernen ). Diese Jungs sind eher wie die übliche CPU - sie haben ihre Register, Cache, etc. Jeder SM hat auch einen eigenen gemeinsamen Speicher (Shared Memory) - eine Art zusätzlicher Cache, auf den alle SPs zugreifen können. Er kann sowohl als Cache für häufig verwendete Daten als auch für die „Kommunikation“ zwischen Threads eines CUDA-Blocks verwendet werden.
  • Die GPU verfügt auch über einen eigenen Speicher, den so genannten Gerätespeicher , der allen CUDA-Streams gemeinsam ist. Mit ihm arbeiten die Funktionen cudaMalloc und cudaMemcpy (auch Schulkinder und GPU-Hersteller messen ihn gern nach Größe) .

Einhaltung des CUDA-Modells und der GPU-Hardware. CUDA-Garantien


Gemäß dem CUDA-Modell teilt ein Programmierer eine Aufgabe in Blöcke und Blöcke in Threads auf. Wie werden diese Softwareeinheiten den oben beschriebenen GPU-Hardwareblöcken zugeordnet?

  • Jeder Block wird auf der zugewiesenen SM vollständig ausgeführt .
  • Die Verteilung von Blöcken durch SM ist eine GPU, kein Programmierer.
  • Alle Abläufe von Block X werden in Warps genannte Gruppen (normalerweise Warps) unterteilt und auf SM ausgeführt . Die Größe dieser Gruppen hängt vom GPU-Modell ab. Bei Modellen mit Fermi- Mikroarchitektur beträgt sie beispielsweise 32. Alle Threads von einem Warp werden gleichzeitig ausgeführt und belegen einen bestimmten Teil der SM- Ressourcen . Darüber hinaus führen sie entweder den gleichen Befehl aus (jedoch mit unterschiedlichen Daten) oder sind inaktiv.

Auf dieser Grundlage bietet CUDA die folgenden Garantien:
  • Alle Threads in einem bestimmten Block werden auf einer SM ausgeführt .
  • Alle Threads eines bestimmten Kernels werden ausgeführt, bevor der nächste Kernel startet.

CUDA garantiert nicht, dass:
  • Einiger Block X wird vor ausgeführt werden / nach / gleichzeitig mit einem gewissen Block Y .
  • Einige Block X wird an einigen spezifischen gebildet werden SM Z .

Synchronisieren


Daher listen wir die wichtigsten von CUDA bereitgestellten Synchronisationsmechanismen auf:
  • Eine Barriere ist ein Punkt im Kernel-Code, bei dessen Erreichen ein Thread nur dann weiter "gehen" kann, wenn alle Threads aus seinem Block diesen Punkt erreicht haben. Nochmals: Mit der Barriere können Sie nur die Flüsse eines Blocks synchronisieren , und im Prinzip nicht alle! Die Einschränkung liegt auf der Hand , da die Anzahl der vom Programmierer festgelegten Blöcke die Anzahl der verfügbaren SMs erheblich überschreiten kann .
  • Atomare Operationen - Ähnlich wie bei atomaren Operationen der CPU finden Sie hier eine vollständige Liste der verfügbaren Operationen .
  • __threadfence ist kein primitives Element der Synchronisation: Nach Erreichen dieser Anweisung kann ein Thread erst dann weiter ausgeführt werden, wenn alle seine Speichermanipulationen für andere Threads sichtbar sind. Tatsächlich wird der Thread gezwungen, den Cache zu leeren.

Grundprinzipien für eine effektive Nutzung von CUDA


  • Das Prinzip der Erhöhung des Verhältnisses (Zeit der nützlichen Arbeit) / (Zeit der Operationen mit Gedächtnis) wurde in einem früheren Artikel erörtert. Der Wert eines Bruchs kann auf zwei Arten erhöht werden: Erhöhen Sie den Zähler, verringern Sie den Nenner, dh, Sie müssen entweder mehr Arbeit leisten oder weniger Zeit für Speicheroperationen aufwenden. Neben der offensichtlichen Lösung - um die Anzahl der Speicherzugriffe so weit wie möglich zu reduzieren, werden die folgenden Prinzipien für eine effektive Arbeit mit dem Speicher verwendet:
    • Bewegen häufig zugegriffenen Daten in dem schnelleren Speicher: Thread - Lokalspeicher > Gesamtspeicherblock >> gemeinsam genutzten Speichervorrichtung >> Speicher des Hosts. Wenn also mehrere Threads dieselben Daten im selben Block verwenden, ist es höchstwahrscheinlich sinnvoll, sie in den gemeinsamen Speicher des Blocks zu verschieben.
    • Sequentieller Zugriff auf den Speicher: Da die Flüsse in Blöcken tatsächlich in Warp-Gruppen ausgeführt werden, sofern die Threads in einem Warp sequentiell mit den im Speicher befindlichen Daten arbeiten, kann CUDA in einem Befehl einen großen Teil des Speichers lesen. Andernfalls erhöht sich die Anzahl der Speicherzugriffe, wenn die Streams in den Warp-Zugriffsdaten im Speicher verstreut sind.

  • Reduzierung der Flussdivergenz: Ein Merkmal von CUDA ist, dass Flüsse in einem Warp immer entweder den gleichen Befehl ausführen oder im Leerlauf sind. Also, wenn es einen Code im Kernel-Code des Formulars gibt
    if (threadIdx.x % 2 == 0)
    {
      ...
    }
    ...
    

    Dann wartet die Hälfte der Warp-Threads (mit einem ungeraden Index), bis die zweite Hälfte den Code innerhalb des if-a ausführt. Daher sollten solche Situationen vermieden werden.

Wir schreiben das zweite Programm auf CUDA


Lass uns weiter üben. Als Beispiel für die vorgestellte Theorie schreiben wir ein Programm, das Gaußsche Bildunschärfe ausführt . Das Funktionsprinzip ist wie folgt: Der Wert der Kanäle R, G, B des Pixels im ausgegebenen unscharfen Bild wird als die gewichtete Summe der Werte der Kanäle R, G, B (jeweils) aller Pixel des Originalbilds in einem bestimmten Muster

berechnet : Die Werte der Gewichte werden unter Verwendung der Gaußschen 2D-Verteilung berechnet, aber wie genau ist dies? getan zu werden ist für unsere aufgabe nicht allzu wichtig.
Wie Sie der Beschreibung der Aufgabe entnehmen können, ist es für die Implementierung dieses Algorithmus ganz natürlich, die Schablonenvorlage zu wählen- Schließlich wird jedes Pixel des Ausgabebildes auf der Grundlage der entsprechenden benachbarten Pixel des Originalbildes berechnet.
Beginnen wir mit dem Grundgerüst des Programms:
main.cpp
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include "openMP.hpp"
#include "CUDA_wrappers.hpp"
#include "common/image_helpers.hpp"
void prepareFilter(float **filter, int *filterWidth, float *filterSigma)
{
  static const int blurFilterWidth = 9;
  static const float blurFilterSigma = 2.;
  *filter = new float[blurFilterWidth * blurFilterWidth];
  *filterWidth = blurFilterWidth;
  *filterSigma = blurFilterSigma;
  float filterSum = 0.f;
  const int halfWidth = blurFilterWidth/2;
  for (int r = -halfWidth; r <= halfWidth; ++r) 
  {
    for (int c = -halfWidth; c <= halfWidth; ++c) 
    {
      float filterValue = expf( -(float)(c * c + r * r) / (2.f * blurFilterSigma * blurFilterSigma));
      (*filter)[(r + halfWidth) * blurFilterWidth + c + halfWidth] = filterValue;
      filterSum += filterValue;
    }
  }
  float normalizationFactor = 1.f / filterSum;
  for (int r = -halfWidth; r <= halfWidth; ++r) 
  {
    for (int c = -halfWidth; c <= halfWidth; ++c) 
    {
      (*filter)[(r + halfWidth) * blurFilterWidth + c + halfWidth] *= normalizationFactor;
    }
  }  
}
void freeFilter(float *filter)
{
  delete[] filter;
}
int main( int argc, char** argv )
{
  using namespace cv;
  using namespace std;  
  using namespace std::chrono;
  if( argc != 2)
  {
    cout <<" Usage: blur_image imagefile" << endl;
    return -1;
  }
  Mat image, blurredImage, referenceBlurredImage;
  uchar4 *imageArray, *blurredImageArray;
  prepareImagePointers(argv[1], image, &imageArray, blurredImage, &blurredImageArray, CV_8UC4);
  int numRows = image.rows, numCols = image.cols;
  float *filter, filterSigma;
  int filterWidth;
  prepareFilter(&filter, &filterWidth, &filterSigma);
  cv::Size filterSize(filterWidth, filterWidth);
  auto start = system_clock::now();
  cv::GaussianBlur(image, referenceBlurredImage, filterSize, filterSigma, filterSigma, BORDER_REPLICATE);
  auto duration = duration_cast(system_clock::now() - start);
  cout<<"OpenCV time (ms):" << duration.count() << endl; 
  start = system_clock::now();
  BlurImageOpenMP(imageArray, blurredImageArray, numRows, numCols, filter, filterWidth);
  duration = duration_cast(system_clock::now() - start);
  cout<<"OpenMP time (ms):" << duration.count() << endl; 
  cout<<"OpenMP similarity:" << getEuclidianSimilarity(referenceBlurredImage, blurredImage) << endl;
  for (int i=0; i<4; ++i)
  {
    memset(blurredImageArray, 0, sizeof(uchar4)*numRows*numCols); 
    start = system_clock::now();
    BlurImageCUDA(imageArray, blurredImageArray, numRows, numCols, filter, filterWidth);
    duration = duration_cast(system_clock::now() - start);
    cout<<"CUDA time full (ms):" << duration.count() << endl;
    cout<<"CUDA similarity:" << getEuclidianSimilarity(referenceBlurredImage, blurredImage) << endl;
  }
  freeFilter(filter);
  return 0;
}


Nach Punkten:
  1. Wir lesen die Bilddatei, bereiten Zeiger auf das Originalbild und das daraus resultierende unscharfe Bild vor. Die prepareImagePointers- Funktion bleibt unverändert. Falls erforderlich, können Sie den Quellcode auf bitbucket anzeigen.
  2. Wir bereiten einen Gaußschen Filter vor - das ist ein Satz unserer Skalen. Wir erinnern uns auch an die verwendeten Filterparameter, damit sie später an OpenCV übertragen werden und ein Muster eines unscharfen Bildes erhalten, um die korrekte Funktionsweise unserer Algorithmen zu überprüfen.
  3. Wir rufen die Gaußsche Unschärfefunktion von OpenCV auf, speichern das resultierende Sample und messen die aufgewendete Zeit.
  4. Wir nennen die mit OpenMP geschriebene Gaußsche Unschärfefunktion, messen die aufgewendete Zeit und vergleichen das Ergebnis mit der Probe. Die Bildähnlichkeitsberechnungsfunktion getEuclidianSimilarity lautet wie folgt:
    getEuclidianSimilarity
    double getEuclidianSimilarity(const cv::Mat& a, const cv::Mat& b)
    {
      double errorL2 = cv::norm(a, b, cv::NORM_L2);
      double similarity = errorL2 / (double) (a.rows * a.cols);
      return similarity;
    }
    


    Tatsächlich wird die durchschnittliche Summe der Quadrate der Unterschiede in den Werten aller Kanäle aller Pixel von zwei Bildern ermittelt.
  5. Wir nennen die CUDA-Version der Gaußschen Unschärfe viermal, wobei jedes Mal die aufgewendete Zeit gemessen und das Ergebnis mit der Stichprobe verglichen wird. Warum 4 mal anrufen? Tatsache ist, dass während des allerersten Anrufs eine bestimmte Zeit für die Initialisierung aufgewendet wird. Daher ist es besser, mehrmals zu starten und die Zeit zu messen, die für nachfolgende Anrufe aufgewendet wurde.

OpenMP-Implementierung des Algorithmus:
openMP.hpp
#include 
#include 
#include 
#include 
void BlurImageOpenMP(const uchar4 * const imageArray, 
					 uchar4 * const blurredImageArray, 
					 const long numRows, 
					 const long numCols, 
					 const float * const filter, 
					 const size_t filterWidth)
{
	using namespace std;
	const long halfWidth = filterWidth/2;
    #pragma omp parallel for collapse(2)
    for (long row = 0; row < numRows; ++row)
    {
        for (long col = 0; col < numCols; ++col)
        {
    	  float resR=0.0f, resG=0.0f, resB=0.0f;
		  for (long filterRow = -halfWidth; filterRow <= halfWidth; ++filterRow) 
		  {
		    for (long filterCol = -halfWidth; filterCol <= halfWidth; ++filterCol) 
		    {
		      //Find the global image position for this filter position
		      //clamp to boundary of the image
		      const long imageRow = min(max(row + filterRow, static_cast(0)), numRows - 1);
		      const long imageCol = min(max(col + filterCol, static_cast(0)), numCols - 1);
		      const uchar4 imagePixel = imageArray[imageRow*numCols+imageCol];
		      const float filterValue = filter[(filterRow+halfWidth)*filterWidth+filterCol+halfWidth];
		      resR += imagePixel.x*filterValue;
		      resG += imagePixel.y*filterValue;
		      resB += imagePixel.z*filterValue;
		    }
		  }
		  blurredImageArray[row*numCols+col] = make_uchar4(resR, resG, resB, 255);
        }
    }
}


Betrachten wir für alle 3 Kanäle jedes Pixels des Quellbildes die beschriebene gewichtete Summe, so wird das Ergebnis an die entsprechende Position des Ausgabebildes geschrieben.
CUDA-Option:
CUDA.cu
#include 
#include 
#include "CUDA_wrappers.hpp"
#include "common/CUDA_common.hpp"
__global__
void gaussian_blur(const uchar4* const d_image,
                   uchar4* const d_blurredImage,
                   const int numRows,
                   const int numCols,
                   const float * const d_filter, 
                   const int filterWidth)
{
  const int row = blockIdx.y*blockDim.y+threadIdx.y;
  const int col = blockIdx.x*blockDim.x+threadIdx.x;
  if (col >= numCols || row >= numRows)
    return;
  const int halfWidth = filterWidth/2;
  extern __shared__ float shared_filter[]; 
  if (threadIdx.y < filterWidth && threadIdx.x < filterWidth)
  {
    const int filterOff = threadIdx.y*filterWidth+threadIdx.x;
    shared_filter[filterOff] = d_filter[filterOff];    
  }
  __syncthreads();
  float resR=0.0f, resG=0.0f, resB=0.0f;
  for (int filterRow = -halfWidth; filterRow <= halfWidth; ++filterRow) 
  {
    for (int filterCol = -halfWidth; filterCol <= halfWidth; ++filterCol) 
    {
      //Find the global image position for this filter position
      //clamp to boundary of the image
      const int imageRow = min(max(row + filterRow, 0), numRows - 1);
      const int imageCol = min(max(col + filterCol, 0), numCols - 1);
      const uchar4 imagePixel = d_image[imageRow*numCols+imageCol];
      const float filterValue = shared_filter[(filterRow+halfWidth)*filterWidth+filterCol+halfWidth];
      resR += imagePixel.x * filterValue;
      resG += imagePixel.y * filterValue;
      resB += imagePixel.z * filterValue;        
    }
  }
  d_blurredImage[row*numCols+col] = make_uchar4(resR, resG, resB, 255);     
}
void BlurImageCUDA(const uchar4 * const h_image, 
                   uchar4 * const h_blurredImage, 
                   const size_t numRows, 
                   const size_t numCols, 
                   const float * const h_filter, 
                   const size_t filterWidth)
{
  uchar4 *d_image, *d_blurredImage;
  cudaSetDevice(0);
  checkCudaErrors(cudaGetLastError());
  const size_t numPixels = numRows * numCols;
  const size_t imageSize = sizeof(uchar4) * numPixels;
  //allocate memory on the device for both input and output
  checkCudaErrors(cudaMalloc(&d_image, imageSize));
  checkCudaErrors(cudaMalloc(&d_blurredImage, imageSize));
  //copy input array to the GPU
  checkCudaErrors(cudaMemcpy(d_image, h_image, imageSize, cudaMemcpyHostToDevice));
  float *d_filter;
  const size_t filterSize = sizeof(float) * filterWidth * filterWidth;
  checkCudaErrors(cudaMalloc(&d_filter, filterSize));
  checkCudaErrors(cudaMemcpy(d_filter, h_filter, filterSize, cudaMemcpyHostToDevice));
  dim3 blockSize;
  dim3 gridSize;
  int threadNum;
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  threadNum = 32;
  blockSize = dim3(threadNum, threadNum, 1);
  gridSize = dim3(numCols/threadNum+1, numRows/threadNum+1, 1);
  cudaEventRecord(start);
  gaussian_blur<<>>(d_image, 
                                                     d_blurredImage, 
                                                     numRows, 
                                                     numCols, 
                                                     d_filter, 
                                                     filterWidth);
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
  float milliseconds = 0;
  cudaEventElapsedTime(&milliseconds, start, stop);
  std::cout << "CUDA time kernel (ms): " << milliseconds << std::endl;
  checkCudaErrors(cudaMemcpy(h_blurredImage, d_blurredImage, sizeof(uchar4) * numPixels, cudaMemcpyDeviceToHost));
  checkCudaErrors(cudaFree(d_filter));
  checkCudaErrors(cudaFree(d_image));
  checkCudaErrors(cudaFree(d_blurredImage));
}


  1. Ordnen Sie auf dem Gerät Speicher für das Quellbild, das Ausgabebild und den Filter zu. Kopieren Sie die entsprechenden Daten vom Host-Speicher in den zugewiesenen Gerätespeicher.
  2. Wir nennen den Kernel. Beachten Sie beim Aufrufen des Kernels den neuen 3. Parameter: <<filterSize >>>- legt die Größe des gemeinsam genutzten Speichers fest, den jeder Block benötigt. In dieser Aufgabe war es möglich, die Verwendung des gemeinsamen Speichers des Blocks auf zwei Arten zu implementieren: entweder die Filterdaten in den gemeinsamen Speicher des Blocks zu verschieben oder ein "Stück" des Bildes dorthin zu verschieben, das nur für diesen Block benötigt wird, da diese Daten von mehreren Blockflüssen gleichzeitig benötigt werden. Die zweite Option ist jedoch etwas komplizierter - Sie müssen berücksichtigen, dass der Teil des Eingabebildes, der für jeden Block benötigt wird, geringfügig größer als der Block selbst ist -, wir führen schließlich die Sammeloperation durch , was bedeutet, dass jeder Stream die Werte eines Pixels des Ausgabebildes unter Verwendung mehrerer benachbarter Pixel des Originalbildes berechnet:

    Daher Ich entschied mich für die erste Option, was bedeutet, dass jeder Block genau benötigtsizeof (float) * filterWidth * filterWidth Speicher zum Speichern aller Filterwerte. Das Verschieben der Filtergewichte aus dem Gerätespeicher in den gemeinsamen Speicher des Blocks geschieht folgendermaßen:
    Versteckter Text
     extern __shared__ float shared_filter[]; 
      if (threadIdx.y < filterWidth && threadIdx.x < filterWidth)
      {
        const int filterOff = threadIdx.y*filterWidth+threadIdx.x;
        shared_filter[filterOff] = d_filter[filterOff];    
      }
      __syncthreads();
    


    Hier __shared__ in der Deklaration des Arrays von Filtergewichtungen besagt, dass diese Daten im gemeinsamen Speicher des Blocks abgelegt werden sollen; extern bedeutet, dass die Größe des zugewiesenen Speichers am Kernel-Aufrufpunkt festgelegt wird; __syncthreads ist eine Barriere, die garantiert, dass alle Filtergewichtungen in den gemeinsamen Speicher eines Blocks übertragen werden, bevor einer der Threads dieses Blocks weiter ausgeführt wird. Ferner werden alle Ablesungen der Filtergewichte bereits aus dem allgemeinen Speicher des Blocks durchgeführt.
  3. Kopieren Sie das Ausgabebild aus dem Speicher des Geräts in den Speicher des Hosts, und geben Sie den zugewiesenen Speicher frei.

Kompilieren, ausführen (Eingabebildgröße - 557 x 313):
OpenCV time (ms):2
OpenMP time (ms):11
OpenMP similarity:0.00287131
CUDA time kernel (ms): 2.93245
CUDA time full (ms):32
CUDA similarity:0.00287131
CUDA time kernel (ms): 2.93402
CUDA time full (ms):4
CUDA similarity:0.00287131
CUDA time kernel (ms): 2.93267
CUDA time full (ms):4
CUDA similarity:0.00287131
CUDA time kernel (ms): 2.93312
CUDA time full (ms):4
CUDA similarity:0.00287131

Wie Sie sehen, haben wir, wenn Sie den allerersten Start der CUDA-Option nicht berücksichtigen, im Vergleich zur OpenMP-Option fast das Dreifache an Zeit gewonnen, obwohl wir die OpenCV-Option - die übrigens OpenCL verwendet - nicht erwischt haben .
Die Konfiguration der Maschine, auf der die Tests durchgeführt wurden:
Versteckter Text
Prozessor: Intel® Core (TM) i7-3615QM-CPU bei 2,30 GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4 GB, 1600 MHz.
OS: OS X 10.9.5.
Compiler: g ++ (GCC) 4.9.2 20141029.
CUDA-Compiler: Cuda-Kompilierungswerkzeuge, Release 6.0, V6.0.1.
Unterstützte Version von OpenMP: OpenMP 4.0.

Das war's für heute, im nächsten Teil werden wir uns einige grundlegende GPU-Algorithmen ansehen.
Der gesamte Quellcode ist auf bitbucket verfügbar .

Jetzt auch beliebt: