几乎在任何地方我了解与CUDA编程存在的所有的经线线程做同样的事情的重要性一提。
在我的代码我有一种情况，我无法避免一定的条件。 它看起来像这样：
// some math code, calculating d1, d2
if (d1 < 0.5)
{
buffer[x1] += 1; // buffer is in the global memory
}
if (d2 < 0.5)
{
buffer[x2] += 1;
}
// some more math code.
一些线程可能会进入一个规定的条件，有些人可能会进入到这两个和其他可能无法进入任一。
现在，为了使条件后，所有的线程回到“做同样的事情”了，我应该使用条件后，它们同步__syncthreads() 或者，这是否在某种程度上自动地发生？
两个线程可以不这样做，由于他们是背后一个操作的一个同样的事情，从而破坏给大家？ 或者是有一些幕后的工作，让他们一个分支后，再次做同样的事情？
在经纱，没有线程就会“出人头地”的任何其他人。 如果有一个条件分支，它是由在经一些线程而不是其他拍摄（又名经“三岔口”），其他线程只会闲置，直到分公司完成，他们都“收敛”重新走到一起在一个共同的指令。 所以，如果你只需要线程内，经同步，即发生“自动地”。
但是，不同的经纱不同步这种方式。 所以，如果你的算法需要某些操作是在许多经线完整的，那么你就需要使用显式同步调用（见CUDA编程指南，第5.4节）。
编辑：重组未来数段，以澄清一些事情。
有真的在这里两个不同的问题：指令同步和存储可见性。
__syncthreads()强制执行指令的同步，确保存储器的可见性，但只能在一个块，而不是跨越块（CUDA编程指南附录B.6）。 它是写，然后读取共享内存有用的，但并不适用于全球同步的存储器访问。
__threadfence()确保全局内存的知名度，但没有做任何指令同步，所以在我的经验是有限的使用（但见附录B.5示例代码）。
全局指令同步是不可能的内核中。 如果你需要f()所有线程完成之前调用g()在任何线程，分割f()和g()分成两个不同的内核，并从主机顺序调用它们。
如果你只需要增加共享或全局计数器，可以考虑使用单位递增函数atomicInc()附录B.10）。 在你的代码的情况下上述情况，如果x1和x2是不是（在网格跨所有的线程）全球唯一的，非原子增量将导致竞争条件，类似于附录B.2.4的最后一段。
最后，请记住，特别是（包括原子能）上的全局内存的任何操作，并同步功能不好的表现。
不知道这个问题你解决这是很难猜测，但也许你可以重新设计你的算法使用共享内存而非全局内存在一些地方。 这将减少同步的需求，并给你的性能提升。
从CUDA最佳实践指南的第6.1节：
任何流控制指令（如果，开关，做，因为，而）可以通过使相同的经纱发散的线程显著影响指令吞吐量; 也就是说，按照不同的执行路径。 如果发生这种情况，不同的执行路径必须是串行化的，因而增加了此经执行的指令的总数量。 当所有不同的执行路径都已完成时，线程将重新汇聚到同一执行路径。
所以，你不需要做什么特别的事情。
在加布里埃尔的回应：
“环球指令同步是不可能的内核之内。如果您需要在调用克（）在任何线程，分割F（）和G（）为两个不同的内核，并从主机连续叫他们前f（）的所有线程完成。 “
如果你需要在同一个线程F（）和g（）的原因是因为你使用的寄存器内存，你想从F寄存器或共享数据让到g？ 也就是说，对于我的问题，整个原因，跨块同步是因为需要以g是F的数据 - 并打破了一个内核将需要大量的额外全球内存是F将寄存器数据克，这是我倒是希望避免
在回答你的问题是没有。 你不需要做什么特别的。 无论如何，你可以解决这个问题，而不是你的代码，你可以这样做：
buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);
您应该检查是否可以使用共享内存和聚结的方式访问全局存储器。 另外要确保你不想写相同的指数超过1个线程。
