http://us.hardware.info/reviews/5419/nvidia-geforce-gtx-titan-z-sli-review-incl-tones-tizair-system 说"GTX Titan-Z"有5760个着色器单元.此外还写道"GTX Titan-Z"拥有2x GK110 GPU.
CUDA exp()expf()和__expf()提到可以用cuda计算指数.
假设我有5亿(5亿)双打数组.我想计算数组中每个值的指数.谁知道会发生什么:5760着色器单元将能够计算exp,或者这个任务只能用两个GK110 GPU完成？性能的差异非常大,所以我需要确定,如果我用CUDA重写我的应用程序,那么它将无法更慢地运行.
换句话说,我可以制作5760个线程来计算5亿个指数吗？
GTX Titan Z是双GPU设备.卡上的两个GK110 GPU中的每一个都通过384位存储器接口连接到自己的6 GB高速存储器.每个内存的理论带宽为336 GB /秒.GTX Titan Z中使用的特定GK110变体由15个执行单元簇组成,称为SMX.每个SMX依次由192个单精度浮点单元,64个双精度浮点单元和各种其他单元组成.
GK110中的每个双精度单元可以在每个时钟周期执行一个FMA(融合乘法 - 加法),或一个FMUL或一个FADD.在705 MHz的基本时钟,因此每个Titan Z上每个GK110 GPU可执行的DP操作的最大总数为705e6*15*64 = 676.8e9.假设所有操作都是FMA,则相当于1.3536双精度TFLOPS.由于该卡使用两个GPU,因此GTX Titan Z的总DP性能为2.7072 TFLOPS.
与CPU一样,GPU通过各种整数和浮点单元提供通用计算.图形处理器还提供特殊功能单元(称为幕阜= 亩 LTI 福基于GK110 nction单元),其可以计算粗单精度近似一些常用的功能,如倒数,倒数平方根,正弦,余弦,指数基体2,和对数2.就指数而言,标准单精度数学函数exp2f()是或多或少直接映射到MUFU指令(MUFU.EX2)的唯一函数.根据编译模式,该硬件指令周围有一个薄的包装器,因为硬件不支持特殊功能单元中的非正规操作数.
CUDA中的所有其他指数都是通过软件子程序执行的.标准的单精度功能expf()是围绕硬件功能的相当重量级的包装器exp2.双精度exp()函数是基于极小极大多项式近似的纯软件程序.它的完整源代码在CUDA头文件中可见math_functions_dbl_ptx3.h(在CUDA 6.5中,DP exp()代码从该文件中的第1706行开始).如您所见,计算主要涉及双精度浮点运算,以及整数和一些单精度浮点运算.您还可以在机器代码通过拆卸调用二进制可执行exp()用cuobjdump --dump-sass.
在性能方面,在CUDA 6.5中,双精度exp()函数在特斯拉K20(1.170 DP TFLOPS)上的吞吐量大约为每秒25e9次函数调用.由于每次调用DP都会exp()消耗8字节的源操作数并产生8字节的结果,因此这相当于大约400 GB /秒的内存带宽.由于Titan Z上的每个GK110比Tesla K20上的GK110提供的性能提高了约15%,因此吞吐量和带宽要求也相应提高.由于所需带宽超过了GPU的理论内存带宽,因此将DP简单地应用于exp()阵列的代码将完全受内存带宽的限制.
GPU中的功能单元的数量和执行的线程数与可以处理的阵列元素的数量无关,但是可以对这种处理的性能产生影响.程序员可以自由选择数组元素到线程的映射.可以一次处理的数组元素的数量是GPU内存大小的函数.请注意,并非设备上的所有原始内存都可用于用户代码,因为CUDA软件堆栈需要一些内存供自己使用,通常大约100 MB左右.exp()此代码段中显示了将DP 应用于阵列的示例性映射:
__global__ void exp_kernel (const double * __restrict__ src,
double * __restrict__ dst, int len)
{
int stride = gridDim.x * blockDim.x;
int tid = blockDim.x * blockIdx.x + threadIdx.x;
for (int i = tid; i < len; i += stride) {
dst[i] = exp (src[i]);
}
}
#define ARRAY_LENGTH (500000000)
#define THREADS_PER_BLOCK (256)
int main (void) {
// ...
int len = ARRAY_LENGTH;
dim3 dimBlock(THREADS_PER_BLOCK);
int threadBlocks = (len + (dimBlock.x - 1)) / dimBlock.x;
if (threadBlocks > 65520) threadBlocks = 65520;
dim3 dimGrid(threadBlocks);
double *d_a = 0, *d_b = 0;
cudaMalloc((void**)&d_a, sizeof(d_a[0]), len);
cudaMalloc((void**)&d_b, sizeof(d_b[0]), len);
// ...
exp_kernel<<<dimGrid,dimBlock>>>(d_a, d_b, len);
// ...
}