Может кто-нибудь, пожалуйста, помогите мне с очень простым примером использования общей памяти? Пример, включенный в руководство по программированию Cuda C, кажется загроможденным нерелевантными деталями.
Например, если я скопирую большой массив в глобальную память устройства и хочу соединить каждый элемент, как можно использовать общую память для ускорения этого? Или это не полезно в этом случае?
В конкретном случае, о котором вы упоминаете, разделяемая память бесполезна по следующей причине: каждый элемент данных используется только один раз. Чтобы общая память была полезной, вы должны использовать данные, перенесенные в общую память несколько раз, используя хорошие шаблоны доступа, чтобы это помогло. Причина этого проста: простое чтение из глобальной памяти требует 1 чтения из глобальной памяти и нулевого чтения из общей памяти; сначала для чтения его в общую память потребуется 1 чтение глобальной памяти и 1 чтение общей памяти, что занимает больше времени.
Вот простой пример, где каждый поток в блоке вычисляет соответствующее значение в квадрате плюс среднее значение для левого и правого соседей в квадрате:
__global__ void compute_it(float *data)
{
int tid = threadIdx.x;
__shared__ float myblock[1024];
float tmp;
// load the thread 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 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).
Подумайте об общей памяти как к явному управляемому кешу - это полезно только в том случае, если вам нужно обращаться к данным более одного раза, либо в пределах одного потока, либо из разных потоков внутри одного и того же блока. Если вы только получаете доступ к данным один раз, тогда общая память не поможет вам.