我編寫了兩個簡單的內核。每個添加兩個類型爲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