roo*_*ody 3 cuda nvidia nvidia-jetson
我目前正在尝试将 Jetson TX1 与 jetson NANO 进行基准测试，根据https://elinux.org/Jetson，它们都具有 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%) ，reached_occupancy (TX1 = 92%, NANO = 88%) 进行了分析。我在这里遗漏了什么，或者我只是没有使用正确的网格和块配置。
PS：我尝试了所有可能的配置，两个平台的最佳配置是一个 (256, 1) 块和相应计算的网格。
我在这里错过了什么吗
是的，你在这里遗漏了一些东西。您的代码不会衡量您的想法：
它们都具有 maxwell 架构，NANO 有 128 个 cuda 内核，TX1 有 256 个 cuda 内核。这意味着通常 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
要计算最大理论 FP32 计算吞吐量，公式为：
# 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