CUDA02中如何实现访存优化及Unified Memory的优化策略？

本文共计8060个文字，预计阅读时间需要33分钟。

CUDA02+ 的内存调度与优化（前篇：传输门）简要介绍了CUDA的底层架构和一些线程调度问题，但这只是CUDA学习的第一步。接下来，我们将探讨数据访问的问题，包括数据以何种形式存在以及如何高效访问。

CUDA02 - 的内存调度与优化

前面一篇（传送门）简单介绍了CUDA的底层架构和一些线程调度方面的问题，但这只是整个CUDA的第一步，下一个问题在于数据的访存：包括数据以何种形式在CPU/GPU之间进行通信、迁移，以及在GPU内部进行存储、访问。

1 global 、shared 、constant、local

通常来讲，待计算的数据都存放在内存或者硬盘（外部存储设备）中，由CPU来进行调度。想要在device上计算、处理数据，就首先需要将数据转移至CUDA，这样的转移操作通常需要经过数据总线实现。这跟叫做PCI-E的总线连接在cpu和gpu之间，负责信号传输，在早期的CUDA版本中程序员必须清楚的知道数据储存在了这根总线的哪一端。当然相比起CUDA的计算速度，总线的访存速度是很慢的。

上一章介绍了CUDA的底层存储结构。在G80中，一个核心计算单元通过访问不同等级的存储设备，来获取计算资源。这些资源有些是属于线程的，有些是属于SM的，还有一些是全局的。下面写一些这些物理结构对应的软件结构，分成了以下几种：

device shared

以__device__ __shared__为关键词声明的变量会被分配至SM上的shared memory， 可以由block内的全部线程所共享，生命周期也随着block的结束而结束。