Nvidia Jetson Tx1 против Jetson NANO (Бенчмаркинг) - PullRequest
Новая
       60

Nvidia Jetson Tx1 против Jetson NANO (Бенчмаркинг)

1 голос
/ 09 июля 2019

В настоящее время я пытаюсь сравнить Jetson TX1 с NANO jetson, согласно https://elinux.org/Jetson, они оба имеют архитектуру maxwell с 128 ядрами cuda для NANO и 256 для TX1.Это означает, что обычно Jetson NANO достигает половины производительности TX1.

Чтобы проверить это, я создал ядро ​​умножения одной операции (с плавающей запятой) следующим образом: 

__global__ void matrixMultiply(float* mat1, float* mat2, int nx, int ny)
    {
        unsigned int ix = threadIdx.x + blockDim.x*blockIdx.x;
        unsigned int iy = threadIdx.y + blockDim.y*blockIdx.y;
        int idx = iy*nx + ix;

        mat1[idx] = mat1[idx]*mat2[idx] ;

    }

Тест:умножение 2 "массив с плавающей запятой размером 15000 * 15000" привело к TX1 = 130 мс и Jetson NANO = 150 мс.Результат кажется странным, как будто я не использую второй SM TX1, поэтому я профилировал с использованием sm_efficiency (TX1 и NANO = 100%), достигается_потребностью (TX1 = 92%, NANO = 88%).Я что-то здесь упускаю или просто не использую правильную конфигурацию сетки и блока.

PS: я перепробовал все возможные конфигурации, и лучшей конфигурацией для обеих платформ был блок (256, 1) исетка рассчитывается соответственно.

1 Ответ

3 голосов
/ 10 июля 2019

Я что-то здесь упускаю

Да, вы что-то здесь упускаете. Ваш код не соответствует тому, что вы думаете:

они оба имеют архитектуру maxwell с 128 ядрами cuda для NANO и 256 для TX1. Это означает, что обычно Jetson NANO достигает половины производительности TX1.

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

Начнем с некоторых спецификаций: 

spec                 | TX1         | Nano     | source
---------------------=-------------=----------=----------
mem bandwidth (GB/s) | 25.6        | 25.6     | 1,2
---------------------=-------------=----------=----------
(FP32) compute cores | 256         | 128      | 1,2
---------------------=-------------=----------=----------
max core clock (MHz) | 998         | 921      | 1,2

источники: 1 , 2

Чтобы вычислить максимальную теоретическую пропускную способность FP32, формула : 

# of SMs * # of FP32 units per SM * 2 * clock rate

Для Jetson NANO: 

128 * 2 * 921MHz = ~236GFlops/s

Для Jetson TX1: 

256 * 2 * 998MHz = ~511GFlops/s

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

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

    mat1[idx] = mat1[idx]*mat2[idx] ;

Мы видим, что для каждой операции умножения FP32 мы должны прочитать две величины (всего 8 байт) и записать одну величину (всего 4 байта). Таким образом, 12 байтов для чтения / записи для каждой операции умножения.

Теперь давайте предположим, что вы можете достичь максимальной пропускной способности на TX1 511GFlops / s. Это 511 000 000 000 операций умножения-добавления в секунду, или ~ 256 000 000 000 операций умножения. Если бы вы могли выполнить операции умножения 256B в секунду, для каждого умножения потребовалось бы 12 байтов операций чтения / записи, поэтому общая требуемая пропускная способность составила бы: 

256,000,000,000 multiply ops              12 bytes        3,072,000,000,000 bytes
----------------------------    *        -----------   =  -----------------------
            sec                          multiply op              sec

Это означает, что потребуется около 3 терабайт в секунду пропускной способности памяти, чтобы ваш код был ограничен пропускной способностью вычислений TX1. Но TX1 имеет только 25,6 гигабайт в секунду пропускной способности памяти. Таким образом, пропускная способность памяти TX1 будет ограничивать пропускную способность вашего кода. Аналогичный расчет показывает, что пропускная способность памяти NANO также ограничивает пропускную способность вашего кода, и поэтому предиктор для соотношения производительности между двумя для вашего кода является отношением пропускной способности памяти: 

25.6GB/s
--------     = 1
25.6GB/s

Поэтому тот факт, что вы наблюдали почти одинаковую производительность между двумя: 

150
---          = 1.15
130

- гораздо более разумный результат для вашего кода, чем ожидалось, что соотношение производительности будет 2: 1.

Если вы хотите увидеть код, который приближается к соотношению 2: 1, вам понадобится код, который выполняет много вычислительных операций, при этом потребляя (условно говоря) почти нет пропускной способности памяти. Возможным примером реального кода такого кода может быть умножение матрицы на матрицу, и вы можете легко написать код CUBLAS Sgemm , чтобы проверить это. Обратите внимание, что ожидание отношения 2: 1 не совсем верно, потому что тактовые частоты ядра не совпадают. Ожидаемое соотношение будет: 

511
--- = ~2.17
236
...