5. 得到正解
得到正确的计算结果是我们的最初目的，但是使用CUDA这种并行编程模型是很容易出错的，这时我们就需要一些方法和工具来帮助我们验证计算结果的正确性，同时，在CUDA编程中也有一些值得我们注意的问题。
5.1 正确性验证
正确性验证主要有两种方法：
1、引用比较（Reference Comparison）
引用比较的核心思想是使用未并行化代码产生的一些具有代表性的结果与并行化后的程序运行结果，当它们的绝对差在可接受范围内时，就认为并行化的结果是正确的。注意，改写前和改写后代码运行结果的不一致是由浮点数表示的不确切性造成的。
第一步完成后，我们使用APOD（Assess、Parallelize、Optimize、Deployment）过程对并行化代码实施进一步优化，我们只要保证每一步优化的引用比较结果正确，那么对于最终的并行化程序，其结果的正确性是可以得到保证的。
2、单元测试（Unit Testing）
单元测试与引用比较的方法是相辅相成的。单元测试是指开发人员在编写代码时就将项目代码组织成单元级别，然后运用一定的技术手段对各个单元分别测试其正确性。在CUDA中，我们可以把内核（kernels）写成一系列小的__device__函数的组合而不是将代码封装到一个庞大的__global__函数中。这样，我们就可以在连接各部分代码之前对各个设备（device）代码进行单独测试。
5.2 调试（Debugging）
调试CUDA需要一些特别的工具：
1. CUDA-GDB
CUDA-GDB是Linux和Mac环境中GNU调试器的一个端口（Port），具体信息参见：GNU-GDB
2. NVIDIA Parallel Nsight
NVIDIA Parallel Nsight调试和性能分析器是可以在Windows环境下作为Microsoft Visual Studio的插件使用，具体参见：NVIDIA Parallel Nsight
3. 一些第三方调试工具
一些第三方工具也支持CUDA的调试，具体参见：Debugging solutions
5.3 数值精度问题
由于CUDA使用的是浮点数进行运算，涉及到一些精确度问题，需要我们在编程时注意一下。
1、单精度与双精度问题
单精度和双精度浮点运算的结果差别是很大的，在CUDA中，只有运算能力大于或等于1.3的硬件才能本地支持双精度运算。执行运算时，程序员务必搞清进行的是哪一种运算以获得正确的结果。nvcc编译命令行中使用-arch=sm_13可开启双精度运算。
2、浮点运算不遵从结合律
对于三个浮点数A、B以及C，需要注意的是，(A+B)+C并不等于A+(B+C)。
3、双精度扩展和单精度截断
由于单精度浮点数和双精度浮点数运算结果的不一致性，CUDA程序员应注意避免一些精度细节问题，如下面这段代码：
float a;
a = a * 1.02;
在C语言中，1.02会被隐含解释为double类型，那么第二个式子右边的a将会被扩展为双精度浮点数而执行double乘法，得到的结果为双精度浮点数，最后再将这一结果截断为单精度浮点数赋值给式子左边的a。这会带来隐患，解决这一问题的方法是用1.02f来表示单精度浮点常数。
6. 性能度量（Performance Metrics）
为了优化CUDA程序的性能，我们需要一种定量的方法来对程序性能进行衡量，同时，我们也需要明晰带宽（bandwidth）在性能衡量中所扮演的角色。下面我们将依依介绍这些概念。
6.1 定时
CUDA调用和内核执行可以使用CPU和GPU定时器来计时，使用这些计时器需要注意相关的问题：
1、使用CPU定时器
使用CPU定时器需要注意很多CUDA API函数是异步的，即这些函数在任务完成之前就会把控制权返回给CPU。它们一般都会在函数名之前带有Async，遇到这类函数时需要特别小心。为了达到精准定时的效果，我们需要同步CPU和GPU，这可以通过在开始和停止定时器前调用函数cudaDeviceSynchronize()来实现。cudaDeviceSynchronize()阻塞当前CPU线程，直到指定流（Stream）中所有之前的CUDA调用运行完毕之后CPU线程才开始继续执行。类似地，CPU线程也可以与GPU流或者GPU事件进行同步，使用cudaStreamSynchronize()函数或者cudaEventSynchronize()函数可以达到这样的效果。然而需要注意的是，上述两个函数不适合在默认流以外的其他流中使用，因为驱动程序可能安排不同流中的代码交错运行，这样就会造成计时错误。而默认流（Stream 0）在设备上表现为串行运行，因此可以使用上面两个函数进行计时。
2、使用CUDA GPU定时器
CUDA event API提供了用于创建和消除以及记录事件（使用timestamp）的调用函数，通过timestamp之间的差值换算，我们可以得到毫秒级的GPU运行时间信息。下面的代码片段展示了这一技术：
How to time code using CUDA events
cudaEvent_t start, stop;这里cudaEventRecord()用于将start和stop事件放置到默认流（stream 0），当设备运行到事件处时，将会为该事件记录一个时间戳（timestamp）。cudaEventElapsedTime()函数返回记录中start和stop事件之间的毫秒级时间差。由于记录是发生在GPU上的，其时钟是独立的，因此时钟分辨率也是操作系统无关的。GPU计时详细请参见：《CUDA Toolkit Reference Manual》.
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
kernel<<<grid, threads>>>(d_odata, d_idata, size_x, size_y, NUM_REPS);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);