Shader Unit是否计算指数

Question

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() and __expf()提到可以在cuda中计算指数。

假设我有 500 000 000（五亿）个双打数组。我想计算数组中每个值的指数。谁知道会发生什么：5760 个着色器单元将能够计算 exp，或者仅使用两个 GK110 GPU 即可完成此任务？性能差异很大，所以我需要确定，如果我用 CUDA 重写我的应用程序，那么它的运行速度不会变慢。

换句话说，我可以创建 5760 个线程来计算 500 000 000 个指数吗？

Answer 1

GTX Titan Z 是双 GPU 设备。卡上的两个 GK110 GPU 中的每一个都通过 384 位内存接口连接到其自己的 6 GB 高速内存。每个内存的理论带宽为336GB/sec。 GTX Titan Z 中使用的特定 GK110 变体由十五个称为 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 通过各种整数和浮点单元提供通用计算。 GPU 还提供特殊功能单元（在 GK110 上称为 MUFU = multifunction 单元）可以计算一些常用的粗略单精度近似值倒数、平方根倒数、正弦、余弦、以 2 为底的指数和以 2 为底的对数等函数。就求幂而言，标准的单精度数学函数 exp2f() 是唯一映射或多或少映射的函数直接到 MUFU 指令 (MUFU.EX2)。根据编译模式，此硬件指令周围有一个薄包装器，因为硬件不支持特殊功能单元中的非规范操作数。

CUDA 中的所有其他求幂都是通过软件子程序执行的。标准单精度函数 expf() 是围绕硬件 exp2 功能的相当重量级的包装器。双精度exp()函数是基于minimax多项式逼近的纯软件例程。它的完整源代码在 CUDA 头文件 math_functions_dbl_ptx3.h 中可见（在 CUDA 6.5 中，DP exp() 代码从该文件的第 1706 行开始）。如您所见，计算主要涉及双精度浮点运算，以及整数和一些单精度浮点运算。您还可以通过反汇编使用 cuobjdump --dump-sass.

调用 exp() 的二进制可执行文件来查看机器码

就性能而言，在 CUDA 6.5 中，双精度 exp() 函数在 Tesla K20 (1.170 DP TFLOPS) 上的吞吐量约为每秒 25e9 次函数调用。由于每次调用 DP exp() 都会消耗一个 8 字节的源操作数并产生一个 8 字节的结果，这相当于大约 400 GB/sec 的内存带宽。由于 Titan Z 上的每个 GK110 提供的性能比 Tesla K20 上的 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);
    // ...
}

Shader Unit是否计算指数

Does Shader Unit calculates exponent

cuda