根据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 =130ms,Jetson NANO =150ms时,2“大小为15000*15000”的浮点数组相乘。结果似乎很奇怪,好像我没有使用TX1的第二个SM,因此我使用了sm_efficiency (TX1和NANO = 100%),achieved_occupancy (TX1 = 92%,NANO =88%)进行分析。我是不是遗漏了什么,或者我只是没有使用正确的网格和块配置。
附言:我尝试了所有可能的配置,两个平台的最佳配置都是(256,1)的块,并相应地计算网格。
发布于 2019-07-10 07:02:16
我是不是漏掉了什么
是的,你在这里遗漏了一些东西。你的代码没有衡量你的想法:
它们都具有麦克斯韦尔架构,具有128个用于NANO的cuda内核和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
要计算最大理论FP32计算吞吐量,the formula为:
# 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
这意味着您的代码将受到TX1计算吞吐量的限制,这将需要大约3 to /秒的内存带宽。但是TX1每秒只有25.6 of的内存带宽。因此,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
https://stackoverflow.com/questions/56953183
复制相似问题