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）。