cuda - Nvidia Jetson Tx1 与 jetson NANO(基准测试)

Question

根据 https://elinux.org/Jetson，我目前正在尝试将 Jetson TX1 与 jetson NANO 进行基准测试。 ，它们都具有 maxwell 架构，NANO 有 128 个 cuda 内核，TX1 有 256 个 cuda 内核。这意味着通常 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] ;

    }

测试:TX1 = 130 ms 和 Jetson NANO = 150 ms 乘以 2“大小为 15000*15000 的浮点数组”。结果看起来很奇怪，就像我没有使用TX1的第二个SM，
因此我使用 sm_efficiency (TX1 and NANO = 100%) 进行了分析，
达到的占用率(TX1 = 92%，NANO = 88%)。我在这里遗漏了什么，或者我只是没有使用正确的网格和块配置。

P.S:我尝试了所有可能的配置，两个平台的最佳配置是 (256, 1) 块和相应计算的网格。

Answer 1

Am I missing something here

是的，你在这里遗漏了一些东西。您的代码不会衡量您的想法:

they both have the maxwell architecture with 128 cuda cores for NANO and 256 for TX1. This means that normally Jetson NANO will achieve half the performance of the 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 计算吞吐量， the formula是:

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

对于杰森纳米:

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 TB 的内存带宽，因为您的代码将受到 TX1 计算吞吐量的限制。但 TX1 每秒只有 25.6 GB 的内存带宽。所以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