Пример реалистичной взаимоблокировки в CUDA/OpenCL

Для учебника, который я пишу, я ищу «реалистичный» и простой пример взаимоблокировки, вызванной незнанием 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.


parallel-processing synchronization cuda simd opencl
person Framester    schedule 21.06.2011    source источник
comment
Это кажется слишком сложным для примера и довольно простым для реалистического примера. Я бы использовал только get_local_id(0) > constant в условии и заменил бизнес-код (назначение) комментариями /* do some stuff */ и /* do another stuff */. Тем не менее, я думаю, что StackOverflow — не лучшее место для дискуссий, это место для вопросов и ответов.   -  person Radim Vansa    schedule 21.06.2011

Ответы (1)


arrow_upward
8
arrow_downward

Простая взаимоблокировка, которую на самом деле легко поймать начинающему программисту 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() рассматриваются как одни и те же барьеры. При таком предположении ваш код фактически завершится без ошибки. Однако не следует не полагаться на такое поведение.

person CygnusX1    schedule 21.06.2011