CUDA C编程（二十二）程序优化指令

  用于优化程序的指令，有很多的选择：单精度或双精度浮点值、标准或内部函数、原子函数或不安全访问。一般情况下，每一个选择在性能、精确度和正确性上都有不同表现。接下来，将比较分析每一类指令的优缺点。

单 精 度 与 双 精 度 的 比 较
  用于存储单精度和双精度数的位数是不同的，因此，双精度变量相较于单精度变量来说，可以在一个更精细的粒度和更广泛的范围上表示不同的数值。为了证实这一点，运行下面的代码，这个程序在主机和设备端将数值12.1分别存储为单精度变量和双精度变量，然后按20个小数位存储的实际值进行输出。

#include <stdio.h>
#include <stdlib.h>

__global__ void kernel(float *F, double *D)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if (tid == 0)
    {
        *F = 12.1;
        *D = 12.1;
    }
}

int main(int argc, char **argv)
{
    float *deviceF;
    float h_deviceF;
    double *deviceD;
    double h_deviceD;

    float hostF = 12.1;
    double hostD = 12.1;

    CHECK(cudaMalloc((void **)&deviceF, sizeof(float)));
    CHECK(cudaMalloc((void **)&deviceD, sizeof(double)));
    kernel<<<1, 32>>>(deviceF, deviceD);
    CHECK(cudaMemcpy(&h_deviceF, deviceF, sizeof(float),
                     cudaMemcpyDeviceToHost));
    CHECK(cudaMemcpy(&h_deviceD, deviceD, sizeof(double),
                     cudaMemcpyDeviceToHost));

    printf("Host single-precision representation of 12.1   = %.20f\n", hostF);
    printf("Host double-precision representation of 12.1   = %.20f\n", hostD);
    printf("Device single-precision representation of 12.1 = %.20f\n", hostF);
    printf("Device double-precision representation of 12.1 = %.20f\n", hostD);
    printf("Device and host single-precision representation equal? %s\n",
           hostF == h_deviceF ? "yes" : "no");
    printf("Device and host double-precision representation equal? %s\n",
           hostD == h_deviceD ? "yes" : "no");

    return 0;
}

  虽然主机和设备上的数值都与12.1近似，但都不是精确值。在这个特殊的例子中，双精度数值比单精度数值稍微更接近于真实数值。
  双精度数值的精确性是以空间和性能消耗为代价的。来自于Wrox.com上floating-point-perf.cu程序随机产生了一个浮点输入向量，将该向量复制到GPU中，在GPU上重复执行大量的数学运算，然后再将结果复制回主机。使用单精度向量和双精度向量执行同样的操作，并对传输和内核所需时间进行测量。整个过程是反复运行的，以减少执行时间中随机变动造成的测量误差。代码如下所示：

#include <stdio.h>
#include <stdlib.h>

/**
 * The computational kernel for single-precision floating-point
 **/
__global__ void lots_of_float_compute(float *inputs, int N, size_t niters,
                                      float *outputs)
{
    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    size_t nthreads = gridDim.x * blockDim.x;

    for ( ; tid < N; tid += nthreads)
    {
        size_t iter;
        float val = inputs[tid];

        for (iter = 0; iter < niters; iter++)
        {
            val = (val + 5.0f) - 101.0f;
            val = (val / 3.0f) + 102.0f;
            val = (val + 1.07f) - 103.0f;
            val = (val / 1.037f) + 104.0f;
            val = (val + 3.00f) - 105.0f;
            val = (val / 0.22f) + 106.0f;
        }

        outputs[tid] = val;
    }
}

/**
 * The computational kernel for double-precision floating-point
 **/
__global__ void lots_of_double_compute(double *inputs, int N, size_t niters,
                                       double *outputs)
{
    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    size_t nthreads = gridDim.x * blockDim.x;

    for ( ; tid < N; tid += nthreads)
    {
        size_t iter;
        double val = inputs[tid];

        for (iter = 0; iter < niters; iter++)
        {
            val = (val + 5.0) - 101.0;
            val = (val / 3.0) + 102.0;
            val = (val + 1.07) - 103.0;
            val = (val / 1.037) + 104.0;
            val = (val + 3.00) - 105.0;
            val = (val / 0.22) + 106.0;
        }

        outputs[tid] = val;
    }
}

/**
 * Runs a full test of single-precision floating-point, including transferring
 * inputs to the device, running the single-precision kernel, and copying
 * outputs back.
 **/
static void run_float_test(size_t N, int niters, int blocksPerGrid,
                           int threadsPerBlock, double *toDeviceTime,
                           double *kernelTime, double *fromDeviceTime,
                           float *sample, int sampleLength)
{
    int i;
    float *h_floatInputs, *h_floatOutputs;
    float *d_floatInputs, *d_floatOutputs;

    h_floatInputs = (float *)malloc(sizeof(float) * N);
    h_floatOutputs = (float *)malloc(sizeof(float) * N);
    CHECK(cudaMalloc((void **)&d_floatInputs, sizeof(float) * N));
    CHECK(cudaMalloc((void **)&d_floatOutputs, sizeof(float) * N));

    for (i = 0; i < N; i++)
    {
        h_floatInputs[i] = (float)i;
    }

    double toDeviceStart = seconds();
    CHECK(cudaMemcpy(d_floatInputs, h_floatInputs, sizeof(float) * N,
                     cudaMemcpyHostToDevice));
    *toDeviceTime = seconds() - toDeviceStart;

    double kernelStart = seconds();
    lots_of_float_compute<<<blocksPerGrid, threadsPerBlock>>>(d_floatInputs,
            N, niters, d_floatOutputs);
    CHECK(cudaDeviceSynchronize());
    *kernelTime = seconds() - kernelStart;

    double fromDeviceStart = seconds();
    CHECK(cudaMemcpy(h_floatOutputs, d_floatOutputs, sizeof(float) * N,
                     cudaMemcpyDeviceToHost));
    *fromDeviceTime = seconds() - fromDeviceStart;

    for (i = 0; i < sampleLength; i++)
    {
        sample[i] = h_floatOutputs[i];
    }

    CHECK(cudaFree(d_floatInputs));
    CHECK(cudaFree(d_floatOutputs));
    free(h_floatInputs);
    free(h_floatOutputs);
}

/**
 * Runs a full test of double-precision floating-point, including transferring
 * inputs to the device, running the single-precision kernel, and copying
 * outputs back.
 **/
static void run_double_test(size_t N, int niters, int blocksPerGrid,
                            int threadsPerBlock, double *toDeviceTime,
                            double *kernelTime, double *fromDeviceTime,
                            double *sample, int sampleLength)
{
    int i;
    double *h_doubleInputs, *h_doubleOutputs;
    double *d_doubleInputs, *d_doubleOutputs;

    h_doubleInputs = (double *)malloc(sizeof(double) * N);
    h_doubleOutputs = (double *)malloc(sizeof(double) * N);
    CHECK(cudaMalloc((void **)&d_doubleInputs, sizeof(double) * N));
    CHECK(cudaMalloc((void **)&d_doubleOutputs, sizeof(double) * N));

    for (i = 0; i < N; i++)
    {
        h_doubleInputs[i] = (double)i;
    }

    double toDeviceStart = seconds();
    CHECK(cudaMemcpy(d_doubleInputs, h_doubleInputs, sizeof(double) * N,
                     cudaMemcpyHostToDevice));
    *toDeviceTime = seconds() - toDeviceStart;

    double kernelStart = seconds();
    lots_of_double_compute<<<blocksPerGrid, threadsPerBlock>>>(d_doubleInputs,
            N, niters, d_doubleOutputs);
    CHECK(cudaDeviceSynchronize());
    *kernelTime = seconds() - kernelStart;

    double fromDeviceStart = seconds();
    CHECK(cudaMemcpy(h_doubleOutputs, d_doubleOutputs, sizeof(double) * N,
                     cudaMemcpyDeviceToHost));
    *fromDeviceTime = seconds() - fromDeviceStart;

    for (i = 0; i < sampleLength; i++)
    {
        sample[i] = h_doubleOutputs[i];
    }

    CHECK(cudaFree(d_doubleInputs));
    CHECK(cudaFree(d_doubleOutputs));
    free(h_doubleInputs);
    free(h_doubleOutputs);
}

int main(int argc, char **argv)
{
    int i;
    double meanFloatToDeviceTime, meanFloatKernelTime, meanFloatFromDeviceTime;
    double meanDoubleToDeviceTime, meanDoubleKernelTime,
           meanDoubleFromDeviceTime;
    struct cudaDeviceProp deviceProperties;
    size_t totalMem, freeMem;
    float *floatSample;
    double *doubleSample;
    int sampleLength = 10;
    int nRuns = 5;
    int nKernelIters = 20;

    meanFloatToDeviceTime = meanFloatKernelTime = meanFloatFromDeviceTime = 0.0;
    meanDoubleToDeviceTime = meanDoubleKernelTime =
                                 meanDoubleFromDeviceTime = 0.0;

    CHECK(cudaMemGetInfo(&freeMem, &totalMem));
    CHECK(cudaGetDeviceProperties(&deviceProperties, 0));

    size_t N = (freeMem * 0.9 / 2) / sizeof(double);
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    if (blocksPerGrid > deviceProperties.maxGridSize[0])
    {
        blocksPerGrid = deviceProperties.maxGridSize[0];
    }

    printf("Running %d blocks with %d threads/block over %lu elements\n",
           blocksPerGrid, threadsPerBlock, N);

    floatSample = (float *)malloc(sizeof(float) * sampleLength);
    doubleSample = (double *)malloc(sizeof(double) * sampleLength);

    for (i = 0; i < nRuns; i++)
    {
        double toDeviceTime, kernelTime, fromDeviceTime;

        run_float_test(N, nKernelIters, blocksPerGrid, threadsPerBlock,
                       &toDeviceTime, &kernelTime, &fromDeviceTime,
                       floatSample, sampleLength);
        meanFloatToDeviceTime += toDeviceTime;
        meanFloatKernelTime += kernelTime;
        meanFloatFromDeviceTime += fromDeviceTime;

        run_double_test(N, nKernelIters, blocksPerGrid, threadsPerBlock,
                        &toDeviceTime, &kernelTime, &fromDeviceTime,
                        doubleSample, sampleLength);
        meanDoubleToDeviceTime += toDeviceTime;
        meanDoubleKernelTime += kernelTime;
        meanDoubleFromDeviceTime += fromDeviceTime;

        if (i == 0)
        {
            int j;
            printf("Input\tDiff Between Single- and Double-Precision\n");
            printf("------\t------\n");

            for (j = 0; j < sampleLength; j++)
            {
                printf("%d\t%.20e\n", j,
                       fabs(doubleSample[j] - (double)floatSample[j]));
            }

            printf("\n");
        }
    }

    meanFloatToDeviceTime /= nRuns;
    meanFloatKernelTime /= nRuns;
    meanFloatFromDeviceTime /= nRuns;
    meanDoubleToDeviceTime /= nRuns;
    meanDoubleKernelTime /= nRuns;
    meanDoubleFromDeviceTime /= nRuns;

    printf("For single-precision floating point, mean times for:\n");
    printf("  Copy to device:   %f s\n", meanFloatToDeviceTime);
    printf("  Kernel execution: %f s\n", meanFloatKernelTime);
    printf("  Copy from device: %f s\n", meanFloatFromDeviceTime);
    printf("For double-precision floating point, mean times for:\n");
    printf("  Copy to device:   %f s (%.2fx slower than single-precision)\n",
           meanDoubleToDeviceTime,
           meanDoubleToDeviceTime / meanFloatToDeviceTime);
    printf("  Kernel execution: %f s (%.2fx slower than single-precision)\n",
           meanDoubleKernelTime,
           meanDoubleKernelTime / meanFloatKernelTime);
    printf("  Copy from device: %f s (%.2fx slower than single-precision)\n",
           meanDoubleFromDeviceTime,
           meanDoubleFromDeviceTime / meanFloatFromDeviceTime);

    return 0;

  这个例子说明了两点。首先，单精度和双精度浮点运算在通信和计算上的性能差异是不可忽略的。在这种情况下，使用双精度数值能够使总的程序运算时间增加近一倍(虽然这个结果可能取决于应用程序是计算密集型还是I/O密集型)。在设备端进行进行数据通信的时间也是使用单精度数值的两倍，这是由双精度数值长度是单精度数值长度的二倍造成的。随着全局内存输入/输出数量和每条指令执行的位操作数量的增加，设备上的计算时间也会增加。这个程序也说明了单精度与双精度的结果有较大的数值差异，这些结果可能在迭代过程中不断被积累，即第一次迭代产生的不精确的结果作为下一次迭代的输入继续参与运算，导致最后的结果偏差很大。因此，考虑到数值精确度，在迭代应用中可能更需要使用双精度变量。
  还需要注意的是，由于双精度数值所占空间是单精度数值的两倍，所以当在寄存器中存储一个双精度数值(在内核中已被声明)时，一个线程块总的共享寄存器区域会比使用浮点数小得多。在声明单精度浮点数值时必须非常谨慎(例如，pi = 3.14159f;)任何不正确的省略尾数f的声明(pi = 3.14159)都会自动地被NVCC编译器转换成双精度数。
  浮点运算对应用程序地性能和数值精确度上的影响并不只是在GPU上才会产生，使用其他架构时，会面对同样的问题。以下是CUDA和GPU独有的特点：1.使用双精度数值增加主机和设备之间的通信；2.使用双精度数值增加全局内存的输入/输出；3.数值精度的损失是由CUDA编译器强制浮点数值优化导致的。一般情况下，如果应用程序精确度要求很高的话，必须使用双精度数值。否则，使用单精度数值可以获得性能提升。下面总结了CUDA中使用浮点数运算的一些经验：

标 准 函 数 与 内 部 函 数 的 比 较
  标准函数和内部函数在数值精确度和性能上的表现是不同的。标准函数支持大部分数学运算。但是，许多等效的内部函数能够使用较少的指令、改进的性能和更低的数值精确度，实现相同的功能。

标准函数和内部函数可视化
  通过学习由CUDA编译器产生的针对每个函数的指令，可以将标准函数和内部函数差异可视化。使用nvcc的–ptr标志能够让编译器在并行线程执行(PTX)和指令集架构(ISA)中生成程序的中间表达式，而不是生成一个最终的可执行文件。PTX类似于x86编程里面的程序集，它提供了一个我们所编写的内核代码之间的中间表达式，该指令在GPU上执行。因此，它对于深入了解内核的低级别执行路径是很有用的。
  例如，我们可以为以下两个CUDA函数生成一个PTX来直观地比较标准函数和内部函数。为此，将这些函数存储到一个命名为foo.cu文件中：

__global__ void intrinsic(float *ptr)
{
   *ptr = __powf(*ptr,2.0f);
}

__global__ void standard(float *ptr){
   *ptr = powf(*ptr,2.0f);
}

  接下来使用以下命令生成一个PTX文件并命名为foo.ptx:$ nvcc --ptx -o foo.ptx foo.cu,nvcc编译器会为这些设备函数生成一个包含PTX指令的文件，可以用文本编辑器打开。
  内部函数_powf实现需要17行代码，并且只有7条指令执行浮点数运算。标准函数powf实现的代码要长的多,这些代码行数并不直接转化为指令或者循环，所以性能上的区别仍很重要。
  然而，区分标准函数和内部函数的不仅有性能，它们的计算精度也是被不同的。为了测试性能和精确度的不同，可以从Wrox.com中下载Intrinsic-standard-comp.cu这个例子，创建并运行相关应用程序。在该程序中的核函数中，先使用标准函数powf，再使用内部函数__powf,利用它们反复计算输入值的平方根。这个例子也使用主机上的C标准数学库来执行相同的计算，并使用主机上的结果作为基准值。intrinsic-standard-comp.cu的示例输出如下所示：

  不出所料，使用内部函数相较于标准函数来说，速度提升了将近24倍，获得了巨大的性能提升。CUDA标准函数和内部蛤属不仅输出结果不同，它们与主机标准数学库计算的结果也不同。但是，当比较内部函数和标准函数的计算结果时，内在结果比主机结果相差一个数量级。

  使用CUDA来执行科学仿真、金融算法和其他要求高精度和高保真度的应用程序通常需要两个步骤：将传统应用从只有CPU的系统移植到CUDA系统中，接着通过比较传统应用结果与使用CUDA的执行结果，来验证程序移植的数值精确性。即使使用数值稳定的CUDA函数，GPU上的运算结果仍与传统的只在CPU上运行的应用结果不同。由于主机和设备上的浮点运算都存在固有的不精确性，有时很难指出一个输出结果与另一个输出结果哪个更精确。因此，必须考虑数值差异并作出恰当的移植计划，而且有必要的话需要设置允许的误差范围。

操纵指令生成
  在大多数情况下，将程序员编写的内核代码转换为GPU指令集这一过程是由CUDA编程器完成的。程序员很少会有检查或手动修改指令的想法。但是，这并不意味着我们无法引导编译器倾向于实现良好的性能或准确性或者达到两者的平衡。CUDA编译器中有两种方法可以控制指令级优化类型：编译器指令、内部或标准函数调用。
  例如，内部函数__fdividef与运算符“/”相比，在执行浮点数除法时速度更快，但数值精确度相对较低。一个个手动调整内核操作的工作量太大了。编译器标志提供了一个更自动、全局化的方式来操纵编译器指令的生成。例如，我们可能想要通过CUDA编译器控制浮点数MAD(FMAD)指令的生成。MAD是一个简单的编译器优化指令，它能将乘法和加法融合到一个指令中，从而使运算时间比使用两个指令缩短一半。但是，这个优化需要以数值精度为代价。所以，一些应用程序会明确限制FMAD指令的使用。nvcc的–fmad选项可全局性地启用或禁止FMAD整个编译单元的优化。默认情况下，nvcc使用“–fmad=true”以启用FMAD指令来优化性能。"–fmad=false"的意思是阻止编译器混合任何乘法和加法，这虽然有损性能但可能提高应用程序的数值精度。注意，除了–fmad，还有许多CUDA编译器指令标志会影响算法指令的生成。完整的列表可在nvcc的–help选项中找到。下表中列出了这些编译器标志：

  除了–fmad选项，CUDA还包含一对用于控制FMAD指令生成的内部函数：__fmul和__dmul,这些函数用于实现单精度浮点型和双精度浮点型乘法。然而这些函数不会影响乘法运算的性能，在有“*”运算符的地方调用它们可阻止nvcc将乘法作为乘加优化的一部分来使用。需要注意的是，不论是指定–fmad=true还是–fmad=false,__fmul和__dmul都阻止MAD指令的生成。因此，当通过有选择地调用__fmul或者__dmul的计算来提升某些数值的健壮性时，可启用MAD编译器优化全局。
  在调用__fmul时，实际上调用的是一个__fmul_rn，许多浮点型内部函数(包括__fadd,__fsub,__fmul等)在函数名中都使用两个后缀字符，这明确指出了浮点四舍五入的模式。回想一下，由于浮点变量只能表示离散的细粒度值，任何不能表示的值必须倍舍入为可表示的值。浮点运算的舍入模式决定了如何将不可表示的值转化为可表示的值。不同的四舍五入模式下__fmul的变体如下表所示：

总结

了 解 原 子 指 令
从头开始
  通过使用一个原子函数，每个由CUDA提供的原子函数可以重复被执行：原子级比较并交换(CAS)运算符。原子级CAS是一个很重要的操作，不仅可以使你在CUDA中定义你自己的原子函数，还能帮助你更深层次理解原子操作。CAS将3个内容作为输入：内存地址、存储在此地址中的期望值，以及实际想要存储在此位置的新值，然后执行以下几步：1.读取目标地址并将该处地址的存储值与预期值进行比较(a.如果存储值与预期值相等，那么新值将存入目标位置；b.如果存储值与预期值不等，那么目标位置不会发生变化。)；2.不管发生什么情况，一个CAS操作总是返回目标地址中的值。注意，使用返回值可以用来检查一个数值是否被替换成功。如果返回值等于传入的预期值，那么CAS操作一定成功了。
  这只是CAS操作。一个原子CAS意味着整个CAS进程是在没有其他任何线程的干扰下完成的。因为这是一个原子操作，如果CAS操作返回值显示写操作成功，那么所执行的数值交换必须对其他所有线程也可见。
  想要学习更多关于原子操作的知识，可以使用CUDA的atomicCAS设备函数从头开始去实现一个原子函数。在这个例子中，你可以进行原子级32位整型加法运算。atomicCAS相关变体的函数签名为：int atomicCAS(int *address, int compare, int val);其中“address”是目标内存地址，"compare"是预期值，"val"是实际想写入的新值。定义成CAS操作。当执行自定义原子操作时，定义目标的起始和结束状态是很有帮助的。在原子加法中，起始状态时递增运算的基值。结束值是起始状态和增值的总和。这个定义直接转换为atomicCAS:预期值是起始状态，实际写入的新值是完成状态。
  若想实现一个自定义的原子加法函数，需要从函数签名开始，它需要一个目的地址存储到该地址的值。__device__ int myAtomicAdd(int *address, int incr){ ... },可以通过读取目标内存的地址，计算出存放在目标地址的预期值。将读取到的值以及传递给myAtomicAdd的incr值定义实际值。使用这些预期值和实际值，可以调用atomicCAS来实现加法运算：

__device__ int myAtomicAdd(int *address, int incr)
{
   int expected = *address;
   int oldValue = atomicCAS(address,expected,expected + incr);
}

  这个myAtomicAdd函数可以实现原子加法。但是只有当执行atomicCAS，读入“expected”的值与存入“address”的值相同时操作才成功。因为目标位置是由多线程共享的(否则不需要原子操作)，所以另一个线程修改“address”的值是有可能的，这个值处于被“expected”读入和atomicCAS修改之间。如果发生这种情况，atomicCAS的执行会因在"address"中的值和“expected”中的值不同而失败。如果“atomicCAS”的返回值与预期值不同则程序会失败。因此，myAtomicAdd可以用来检查失败并在一个循环中重试CAS直到atomicCAS成功。

__device__ int myAtomicAdd(int *address, int incr)
{
   int expected = *address;
   int oldValue = atomicCAS(address,expected,expected + incr);

   while(oldValue != expected)
   {
      expected = oldValue;
      oldValue = atomicCAS(address,expected,expected + incr);
   }
   return oldValue;
}

  该函数的前三行和之前的相同，如果第一个atomicCAS失败了，那么myAtomicAdd就会循环执行直到atomicCAS最后的返回值与预期值不同。一旦条件失败，交换必定已经成功，并且myAtomicAdd退出循环。另一方面，预期值重置为最近读取的值并重试。为了匹配其他CUDA原子函数的语义，通过atomicCAS最近的返回值，myAtomicAdd也返回目标地址中的数值。
内置的CUDA原子函数
  下表列出了CUDA支持的原子函数的原子操作，包括相关的CUDA设备函数和支持的数值类型：

原子操作的成本
  原子函数在一些应用中很有帮助且很有必要，但可能要付出很高的性能代价。导致这种局面的原因有如下几个方面：
  1.当在全局或共享内存中执行原子操作时，能保证所有的数值变化对所有线程都是立即可见的。因此，在最低限度下，一个原子操作指令将通过任何方式进入到全局或共享内存中读取当前存储的数值而不需要缓存。如果原子指令，那么必须把实际需要的值写入到全局或共享内存中。
  2.共享地址冲突的原子访问可能要求发生冲突的线程不断地进行重试，类似于运行多个myAtomicAdd循环的迭代。尽管内置原子函数建立过程的可见性是有限的，但对你所实现的任何自定义原子操作来说都是真实的。如果你的应用程序反复循环而致使I/O开销较大，相应地性能会降低。
  3.当线程在同一个线程束时必须执行不同的指令，线程束执行是序列化的。如果一个线程束中的多个线程在相同的内存地址发出一个原子操作，就会产生类似于线程冲突的问题。因为只有一个线程的原子操作可以成功，所以所有其他的线程必须重试。如果一个原子指令需要n个循环，并且需要同一线程束中的t个线程在相同的内存地址上执行该原子指令，那么运行的时间将会是t×n，因为每次重试时只有一个线程会成功。记住，线程束中剩下的那些线程会等待所有原子操作的完成，并且一个原子操作也意味着一个全局的读取和写入。

  当原子操作是必要的而不安全访问是一个选择项时，这将在很大程度上降低性能和正确性。当做这个决定时必须非常小心，并不推荐使用不安全访问，应当只有在能保证正确性的情况下才尝试使用不安全访问。
限制原子操作的性能成本
  幸运的是，当必须执行原子操作时，使用有些方法可以减少性能损失。我们可以使用局部操作来增强全局原子操作，这些局部操作能从同一线程块的线程中产生一个中间结果。这需要使用本地较低延迟的资源，如shuffle指令或共享内存，在使用原子操作把局部结果结合到最终全局结果之前，需要先从每个线程块产生局部结果。当然，为使其有效，这些操作必须是可替换的(也就是操作的顺序不能影响最后的结果)。下图展示了局部还原产生部分结果，然后是原子操作去计算最终的输出。

原子级浮点支持
  原子函数中要注意的一点是它们大多被声明在整型数值上操作，如int、unsigned int 或unsigned long long int。纵观所有原子函数，只有atomicExch和atomicAdd支持单精度浮点数。所有原子函数都不支持双精度数值的运算。所幸，如果我们的应用程序需要管理多个线程访问共享浮点变量，那么有些方法可以让我们实现自己的浮点原子操作。在高级别上，有一个方法是用一个变量中支持的类型存储浮点数的原始比特位，并使用所支持的类型执行原子CAS操作。

__device__ float myAtomicAdd(float *address, float incr)
{
   unsigned int *typedAddress = (unsigned int *)address;
   
   float currentVal = *address;
   unsigned int expected = __float2uint_rn(currentVal);
   unsigned int desired = __float2uint_rn(currentVal + incr);
    
   int oldIntValue = atomicCAS(typeAddress,expected,desired);
   while(oldIntValue != expected)
   {
      expected = oldIntValue;
      
      desired = __float2uint_rn(__uint2float_rn(oldIntValue) + incr);
      oldIntValue = atomicCAS(typeAddress,expected,desired);
   }
   return __unit2float_rn(oldIntValue);
}

  这段代码中的大部分内容与之前的myAtomicAdd示例类似。主要不同的是atomicCAS数值转换的传入和传出，这个过程使用的是CUDA提供的各种类型的转换函数。这个特例使用了：
  1.一个cast改变了address指针的类型，使其从float型转换为unsigned int型。
  2.使用__float2uint_rn将期望值、address、期望值以及address+incr的类型转换为包含相同比特位的unsigned int 类型。
  3.如果操作失败了，使用__unit2float_rn检索一个从atomicCAS返回的unsigned int 浮点数并计算新的期望值。
  所有这些类型转换都是必要的，因为应用程序要求的类型(float)和atomicCAS函数要求的类型(unsigned int)不同的。CUDA提供了一个有很大范围的有其他特定类型转换的函数，包括__double_as_longlong、__longlong_as_double、__double-2float_rn等。这些函数对实现CUDA中大范围的自定义浮点原子函数是很有用的。完整的列表可以在CUDA MATH API文档中找到。
总结