来自 CUDA 编程指南(v. 5.5)：
CUDA 编程模型假设设备具有弱序 内存模型，即：
- CUDA 线程将数据写入共享内存、全局内存、页面锁定主机内存或对等设备内存的顺序 不一定是观察数据的顺序 由另一个 CUDA 或主机线程编写;
- CUDA 线程从共享内存、全局内存、页面锁定主机内存或对等设备的内存读取数据的顺序 不一定是读取指令出现的顺序 用于彼此独立的指令的程序
但是，我们是否可以保证从单个线程中看到的(依赖)内存操作实际上是一致的？如果我这样做 - 说：
arr[x] = 1;
int z = arr[y];
如果x恰好等于y，并且没有其他线程接触内存，我能保证z是 1 吗？还是我仍然需要在这两个操作之间设置一些volatile或障碍？
回应奥佩多的回答。
如果您的编译器没有将代码中所述的功能编译为机器代码中的相同功能，则编译器要么损坏，要么您没有考虑优化......
我的问题是允许哪些优化(由编译器或硬件完成)？ 例如---store指令是非阻塞的，并且不知何故，内存控制器管理load指令的速度比已经排队的store更快，---可能会发生这种情况。
我不知道 CUDA 硬件。我能保证上述情况永远不会发生吗？
CUDA 编程指南简单地说，您无法预测线程的执行顺序，但每个线程仍将作为顺序线程运行。 在您陈述的示例中，x 和 y 相同并且没有其他线程接触内存，您确实可以保证 z = 1。 这里的重点是，如果你有几个线程对相同的数据(例如一个数组)执行操作，你不能保证线程 #9 在 #10 之前执行。
举个例子：
__device__ void sum_all(float *x, float *result, int size N){
x[threadId.x] = threadId.x;
result[threadId.x] = 0;
for(int i = 0; i < N; i++)
result[threadId.x] += x[threadID.x];
}
这里我们有一些愚蠢的函数，它应该用 m 中的数字填充共享数组 (x) ...n(从一个数字读到另一个数字)，然后将已经放入数组中的数字相加并将结果存储在另一个数组中。 假设您的最低索引线程是枚举线程 #0，您会期望代码第一次运行时此代码 x 应包含
x[] = {0， 0， 0 ...0} 和结果 [] = {0， 0， 0 ...0}
接下来是线程 #1
x[] = {0， 1， 0 ...0} 和结果 [] = {0， 1， 0 ...0}
接下来是线程 #2
x[] = {0， 1， 2 ...0} 和结果 [] = {0， 1， 3 ...0}
等等。 但这并不能保证。例如，您无法知道线程 #3 是否先运行，因此在线程 #0 运行之前更改数组 x[]。实际上，您甚至不知道在执行代码时数组是否被其他线程更改。
我不确定，这是否在 CUDA 文档中明确说明(我不希望如此)，因为这是计算的基本原理。基本上你要问的是，在GFX上运行代码是否会改变代码的功能。
GPU 的核心通常与 CPU 的内核相同，只是控制算术较少，指令集较小，并且通常仅支持单精度。 在 CUDA-GPU 中，每个 Warp 有 1 个程序计数器(32 个同步内核的部分)。与 CPU 一样，程序计数器在每条指令后增加一个地址元素的数量级，除非您有分支或跳转。这给出了程序的顺序流，这是无法更改的。 分支和跳转只能由内核上运行的软件引入，因此由编译器确定。编译器优化实际上可以更改代码的功能，但仅在代码相对于编译器实现"错误"的情况下 简而言之 - 您的代码将始终按照它在内存中的顺序执行，无论它是在 CPU 还是 GPU 上执行。如果您的编译器没有将代码中所述的功能编译为机器代码中的相同功能，则编译器要么损坏，要么您没有考虑优化......
希望这足够清楚:)
据我了解，您基本上是在询问 CUDA 编译器中是否尊重内存依赖关系和别名分析信息。
这个问题的答案是，假设 CUDA 编译器没有错误，是的，因为正如 Robert 指出的那样，CUDA 编译器在引擎盖下使用 LLVM 和两个基本模块(目前，我真的不认为它们可以被管道排除)是：
- 内存依赖性分析
- 别名分析
这两个传递检测可能指向同一地址的内存位置，并对变量使用实时分析(即使在块范围之外)以避免危险的优化(例如，您无法在下次读取之前写入实时变量，数据可能仍然有用)。
我不知道编译器的内部结构，但假设(与任何其他合理信任的编译器一样)它将尽最大努力实现无错误，其中发生的分析真的不应该打扰您，并向您保证至少在理论上您刚刚作为示例呈现的内容(即依赖加载比存储更快)不会发生。
什么保证你这样做？无非是该公司正在提供编译器以供使用的事实，并且有免责声明，以防万一不是特殊情况:)
另外：除了编译器主题之外，指令执行还取决于硬件规范。在这种情况下，SIMT 硬件指令发出单元 http://www.csl.cornell.edu/~cbatten/pdfs/kim-simt-vstruct-isca2013.pdf 和所有参考论文以获取更多信息