要解释 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) 是什么,如何简单易懂地介绍比特币?

分享到