CUDA 中的原子操作是否保证按 warp 进行调度？

Question

假设我在 GTX 970 上有 8 个块，每个块有 32 个线程运行。每个块将全 1 或全 0 写入全局内存中长度为 32 的数组，其中块中的线程 0 写入在数组中定位 0。

现在要写入实际值，使用 atomicExch，将数组中的当前值与块尝试写入的值交换。由于 SIMD、原子操作以及 warp 以锁步方式执行的事实，我希望数组在任何时间点都只包含 1 或 0。但绝不能将两者混为一谈。

但是，虽然运行代码是这样的，但在某些情况下，在某些时间点数组包含 0 和 1 的混合。这似乎表明原子操作不是按 warp 执行的，而是使用其他一些方案进行调度的。

从其他来源我还没有真正找到关于跨不同 warp 的原子操作调度的结论性文章（如果我错了请纠正我），所以我想知道是否有关于这个主题的任何信息.由于我需要将许多由几个32位整数组成的小向量原子写入全局内存，而保证原子写入单个向量的原子操作显然非常重要。

对于那些想知道的人，我编写的代码是在 GTX 970 上执行的，在计算能力 5.2 上编译，使用 CUDA 8.0。

Answer 1

像所有指令一样，原子指令按 warp 进行调度。然而，有一个与原子关联的未指定管道，并且不能保证通过管道的预定指令流对于每个线程，对于通过管道的每个阶段都以锁步方式执行。这为您的观察提供了可能性。

我相信一个简单的思想实验将证明这一定是真的：如果同一个 warp 中的 2 个线程针对同一个位置会怎样？显然，处理的每个方面都无法同步进行。我们可以将这个思想实验扩展到我们在一个 SM 内甚至跨 SM 每个时钟有多个问题的情况，作为其他示例。

如果向量长度足够短（16 字节或更少），那么应该可以通过让 warp 中的线程写入适当的向量类型数量来完成此操作 ("atomic update")，例如int4。只要所有线程（无论它们在网格中的哪个位置）都在尝试更新自然对齐的位置，写入就不应被其他写入破坏。

然而，经过评论中的讨论，OP 的目标似乎是能够让一个 warp 或 threadblock 更新一定长度的向量，而不受其他 warp 或 threadblock 的干扰。在我看来，真正需要的是访问控制（这样一次只有一个 warp 或 threadblock 更新特定向量）并且 OP 有一些代码没有按预期工作。

可以使用普通的原子操作（下例中的atomicCAS）强制执行此访问控制，一次只允许一个 "producer" 更新向量。

接下来是一个示例生产者-消费者代码，其中有多个线程块正在更新一系列向量。每个向量 "slot" 都有一个 "slot control" 变量，该变量自动更新以指示：

向量为空
向量正在填充
向量已填满，准备好"consumption"

使用这个 3 级方案，我们可以允许消费者和多个生产者工作人员使用单个普通原子变量访问机制对向量进行普通访问。这是一个示例代码：

#include <assert.h>
#include <iostream>
#include <stdio.h>

const int num_slots = 256;
const int slot_length = 32;
const int max_act = 65536;
const int slot_full = 2;
const int slot_filling = 1;
const int slot_empty = 0;
const int max_sm = 64;  // needs to be greater than the maximum number of SMs for any GPU that it will be run on
__device__ int slot_control[num_slots] = {0};
__device__ int slots[num_slots*slot_length];
__device__ int observations[max_sm] = {0}; // reported by consumer
__device__ int actives[max_sm] = {0};      // reported by producers
__device__ int correct = 0;
__device__ int block_id = 0;
__device__ volatile int restricted_sm = -1;
__device__ int num_act = 0;

static __device__ __inline__ int __mysmid(){
  int smid;
  asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
  return smid;}


// this code won't work on a GPU with a single SM!
__global__ void kernel(){

  __shared__ volatile int done, update, next_slot;
  int my_block_id = atomicAdd(&block_id, 1);
  int my_sm = __mysmid();
  if (my_block_id == 0){
    if (!threadIdx.x){
      restricted_sm = my_sm;
      __threadfence();
      // I am "block 0" and process the vectors, checking for coherency
      // "consumer"
      next_slot = 0;
      volatile int *vslot_control = slot_control;
      volatile int *vslots = slots;
      int scount = 0;
      while(scount < max_act){
        if (vslot_control[next_slot] == slot_full){
          scount++;
          int slot_val = vslots[next_slot*slot_length];
          for (int i = 1; i < slot_length; i++) if (slot_val != vslots[next_slot*slot_length+i]) { assert(0); /* badness - incoherence */}
          observations[slot_val]++;
          vslot_control[next_slot] = slot_empty;
          correct++;
          __threadfence();
          }
        next_slot++;
        if (next_slot >= num_slots) next_slot = 0;
        }
      }}
  else {
    // "producer"
    while (restricted_sm < 0);  // wait for signaling
    if (my_sm == restricted_sm) return;
    next_slot = 0;
    done = 0;
    __syncthreads();
    while (!done) {
      if (!threadIdx.x){
        while (atomicCAS(slot_control+next_slot, slot_empty,  slot_filling) > slot_empty) {
          next_slot++;
          if (next_slot >= num_slots) next_slot = 0;}
        // we grabbed an empty slot, fill it with my_sm
        if (atomicAdd(&num_act, 1) < max_act)   update = 1;
        else {done = 1; update = 0;}
        }
      __syncthreads();

      if (update) slots[next_slot*slot_length+threadIdx.x] = my_sm;
      __threadfence(); //enforce ordering
      if ((update) && (!threadIdx.x)){
        slot_control[next_slot] = 2; // mark slot full
        atomicAdd(actives+my_sm, 1);}
      __syncthreads();
    }
  }
}

int main(){

  kernel<<<256, slot_length>>>();
  cudaDeviceSynchronize();
  cudaError_t res= cudaGetLastError();
  if (res != cudaSuccess) printf("kernel failure: %d\n", (int)res);
  int *h_obs = new int[max_sm];
  int *h_act = new int[max_sm];
  int h_correct;
  cudaMemcpyFromSymbol(h_obs, observations, sizeof(int)*max_sm);
  cudaMemcpyFromSymbol(h_act, actives, sizeof(int)*max_sm);
  cudaMemcpyFromSymbol(&h_correct, correct, sizeof(int));
  int h_total_act = 0;
  int h_total_obs = 0;
  for (int i = 0; i < max_sm; i++){
    std::cout << h_act[i] << "," << h_obs[i] << " ";
    h_total_act += h_act[i];
    h_total_obs += h_obs[i];}
  std::cout << std::endl << h_total_act << "," << h_total_obs << "," << h_correct << std::endl;
}

我不声称此代码在任何用例中都没有缺陷。它用于演示概念的可行性，而不是作为生产就绪代码。它似乎适用于 linux，在我测试过的几个不同系统上。它不应该在只有一个 SM 的 GPU 上运行，因为一个 SM 是为消费者保留的，其余的 SM 由生产者使用。

CUDA 中的原子操作是否保证按 warp 进行调度？

Are atomic operations in CUDA guaranteed to be scheduled per warp?

cuda

gpu-atomics