CUDA學習（八十八）

3.雖然<code>__syncthreads（）</code>一直被記錄為同步線程塊中的所有線程，但Pascal和以前的體系結構隻能在warp級别強制執行同步。 在某些情況下，隻要每條經線中至少有一條線達到屏障，就可以在不被每條線執行的情況下成功實作屏障。 從Volta開始，CUDA内置的<code>__syncthreads（）</code>和PTX指令bar.sync（及其衍生物）是針對每個線程實施的，是以在該塊中的所有非退出線程到達之前都不會成功。 利用先前行為的代碼可能會死鎖，并且必須進行修改，以確定所有非退出的線程都到達障礙。

cuda-memcheck提供的racecheck和synccheck工具可以幫助定位第2點和第3點。

為了在實施上述糾正措施時幫助遷移，開發人員可以選擇不支援獨立線程排程的Pascal排程模型。 詳情請參閱應用程式相容性 。

全局記憶體：

全局記憶體的運作方式與計算能力5.x的裝置相同。

共享記憶體：

為共享記憶體保留的128 KB資料高速緩存的比率稱為共享記憶體分塊。 可以在每個核心的基礎上配置分區。 支援的共享記憶體容量為0,8,16,32,64或96 KB。

不是擴充用于Kepler架構的現有cudaFuncSetCacheConfig（）API以支援一組擴充的共享容量，而是設計了一個新的運作時API。 新的API解決了舊API中的兩個關鍵缺陷。 首先，為每個容量比定義單獨的枚舉不會優雅地适用于許多選項。 其次，因為傳統API将共享記憶體容量視為核心啟動的硬性要求，是以将核心與不同的共享記憶體請求交錯将不必要地在共享記憶體重新配置後序列化啟動。

新API使用函數cudaFuncSetAttribute（），如下所示。

在這裡，整數分割指定共享記憶體分割首選項占總資源的百分比。 這隻是一個提示，如果需要執行該功能，驅動可以選擇不同的比例。

計算能力7.0裝置允許單個線程塊尋址完整的96 KB共享記憶體。 依賴每塊超過48 KB的共享記憶體配置設定的核心是特定于體系結構的，是以它們必須使用動态共享記憶體（而不是靜态大小的數組），并且需要使用cudaFuncSetAttribute（）進行顯式選擇，如下所示：

其他共享記憶體的行為與計算能力5.x的裝置相同。