cuda
Межблочная связь

замечания

Защитник последнего блока

Рассмотрим сетку, работающую над некоторой задачей, например параллельное сокращение. Первоначально каждый блок может выполнять свою работу независимо, производя некоторый частичный результат. В конце концов, однако, частичные результаты необходимо объединить и объединить. Типичным примером является алгоритм сокращения больших данных.

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

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

С таким защитником последний блок гарантированно будет видеть все результаты, полученные всеми другими блоками, и может выполнять слияние.

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

Предположения:

• Счетчик должен быть глобальным указателем памяти, инициализированным до 0 до вызова ядра.
• Функция lastBlock вызывается равномерно всеми потоками во всех блоках
• Ядро вызывается в одномерной сетке (для простоты примера)
• T называет любой тип, который вам нравится, но этот пример не предназначен для шаблона в смысле C ++

Глобальная рабочая очередь

Рассмотрим множество рабочих элементов. Время, необходимое для завершения каждого рабочего элемента, сильно различается. Чтобы сбалансировать распределение работы между блоками, для каждого блока может быть разумным выборку следующего элемента только тогда, когда предыдущий закончен. Это контрастирует с априорным назначением элементов блокам.

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

Предположения:

• Объект WorkQueue, а также массив gItem находятся в глобальной памяти
• Никакие новые рабочие элементы не добавляются в объект WorkQueue в ядре, которое извлекает из него
• WorkItem - небольшое представление задания, например указатель на другой объект
• Функция WorkItem::none() static member создает объект WorkItem который не представляет никакой работы вообще
• WorkQueue::fetch() должен быть вызван равномерно всеми потоками в блоке
• Нет никаких 2 вызовов WorkQueue::fetch() без другого __syncthreads() между ними. В противном случае появится состояние гонки!

В примере не указано, как инициализировать WorkQueue или заполнить его. Это делается другим кодом ядра или ЦП и должно быть довольно простым.