Гарантированы ли атомные операции в CUDA запланированными для каждой деформации?

Question

Предположим, у меня есть 8 блоков по 32 потока, каждый из которых работает на GTX 970. Каждый blcok записывает все 1 или все 0 в массив длиной 32 в глобальной памяти, где поток 0 в блоке записывает в позицию 0 в массиве.

Теперь для записи фактических значений используется atomicExch, заменяя текущее значение в массиве значением, которое блок пытается записать. Из-за SIMD, атомарной операции и того факта, что деформация выполняется в режиме lockstep, я ожидаю, что массив в любой момент времени будет содержать только 1 или 0. Но никогда не смешивать два.

Однако при выполнении кода, подобного этому, есть несколько случаев, когда в какой-то момент времени массив содержит комбинацию из 0 и 1. Который, кажется, указывает на тот факт, что атомарные операции не выполняются за основу, а вместо этого планируются с использованием какой-то другой схемы.

Из других источников я на самом деле не нашел убедительной рецензии, подробно описывающей планирование атомарных операций в разных варпах (поправьте меня, если я ошибаюсь), поэтому мне было интересно, есть ли какая-либо информация по этой теме. Поскольку мне нужно записать много маленьких векторов, состоящих из нескольких 32-битных целых чисел атомарно, в глобальную память, и атомарная операция, которая гарантированно записывает один вектор атомарно, очевидно, очень важна.

Для тех, кто интересуется, код, который я написал, был выполнен на GTX 970, скомпилирован на вычислительных возможностях 5.2 с использованием CUDA 8.0.

Answer 1

Атомарные инструкции, как и все инструкции, запланированы на основе деформации. Однако существует неопределенный конвейер, связанный с атомарностью, и не гарантируется, что запланированный поток команд через конвейер будет выполняться в режиме блокировки для каждого потока, для каждой стадии конвейера. Это дает возможность для ваших наблюдений.

Я полагаю, что простой мысленный эксперимент продемонстрирует, что это должно быть правдой: что если 2 потока в одной и той же деформации нацелены на одно и то же место? Очевидно, что каждый аспект обработки не может идти в ногу со временем. Мы могли бы распространить этот мысленный эксперимент на случай, когда у нас есть многократная проблема за такт внутри SM и даже между SM, в качестве дополнительных примеров.

Если бы длина вектора была достаточно короткой (16 байт или меньше), то это можно было бы сделать ("атомарное обновление") просто с помощью потока в деформации, записав соответствующую величину типа вектора, например int4, Пока все потоки (независимо от того, где они находятся в сетке) пытаются обновить естественно выровненное местоположение, запись не должна быть повреждена другими записями.

Однако, после обсуждения в комментариях, кажется, что цель OP состоит в том, чтобы иметь возможность обновлять вектор некоторой длины деформацией или блоком потоков без вмешательства других деформаций или блоков потоков. Мне кажется, что действительно требуется управление доступом (так, чтобы только один деформационный блок или блок потока обновлял определенный вектор за раз), и у OP был какой-то код, который не работал должным образом.

Этот контроль доступа может быть осуществлен с использованием обычной атомарной операции (atomicCAS в приведенном ниже примере) разрешить только одному "производителю" обновлять вектор за один раз.

Ниже приведен пример кода производителя-потребителя, где есть несколько потоковых блоков, которые обновляют диапазон векторов. Каждый векторный "слот" имеет переменную "слот-контроль", которая атомарно обновляется, чтобы указать:

1. вектор пуст
2. вектор заполняется
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, так как один SM зарезервирован для потребителя, а остальные SM используются производителями.