着色单位计算指数吗？

Question

http://us.hardware.info/reviews/5419/nvidia-geforce-gtx-titan-z-sli-review-incl-tones-tizair-system说"GTX泰坦-Z“有5760个着色单位。这里还写着"GTX泰坦-Z“有2x GK110 GPU。

expf()提出了计算库达指数的可能性。

假设我有50万(5亿)双倍。我要计算数组中每个值的指数。谁知道会发生什么: 5760个着色器单位将能够计算exp，或者这个任务只能用两个GK110 GPU完成？性能上的差异是可怕的，所以我需要确定，如果我用CUDA重写我的应用程序，那么它将不会工作得更慢。

换句话说，我是否可以制造5760个线程来计算50万指数？

Answer 1

GTX泰坦Z是一个双GPU设备。卡上的两个GK110 GPU中的每一个都通过一个384位内存接口连接到它自己的6GB高速内存上。每个内存的理论带宽为336 GB/秒。GTX中使用的特定GK110变体由15个称为SMX的执行单元组成。每个SMX依次由192个单精度浮点单元、64个双精度浮点单元和各种其他单元组成。

在GK110中，每个双精度单元都可以执行一个FMA (融合乘法加法)，或者一个FMUL，或者每个时钟周期一个FADD。在基准时钟为705 MHz时，泰坦Z上每个GK110 GPU每秒可执行的最大DP操作总数为705e6 * 15 * 64 = 676.8e9。假设所有操作都是FMAs，则等于1.3536双精度TFLOPS.由于该卡使用两个GPU，因此GTX泰坦Z的总DP性能为2.7072 TFLOPS。

与CPU一样，GPU通过各种整数和浮点单元提供通用计算.GPU还提供了特殊的函数单元(在multifunction上称为MUFU = GK110)，它可以计算一些常用函数的粗略单精度逼近，如倒数、倒数平方根、正弦、余弦、指数基2和基于对数的2。就指数计算而言，标准的单精度数学函数exp2f()是唯一或多或少直接映射到MUFU指令(MUFU.EX2)的函数。根据编译模式的不同，这个硬件指令有一个很薄的包装器，因为硬件不支持特殊功能单元中的非正常操作数。

数据自动化系统中的所有其他指数都是通过软件子程序实现的。标准的单精度函数expf()是硬件的exp2功能的一个相当重的包装器.双精度exp()函数是基于极小极大多项式逼近的纯软件程序.它的完整源代码可以在CUDA头文件math_functions_dbl_ptx3.h中看到(在CUDA 6.5中，DP exp()代码从该文件的第1706行开始)。如您所见，计算主要涉及双精度浮点运算，以及整数运算和一些单精度浮点运算。您还可以通过分解一个用exp()调用cuobjdump --dump-sass的二进制可执行文件来查看机器代码。

在性能方面，在CUDA6.5中，双精度exp()函数的吞吐量是每秒25e9个函数调用在特斯拉K20 (1.170 DP TFLOPS)。因为每次对DP exp()的调用都会消耗一个8字节的源操作数并产生一个8字节的结果，这相当于大约400 GB/秒的内存带宽。由于泰坦Z上的每个GK110比特斯拉GK110上的GK110提供了大约15%的性能，因此吞吐量和带宽需求也相应增加。由于所需带宽超过了GPU的理论内存带宽，因此将DP exp()应用于数组的代码将完全受内存带宽的约束。

GPU中的功能单元数和执行线程的数量与可以处理的数组元素的数量没有关系，但会对这种处理的性能产生影响。编程人员可以自由选择数组元素到线程的映射。可以在一次访问中处理的数组元素数是GPU内存大小的函数。请注意，并非设备上的所有原始内存都可供用户代码使用，因为CUDA软件栈需要一些内存供其自己使用，通常大约为100 MB左右。用于将DP exp()应用于数组的示例性映射在以下代码片段中显示：

__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);
    // ...
}