3

根据https://elinux.org/Jetson的说法，我目前正在尝试将 Jetson TX1 与 Jetson NANO 进行基准测试，它们都具有 maxwell 架构，NANO 具有 128 个 cuda 内核，TX1 具有 256 个内核。这意味着通常 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%) 、atained_occupancy (TX1 = 92%, NANO = 88 %) 进行了分析。我在这里遗漏了什么，或者我只是没有使用正确的网格和块配置。

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

4

1 回答 1

7

我在这里错过了什么吗

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

它们都具有 maxwell 架构，为 NANO 提供 128 个 cuda 内核，为 TX1 提供 256 个内核。这意味着通常 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 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
于 2019-07-09T23:02:15.927 回答