1.4 并行执行和竞争条件
OpenACC并行化for循环（Fortran中是do循环），因此循环内的代码使用并发硬件执行线程并行执行。
循环内的变量i似乎是顺序递增的，但实际上在这个for循环中使用多个i变量的线程可能同时并行执行，这可能有点令人困惑。OpenACC不保证线程执行的顺序，注意这点非常重要。实际上，甚至不可能假设单调性。例如，很有可能第nCount―1次迭代实际上先于第0次迭代执行完。
OpenACC不保证线程执行的顺序，注意这点非常重要。
总之，OpenACC编程人员不能也不应该对线程执行的顺序做任何假设，如图1-15所示。
竞争条件
并行循环中线程之间的数据依赖性可能会给OpenACC编程人员带来问题，尤其是OpenACC除了原子操作外不提供任何锁机制来防止竞争条件。当多个线程相竞对一个共享数据项执行某个操作，这会导致竞争条件。除非在另一个线程可以访问这个共享数据项之前开始并且执行完操作，例如原子处理写或写后读操作，否则操作结果是未定义的。（注意：对只读数据项不存在竞争条件。）
如图1-16所示，以粗体显示的并行更新计数器发生共享数据依赖。当通过counter++更新时，示例使用#pragma acc atomic update来保护counter变量。在OpenACC中，使用原子读、写或者更新语法可以保护单个变量。示例使用条件编译，所以可以看到如果忽略原子操作结果会是什么。
图1-16 accCounter.cpp：包含竞争条件的示例
因为编程人员不能控制线程的执行顺序，因此不包含原子语法是一个常见的错误。这将产生竞争条件，从而导致垃圾和不确定性结果。非确定性行为是非常危险的，因为程序可能在调试时报告看起来是正确的结果，但是在产品中失败了。此外，应用程序可能在一个平台上正常执行，但在另一个平台上失败，或者使用特定软件版本时正确执行，然后软件更新后执行失败。
当遗漏原子操作时，PGI OpenACC编译器可以智能地检测示例中的错误。如图1-17中的粗体所示，编译器决定使用一个线程来生成串行代码，以确保正确性。下面的信息给出了原因，即变量counter在循环外生存。简单地说：变量保存了将来可能需要的值且编译器的数据流分析确定退出代码块时其他语句需要这个变量，这种情况下变量还生存着。精确的生存时数学定义超出了本章的讨论范围，感兴趣的读者可以阅读关于编译器数据流分析的文章以作了解。
定义USE_ATOMIC宏，在编译时包含OpenACC原子语法。在accCounter.cpp示例中，PGI编译器为GPU生成一个并行内核，为CPU生成多核循环。命令行参数nvidia:cc35告诉编译器目标设备为NVIDIA 3.5计算能力。这么做是为了看到更高效的现代GPU原子操作对运行时的影响，因为计算能力2.x（例如，费米）和更高版本的GPU具有更高效的原子操作。虽然原子操作方便，但还是要避免使用它们。因为每次只有一个线程可以访问原子变量，所以它们在运行时强制串行化。accCounter.cpp示例是一个糟糕的应用场景，因为每个线程必须排队来执行counter++的操作。
如图1-18所示，编译器报告生成了并行代码。由于对counter变量的竞争操作致使并行代码性能随着问题规模线性增长。如图1-19所示，相比串行代码，由于线程和原子操作的间接开销，并行代码运行时间更长。
如图1-20所示，在NVIDIA Tesla K40（计算能力3.5）GPU上串行版本运行时随着问题规模线性增长。与此同时，并行版本运行时比较稳定。这是因为NVIDIA对硬件做了优化，这些优化使得一些原子操作适用于大规模并行计算环境。更多信息，建议阅读《GPU大规模并行处理的大规模原子操作》（Egielski, Huang, &Zhang, 2014）和《NVIDIA费米：第一个完全GPU计算架构》（Glaskowshy, 2009）。
