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的设备相同。