Когда полезна память CUDA __shared__?

Question

Может кто-нибудь помочь мне с очень простым примером того, как использовать разделяемую память?Пример, включенный в руководство по программированию на Cuda C., кажется загроможденным несущественными деталями.

Например, если я скопирую большой массив в глобальную память устройства и захочу привести в квадрат каждый элемент, как можно использовать общую память для ускоренияэто вверх?Или это не полезно в этом случае?

Answer 1

В конкретном случае, о котором вы упоминаете, разделяемая память бесполезна по следующей причине: каждый элемент данных используется только один раз.Чтобы общая память была полезной, вы должны использовать данные, переданные в общую память несколько раз, используя хорошие шаблоны доступа, чтобы это помогло.Причина этого проста: простое чтение из глобальной памяти требует 1 чтения из глобальной памяти и нулевого чтения из общей памяти;сначала для чтения его в совместно используемую память потребуется 1 чтение глобальной памяти и 1 чтение общей памяти, что занимает больше времени.

Вот простой пример, где каждый поток в блоке вычисляет соответствующее значение в квадрате плюс среднее значениекак его левые, так и правые соседи в квадрате:

  __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] = myblock[tid];
  }

Обратите внимание, что предполагается работать с использованием только одного блока.Расширение на большее количество блоков должно быть простым.Предполагается размер блока (1024, 1, 1) и размер сетки (1, 1, 1).

Answer 2

Думайте о разделяемой памяти как о явно управляемом кеше - это полезно, только если вам нужен доступ к данным более одного раза, либо в одном и том же потоке, либо из разных потоков в одном и том же блоке .Если вы обращаетесь к данным только один раз, то общая память вам не поможет.