要解释 GPU 的体系结构对于挖矿程序性能的影响,需要从算法本身入手。但简单来说我们可以认为 GPU 挖矿的性能只与两点有关:1)GPU 整数运算单元的数量。2)程序能在目标体系结构上达到多高的利用率。因此直接回答题主的问题:AMD 的 VLIW5 和 VLIW4 在整数运算单元数量一致的情况下,谁的利用率高,谁的挖矿性能就越高。
挖矿程序的实质是计算 Hash,包括大量的整数位运算。后期的币种,如 LiteCoin 所使用的 Scrypt 算法还引入了大量相互依赖的、随机的访存指令,以弱化矿机(ASIC/FPGA)相较于 GPU 在整数运算性能上的优势。
以 bfgminer 中 Scrypt 算法的 OpenCL 实现为例,
首先是 Kernel 的主体结构,大致逻辑是并行的计算多个 Hash 值,并保存于显存:
void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
{
shittify(X);
const uint zSIZE = 8;
const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
const uint xSIZE = CONCURRENT_THREADS;
uint x = get_global_id(0)%xSIZE;
for(uint y=0; y<1024/LOOKUP_GAP; ++y)
{
#pragma unroll
for(uint z=0; z<zSIZE; ++z)
lookup[CO] = X[z];
for(uint i=0; i<LOOKUP_GAP; ++i)
salsa(X);
}
#if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
{
uint y = (1024/LOOKUP_GAP);
#pragma unroll
for(uint z=0; z<zSIZE; ++z)
lookup[CO] = X[z];
for(uint i=0; i<1024%LOOKUP_GAP; ++i)
salsa(X);
}
#endif
for (uint i=0; i<1024; ++i)
{
uint4 V[8];
uint j = X[7].x & K[85];
uint y = (j/LOOKUP_GAP);
#pragma unroll
for(uint z=0; z<zSIZE; ++z)
V[z] = lookup[CO];
#if (LOOKUP_GAP == 1)
#elif (LOOKUP_GAP == 2)
if (j&1)
salsa(V);
#else
uint val = j%LOOKUP_GAP;
for (uint z=0; z<val; ++z)
salsa(V);
#endif
#pragma unroll
for(uint z=0; z<zSIZE; ++z)
X[z] ^= V[z];
salsa(X);
}
unshittify(X);
}
Scrypt 算法的核心是 Salsa20/8,通过此 Hash 算法产生大量随机的访存。在 Footprint 足够大的条件下,GPU 会在 L2 级别、甚至 TLB 级别产生大量的缓存失效,从而产生大量的 DRAM 访问。
void salsa(uint4 B[8])
{
uint4 w[4];
#pragma unroll
for(uint i=0; i<4; ++i)
w[i] = (B[i]^=B[i+4]);
#pragma unroll
for(uint i=0; i<4; ++i)
{
w[0] ^= rotl(w[3] +w[2] , 7U);
w[1] ^= rotl(w[0] +w[3] , 9U);
w[2] ^= rotl(w[1] +w[0] ,13U);
w[3] ^= rotl(w[2] +w[1] ,18U);
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
}
#pragma unroll
for(uint i=0; i<4; ++i)
w[i] = (B[i+4]^=(B[i]+=w[i]));
#pragma unroll
for(uint i=0; i<4; ++i)
{
w[0] ^= rotl(w[3] +w[2] , 7U);
w[1] ^= rotl(w[0] +w[3] , 9U);
w[2] ^= rotl(w[1] +w[0] ,13U);
w[3] ^= rotl(w[2] +w[1] ,18U);
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
}
#pragma unroll
for(uint i=0; i<4; ++i)
B[i+4] += w[i];
}
理想的情况下 Scrypt 算法的瓶颈会出现在内存带宽上,而这也是 ASIC/FPGA 也无法规避的问题,因此理论上 GPU 与矿机的性能应该基本一致。当然了,实际情况下 GPU 的访存带宽受限于种种条件(如整数运算单元不足,寄存器数量有限等),往往是无法将访存带宽利用满的。
再回到题主的问题,从 GPU 体系结构的角度来看,从早期的 VLIW、Fermi 到现在的 GCN 、Maxwell,主要的改进方向就是尽可能的挖掘计算单元和访存单元的利用率。简单来看,越新的架构在同样的芯片面积下性能也会越好。同样的,新架构带来的高利用率还可以降低功耗,带来更高的能耗比(Perf/Watt)。
所以买新卡总不会错的 ^_^
— 完 —
本文作者:吕超
【知乎日报】
你都看到这啦,快来点我嘛 Σ(▼□▼メ)
此问题还有 1 个回答,查看全部。
延伸阅读:
挖矿是如何产生比特币的?
比特币 (Bitcoin) 是什么,如何简单易懂地介绍比特币?