Реалистичный пример тупика в CUDA / OpenCL - PullRequest
Новая
       23

Реалистичный пример тупика в CUDA / OpenCL

2 голосов
/ 21 июня 2011

Для учебника, который я пишу, я ищу "реалистичный" и простой пример тупика, вызванного незнанием SIMT / SIMD.

Я придумал этот фрагмент, который, кажется, хороший пример.

Любой вклад будет оценен. 

…
int x = threadID / 2;
if (threadID > x) {
    value[threadID] = 42;
    barrier();
    }
else {
    value2[threadID/2] = 13
    barrier();
}
result = value[threadID/2] + value2[threadID/2];

Я знаю, это не является ни CUDA C, ни OpenCL C.

1 Ответ

8 голосов
/ 21 июня 2011

Простой тупик, который на самом деле легко обнаружить начинающему программисту CUDA, - это когда кто-то пытается реализовать критическую секцию для одного потока, которая в конечном итоге должна выполняться всеми потоками. Это выглядит примерно так: 

__global__ kernel() {
  __shared__ int semaphore;
  semaphore=0;
  __syncthreads();
  while (true) {
    int prev=atomicCAS(&semaphore,0,1);
    if (prev==0) {
      //critical section
      semaphore=0;
      break;
    }
  }
}

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

Проблема в том, что, хотя 1 поток получает prev = 0, 31 поток, принадлежащий одному и тому же SIMD-блоку, получает значение 1. В операторе if CUDA планировщик переводит этот единственный поток в режим удержания (маскирует его) и пусть другие 31-нитки продолжат свою работу. В нормальных условиях это хорошая стратегия, но в данном конкретном случае вы получите 1 поток критической секции, который никогда не выполняется, и 31 поток, ожидающий бесконечности. Тупик.

Также обратите внимание на существование break, которое выводит поток управления за пределы цикла while. Если вы пропустите инструкцию break и после блока if выполняете еще несколько операций, которые должны выполняться всеми потоками, это может фактически помочь планировщику избежать тупика.

Что касается вашего примера, приведенного в вопросе: в CUDA категорически запрещено вводить __syncthreads() в SIMD-расходящийся код. Компилятор не поймает его, но в руководстве говорится о «неопределенном поведении». На практике на устройствах до Ферми все __syncthreads() рассматриваются как одинаковые барьеры. При таком предположении ваш код фактически завершится без ошибки. Однако следует не полагаться на это поведение.

...