Какова связь между занятостью потоков графического процессора и остановками синхронизации?

Повлияет ли количество потоков на блок на среднее время, которое поток проводит в ожидании у барьера?

Вероятный. Большее количество потоков на блок, безусловно, увеличивает худший случай. Допустим, вы запускаете 512 потоков на блок на RTX 3080 (CUDA 8.6, 1536 потоков на SM), затем три варпа (по одному на блок) могут заблокировать продвижение вперед 45 других варпов на том же SM из-за опоздания.

Вам следует провести собственные тесты, но в прошлый раз, когда я тестировал это, я использовал блоки меньшего размера, даже если это немного снизило эффективность памяти.

Общее количество времени, которое все потоки проводят в ожидании? Что делать, если у меня ядро ​​с низкой занятостью (т. е. много регистров на поток, малое количество потоков на блок)

Что ж, низкая занятость, вероятно, уменьшит время ожидания каждого потока просто потому, что потокам придется тратить меньше времени на ожидание доступности исполнительного модуля. Но нужно быть осторожным, когда вы это делаете. Если ваши оставшиеся ядра не смогут использовать всю вычислительную производительность графического процессора или полностью скрыть задержку, вы, очевидно, проиграете.

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

существуют ли какие-либо стратегии, которые могут помочь уменьшить ожидание синхронизации в этом случае?

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

__global__ void kernel()
{
  __shared__ float data[N];
  for (int i = 0; i < NUM_ITERATIONS; i++)
  {
    data[threadIdx.x] = load_global();
    __syncthreads();
    float loc = data[y];
    __syncthreads();
  }
}

сделай это:

__global__ void kernel()
{
  __shared__ float data[2][N];
  for (int i = 0; i < NUM_ITERATIONS; i++)
  {
    data[i & 1][threadIdx.x] = load_global();
    __syncthreads();
    float loc = data[i & 1][y];
  }
}

Вы можете расширить это, используя Асинхронные копии данных . Примеры CUDA показывают, как это сделать, в примере globalToShmemAsyncCopy. Я думаю, что ядро ​​MatrixMulAsyncCopyMultiStage — это то, что вам нужно.