Как синхронизировать Cuda по разным веткам? - PullRequest
Новая
       37

Как синхронизировать Cuda по разным веткам?

0 голосов
/ 20 марта 2019

У меня есть данные, которые я хочу обработать с помощью Cuda C ++, где я работаю с четырьмя пикселями одновременно, все четыре пикселя имеют общий угол.Например, я бы работал с четырьмя + пикселями вместе для всех мест в поле: 

------------
--------++--
--------++--
------------
------------

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

// We break the threads up into four flights:
//
//     0: even X and even Y
//     1:  odd X and even Y
//     2: even X and  odd Y
//     3:  odd X and  odd Y
const int flight = ( threadIdx.x % 2 + ( ( threadIdx.y % 2 ) << 1 ) );

for (int idx = 0; idx < flight; ++idx) {
  __syncthreads();
 }

doWork( pixel[ threadIdx.x +     threadIdx.y * blockDim.x ],
        pixel[ threadIdx.x + 1 + threadIdx.y * blockDim.x ],
        pixel[ threadIdx.x + 1 + (threadIdx.y + 1) * blockDim.x ],
        pixel[ threadIdx.x +     (threadIdx.y + 1) * blockDim.x ]);

for (int idx = 3; idx > flight; --idx) {
  __syncthreads();
 }

Цель состоит в том, чтобы разбить работу на четыре полета, которые синхронизируются сдруг друга, что-то вроде этого:

Четный X / Четный Y полет: 

doWork(...);
__syncthreads("one");
__syncthreads("two");
__syncthreads("three");

Нечетный X / Четный Y полет: 

__syncthreads("one");
doWork(...);
__syncthreads("two");
__syncthreads("three");

Четный X / Нечетный Yflight: 

__syncthreads("one");
__syncthreads("two");
doWork(...);
__syncthreads("three");

Odd X / Odd Y flight: 

__syncthreads("one");
__syncthreads("two");
__syncthreads("three");
doWork(...);

Однако я думаю, что __syncthreads() не будет делать то, что я хочу, так как мое приложение не появляетсядля правильной работы.

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

Конфигурация запуска выглядит следующим образом: 

const int32_t pixelBlockSize = <argument to function>;
const int32_t pixelGridSize  = <argument to function>;

const size_t scratch   = (pixelBlockSize * pixelBlockSize + 2) * sizeof( float );
const dim3 dimBlock( pixelBlockSize, pixelBlockSize );
const dim3 dimGrid( pixelGridSize, pixelGridSize );

CallKernel<<< dimGrid, dimBlock, scratch >>> ( ... )

Ответы [ 2 ]

1 голос
/ 20 марта 2019

__syncthreads() является точкой синхронизации. Невозможно синхронизировать несколько отдельных потоков с помощью __syncthreads(). Каждый __syncthreads() является барьером, который заставляет каждый поток в блоке ждать, пока все потоки блока не достигнут точки __syncthreads(). Вы не можете иметь __syncthreads() в расходящихся ветвях. Все (не завершенные) потоки блока должны прибыть в каждый __syncthreads(); в противном случае поведение не определено. В то время как на уровне PTX были бы способы для более точной синхронизации барьера, я не думаю, что это действительно ответ. Если я правильно понимаю вашу проблему, то все, что вы ищете, будет 

doWork(…)
__syncthreads();
doWork(…)
__syncthreads();
doWork(…)
__syncthreads();
doWork(…)

Все потоки каждого блока запускают начальный doWork(…) параллельно. Вы ждете, пока все темы не закончатся с этим. Затем вы запускаете следующую doWork(…) и т. Д. & Hellip;

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

0 голосов
/ 22 марта 2019

За ответ Михаэля Кензеля: 

// We break the threads up into four flights:
//
//     0: even X and even Y
//     1:  odd X and even Y
//     2: even X and  odd Y
//     3:  odd X and  odd Y
const int flight = ( threadIdx.x % 2 + ( ( threadIdx.y % 2 ) << 1 ) );

auto Process = [&](const bool run) {
  if ( run )
    {
      doWork( pixel[ threadIdx.x +     threadIdx.y * blockDim.x ],
              pixel[ threadIdx.x + 1 + threadIdx.y * blockDim.x ],
              pixel[ threadIdx.x + 1 + (threadIdx.y + 1) * blockDim.x ],
              pixel[ threadIdx.x +     (threadIdx.y + 1) * blockDim.x ]);
    }
};

Process( 0 == flight ); __syncthreads();
Process( 1 == flight ); __syncthreads();
Process( 2 == flight ); __syncthreads();
Process( 3 == flight );
...