CUDA开发笔记---向量运算对比

2020-08-27

最近开始学习CUDA。正式开始写代码前查了一些资料,对CUDA的印象就是擅长浮点数运算,缺点就是如果数据量不足够大,就不一定比CPU快。自己尝试了一些简单的对比后稍微有些理解了。

电脑配置:i7-6700HQ,GTX970m,16G。

测试使用三个CUDA官方例子中的三维向量计算函数:

inline __host__ __device__ float3 normalize(float3 v)
{
    float invLen = rsqrtf(dot(v, v));
    return v * invLen;
}
inline __host__ __device__ float length(float3 v)
{
    return sqrtf(dot(v, v));
}
inline __host__ __device__ float dot(float3 a, float3 b)
{
    return a.x * b.x + a.y * b.y + a.z * b.z;
}

其中用到的其他函数是CUDA内置的sqrtf()和基于sqrtf()的rsqrtf():平方根和平方根倒数

extern __CUDA_MATH_CRTIMP __DEVICE_FUNCTIONS_DECL__ __device_builtin__ float __cdecl sqrtf(float) __THROW; 
inline float rsqrtf(float x)
{
    return 1.0f / sqrtf(x);
}

输入: float3 a :(99, 99, 99); float3 b :(999, 999, 999);

时间单位ms。

第一次测试:使用以上三个函数计算100万次:normalize(a); length(a-b); dot(a,b).

  normalize(a) length(a-b) dot(a,b)
CPU用时 132 141 32
CUDA用时 None None None

CUDA下直接爆内存卡死。可能是我的配置扛不住连续一百万次cudaMemcpy()。

第二次测试:换成100次。

  normalize(a) length(a-b) dot(a,b)
CPU用时 0 0 0
CUDA用时 109 14 16

结论:差距巨大,按照第一次测试的结果估算,这种情况下CPU比CUDA快5333倍。

第三次测试:把单独输入换成数组一次输入。以下是不同大小数组的结果。

  CPU nor CUDA nor CPU len CUDA len CPU dot CUDA dot
100 0 1 0 0 0 0
1000 0 1 0 0 0 0
10000 1 2 1 2 0 1
100000 14 21 14 11 3 13
1000000 139 143 146 103 34 107
10000000 1395 1455 1460 1064 361 1056
100000000 14254 17737 14661 14809 3420 10722

这次的结果就有意思了。10000长度以下差别不是很大,10万次以上CUDA的length(a-b)明显比CPU快,但其他的还是比不过CPU。一亿次时爆了内存,程序卡了,CUDA的length(a-b)优势也因为爆内存而消失了。

这就说明了CUDA的好几个特点。首先是输入量大的时候,把大量变量一起传给CUDA运算会比单个传快得多,因为单个传只用了CUDA的一个线程。在数据运算量不大的时候,CUDA和CPU差距也不大,甚至CPU还全面赶超CUDA。当数据量足够大时,CUDA的length(a-b)比CPU快了差不多一半,但其他的还是慢。

从函数上看,比dot的话,整数运算CUDA总是比CPU慢。 length函数中也是调用了dot,只不过是dot的两个输入相同。CUDA算dot应该比CPU慢,可是sqrtf()让CPU的length比它自己的dot运算慢了差不多四倍。输入a,b与a-b的量级差不多,所以差别就体现在sqrtf()上。CUDA的length就和它的dot差不多,应该是因为CUDA擅长浮点数运算,所以它的length就比CPU的快。

可是刚刚的结论似乎没法解释normalize上CUDA和CPU差距不是很大的事实。normalize只是比length多了一个倒数和乘法运算。理论上应该和length差不多,实际却比length慢一半。换个角度,对比两者的dot和normalize,CUDA的这两个运算差距不是很大,normalize只比dot慢了一半。而CPU的normalize几乎用了它dot的四倍时间。其中的差异应该还是rsqrtf()的sqrtf()造成的。

试了试嵌套调用normalize,第一次嵌套就比CPU快了一点。CPU/CUDA的时间比随着嵌套次数开始增加,1.04,1.07,1.18,1.13,1.17。normalize嵌套意义可能不大,因为嵌套之后每次运算过程都一样。

以上结果都没有取平均值,不过还是能得出一些初步的结论。比如在牵扯到输入量巨大的浮点数运算时,CUDA才会展现出优势。