0

最近我发布了这个问题，关于一个关键部分。这是一个类似的问题。在这些问题中，给出的答案说，代码是否“有效”取决于编译器，因为各种执行路径的顺序取决于编译器。

为了详细说明问题的其余部分，我需要从CUDA 编程指南中摘录以下内容：

  1. ...组成warp的各个线程一起从同一个程序地址开始，但它们有自己的指令地址计数器和寄存器状态，因此可以自由地分支和独立执行......
  2. 一个 warp 一次执行一条公共指令，因此当一个 warp 的所有 32 个线程都同意它们的执行路径时，就可以实现完全的效率。如果 warp 的线程通过依赖于数据的条件分支发散，则 warp 会串行执行所采用的每个分支路径，禁用不在该路径上的线程，当所有路径完成时，线程会收敛回相同的执行路径...... .
  3. 多处理器处理的每个 warp 的执行上下文（程序计数器、寄存器等）在 warp 的整个生命周期内都在芯片上维护。因此，从一个执行上下文切换到另一个执行上下文是没有成本的，并且在每个指令发出时间，warp 调度程序都会选择一个线程准备好执行其下一条指令（warp 的活动线程）并将指令发布给这些线程.

我从这三个摘录中了解到，线程可以与其他线程自由分歧，如果线程之间存在分歧，所有分支的可能性都将被序列化，如果采用分支，它将执行直到完成。这就是为什么上面提到的问题以死锁结束的原因，因为编译器强加的执行路径的顺序导致采用没有获得锁的分支。

现在的问题是：编译器不应该总是按照用户编写的顺序放置分支？有没有一种高级方法来强制执行顺序？我知道，编译器可以优化，对指令进行重新排序等，但它不应该从根本上改变代码的逻辑（是的，有一些例外，比如没有 volatile 关键字的内存访问，但这就是关键字存在的原因，将控制权交给用户）。

编辑

这个问题的重点不是关于关键部分，而是关于编译器，例如在第一个链接中，编译标志彻底改变了代码的逻辑。一个“工作”，另一个没有。困扰我的是，在所有参考资料中，它只说要小心，没有关于 nvcc 编译器的未定义行为。

4

1 回答 1

1

我相信 CUDA 编译器没有设置或保证执行顺序。这是设置它的硬件——据我所知。

因此，

编译器不应该总是按照用户编写的顺序放置分支？

无论如何它不控制执行顺序

是否有执行命令的高级方法？

只是同步指令，如__syncthreads().

编译器......不应该从根本上改变代码的逻辑

CUDA 代码的语义与 C++ 代码的语义不同...... if 分支的顺序执行不是语义的一部分。

我意识到这个答案可能会让你不满意，但这就是事情的现状，无论好坏。

于 2019-01-09T14:29:12.460 回答