Wann ist der __shared__-Speicher von CUDA nützlich?

Kann mir bitte jemand mit einem sehr einfachen Beispiel helfen, wie man Shared Memory verwendet? Das im Cuda C-Programmierhandbuch enthaltene Beispiel scheint mit irrelevanten Details überladen zu sein.

Wenn ich beispielsweise ein großes Array in den globalen Speicher des Geräts kopiere und jedes Element quadrieren möchte, wie kann der gemeinsame Speicher verwendet werden, um dies zu beschleunigen? Oder ist es in diesem Fall nicht sinnvoll?

In dem von Ihnen erwähnten speziellen Fall ist Shared Memory aus folgendem Grund nicht sinnvoll: Jedes Datenelement wird nur einmal verwendet. Damit Shared Memory nützlich ist, müssen Sie Daten, die an Shared Memory übertragen werden, mehrmals verwenden und gute Zugriffsmuster verwenden, damit es hilft. Der Grund dafür ist einfach: Allein das Lesen aus dem globalen Speicher erfordert 1 Lesen des globalen Speichers und null Lesevorgänge des gemeinsam genutzten Speichers; es zuerst in den gemeinsam genutzten Speicher zu lesen, würde 1 Lesen des globalen Speichers und 1 Lesen des gemeinsam genutzten Speichers erfordern, was länger dauert.

Hier ist ein einfaches Beispiel, bei dem jeder Thread im Block den entsprechenden Wert zum Quadrat plus den Durchschnitt seiner linken und rechten Nachbarn zum Quadrat berechnet:

  __global__ void compute_it(float *data)
  {
     int tid = threadIdx.x;
     __shared__ float myblock[1024];
     float tmp;

     // load the thread's data element into shared memory
     myblock[tid] = data[tid];

     // ensure that all threads have loaded their values into
     // shared memory; otherwise, one thread might be computing
     // on unitialized data.
     __syncthreads();

     // compute the average of this thread's left and right neighbors
     tmp = (myblock[tid > 0 ? tid - 1 : 1023] + myblock[tid < 1023 ? tid + 1 : 0]) * 0.5f;
     // square the previousr result and add my value, squared
     tmp = tmp*tmp + myblock[tid] * myblock[tid];

     // write the result back to global memory
     data[tid] = tmp;
  }

Beachten Sie, dass dies mit nur einem Block funktionieren soll. Die Erweiterung auf weitere Blöcke sollte unkompliziert sein. Geht von Blockmaß (1024, 1, 1) und Rastermaß (1, 1, 1) aus.

Stellen Sie sich Shared Memory als einen vor explizit verwalteter Cache – Es ist nur nützlich, wenn Sie mehr als einmal auf Daten zugreifen müssen, entweder innerhalb desselben Threads oder von verschiedenen Threads innerhalb desselben Block. Wenn Sie nur einmal auf Daten zugreifen, hilft Ihnen Shared Memory nicht weiter.