我编写了两个简单的内核。每个添加两个类型为int(32位)/ long int(64位)的向量。事实证明,在我的GPU(特斯拉K80)上,这款产品刚好相当新颖,内核只有32位。
随着矢量大小的增加,时间大致加倍。
的内核如下:
__global__ void add_32(int * c, int * a, int * b)
{
int gid = blockIdx.x * blockDim.x + threadIdx.x;
c[gid] = a[gid] + b[gid];
}
typedef long int int64;
__global__ void add_64(int64 * c, int64 * a, int64 * b)
{
int gid = blockIdx.x * blockDim.x + threadIdx.x;
c[gid] = a[gid] + b[gid];
}
当矢量大小为1兆个元件,add_32约需102.911微秒,而add_64需要192.669微秒。 (运行释放模式二进制时,使用Nvidia分析器报告执行时间)。
看来,64位指令只是通过32位指令仿真!
这可能是一个蛮力的解决方案,以找出什么样的机器是GPU核心,但绝对不是一个优雅的。
更新:
感谢@保罗A.克莱顿评论,似乎数据的大小在64位的情况下加倍上述解决方案是不公平的比较。所以我们不应该使用相同数量的元素来启动这两个内核。正确的原则是启动64位版本的元素数量的一半。
为了更加确定,我们来考虑一下元素向量乘法而不是加法。如果GPU通过32位指令模拟64位指令,那么它至少需要3个32位乘法指令,例如使用Karatsuba算法来乘以2个64位数字。这意味着如果我们启动具有N/2个元素的64位向量乘法内核,如果只是模拟64位乘法,则需要比具有N个元素的32位内核更长的时间。
下面是内核:
__global__ void mul_32(int * c, int * a, int * b)
{
int gid = blockIdx.x * blockDim.x + threadIdx.x;
c[gid] = a[gid] * b[gid];
}
typedef long int int64;
__global__ void mul_64(int64 * c, int64 * a, int64 * b)
{
int gid = blockIdx.x * blockDim.x + threadIdx.x;
c[gid] = a[gid] * b[gid];
}
而这里是实验细节: 时报报道这里是从NVIDIA分析器上释放模式二进制: 1-内核mul_32与矢量大小为N = 256兆个元件,需要25.608毫秒。 2-矢量大小N = 128的mul_64兆元素,需要24.153毫秒。
我知道这两个内核产生不正确的结果,但我认为这与计算方式没有任何关系。
我对Nvidia GPU并不十分熟悉,但这些信息应该在数据手册或其他手册中。如果公众中没有信息,那么您可能需要与Nvidia签署NDA才能获得此信息。那么你有权访问你所瞄准的GPU的文档吗? – fsasm
GPU规格不显示这些信息。我认为应该有一个API /命令可以告诉这样的信息! – caesar
一般而言,数据表应显示此类信息,因为这是他们的目的。如果供应商不公布信息,那么你不需要它。驱动程序与PTX一起隐藏硬件的所有细节以增加可移植性。如果你真的需要这些信息,你应该联系Nvidia。 – fsasm