我看到大约有10个问题，这使得回答起来相当困难。建议您每次只提一个问题。
一般而言，warp（线程束）中的所有线程都在执行相同的指令流。因此，我们可以考虑两种情况：
1. 没有条件语句（例如if...then...else）。在这种情况下，所有线程都执行相同的指令，这恰好是原子操作。然后，所有32个线程都将执行一个原子操作，尽管不一定在相同的位置上。所有这些原子操作都将由SM处理，并且在某种程度上会序列化（如果它们更新相同的位置，则会完全序列化）。
2. 有条件语句。例如，假设我们有if (!threadIdx.x) AtomicAdd(*data, 1);那么线程0将执行原子操作，其他线程则不会。可能看起来我们可以让其他线程执行其他操作，但锁步warp执行不允许这样做。Warp执行被序列化，以便所有走“if(true)”路径的线程将一起执行，所有执行“if(false)”路径的线程也将一起执行，但真和假的路径将被序列化。因此，在一个warp中，我们实际上不能让不同的线程同时执行不同的指令。
总之，在一个warp中，我们不能让一个线程执行原子操作，而其他线程同时执行其他操作。
您的其他一些问题似乎期望内存事务在它们发起的指令周期结束时完成。这并不是事实。对于全局和共享内存，我们必须在代码中采取特殊步骤，以确保先前的写入事务对其他线程可见（这可以被认为是事务已经完成的证据）。一种典型的方法是使用屏障指令，例如__syncthreads()或__threadfence()。但是，如果没有这些屏障指令，线程将不会“等待”写入完成。读取可能会阻塞线程，而写入通常不会阻塞线程。
现在让我们看看您的问题：
如果其中一个线程开始原子操作，其他线程会怎么做？等待吗？
不，它们不会等待。原子操作被分派到处理原子操作的SM上的一个功能单元中，并且所有线程一起锁定。由于原子通常意味着读取，因此读取可能会阻塞warp。但是，线程不会等待原子操作完成（即写入）。但是，对该位置的后续读取可能会再次阻塞warp，等待原子操作（写入）完成。在全局原子操作的情况下，它将使原始SM中的L1（如果启用）以及L2失效，如果它们包含该位置作为条目。
“有没有办法让其他线程做一些工作，比如读取一些内存，以便原子操作可以隐藏其延迟？”
实际上没有，出于我在开头所述原因。
“原子操作真的很慢吗？我该如何加速它？有使用它的模式吗？”
是的，如果原子操作占主导地位（例如，天真的约简或天真的直方图），那么原子操作可能会使程序运行得更慢。通常，加速原子操作的方法是不使用它们或谨慎使用它们，以不支配程序活动的方式。例如，天真的约简将使用原子操作将每个元素添加到全局总和中。聪明的并行约简将在线程块中的工作中根本不使用任何原子操作。在线程块约简完成后，可能使用单个原子将线程块部分总和更新到全局总和中。这意味着，我可以使用大约32个或更少的原子加法快速并行约简任意数量的元素。这种节省使用原子操作基本上在整个程序执行中不会被注意到，只是使并行约简能够在单个内核调用中完成而不是2次。
"共享内存：它们运行并行还是排队？"
它们将排队。原因是处理共享内存上的原子操作的功能单元数量有限，无法在单个周期中服务于warp的所有请求。
我避免回答与原子操作吞吐量有关的问题，因为据我所知，文档中没有明确规定此类数据。如果您发出足够同时或几乎同时的原子操作，则某些warp将在原子指令上停顿，因为馈送原子功能单元的队列已满。我不知道这是真的，也不能回答相关问题。
int x = *globalPtr; int y = kernelArg1 * kernelArg2; /*some other calculations*/ int z = x * 3;- 第二个：
- Nexenint y = kernelArg1 * kernelArg2; /*some other calculations*/ int z = *globalPtr * 3;第一个更可取，对吧？