2013-03-17 7 views
1

Я измерил пропускную способность при выполнении ядра делает сложение между двумя векторами:CUDA: векторы сложения и векторов размер

__global__ void add(float *a, float *b, float *c, int n) 
{ 
    int tid = blockIdx.x*blockDim.x + threadIdx.x; 

    while (tid < n) 
    { 
     c[tid] = a[tid] + b [tid]; 
     tid += blockDim.x * gridDim.x; 
    } 
} 

я во-первых, запустить ядро ​​один раз, так что он будет загружен на устройстве, а рядом , Я измеряю 10 итераций выполнения ядра.

Пропускная способность намного лучше, когда мой vectors'length является +1000000 1000.

Почему?

Спасибо.

+2

Графический процессор может скрывать задержку памяти, когда у него больше работы. –

ответ

5

Графические процессоры имеют очень высокую пропускную способность памяти (хорошо) и очень высокая латентность памяти (плохо). Требуется несколько сотен циклов, чтобы запрос на память был удовлетворен, но может быть много запросов в полете, которые конвейерны.

Если вы добавляете только 1000 элементов, то вы действительно измеряете латентность памяти.

С гораздо большими векторами будет достаточно блоков потоков, когда при остановке в памяти графический процессор может переключиться на другой блок потока, запрос на память которого выполняется. С ними вы измеряете пропускную способность.

+0

Но из того, что я читал, данные должны считываться блоками по 128 байт (перекосы, коалесцирующие ...). Если это так, для каждого блока должна быть латентность, и поскольку я сравниваю 1000 и 1000000 элементов, разница должна быть линейной, нет?
Было бы различие, если бы данные были прочитаны один раз в двух случаях. Чем я пойму, почему латентность относительно важна с 1000 элм. – Oodini

+0

Он не линейный из-за конвейерной обработки. Предположим, что каждый блок считывает 128 байтов и занимает 300 циклов латентности для запроса памяти. Предположим, у вас 300 блоков. При тактовом галочке 0 первый блок делает запрос на память, который будет удовлетворяться с помощью тактового сигнала 300. В такте 1 галочка выполняется другой блок (контекстные переключатели свободны на графических процессорах), и он останавливается до отметки 301. Последний блок выдает запрос на память на отметке 299, и он будет удовлетворен тиком 599. Таким образом, все 300 блоков будут завершены в 600 циклов, а не 300 * 300 = 90 000. Фактическая задержка зависит от вашей конкретной карты. –

Смежные вопросы