Синхронизация ядер

Question

Я новичок в программировании на Cuda и реализую классический алгоритм Floyd APSP.Этот алгоритм состоит из 3 вложенных циклов, и весь код внутри двух внутренних циклов может выполняться параллельно.

В качестве основных частей моего кода приведен код ядра:

__global__ void dfloyd(double *dM, size_t k, size_t n)
{
    unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int index = y * n + x;
    double d;

    if (x < n && y < n)
    {
        d=dM[x+k*n] + dM[k+y*n];
        if (d<dM[index])
            dM[index]=d;
    }
}

, а здесь - часть основной функции, в которой запускаются ядра (для удобства чтения я опустил код обработки ошибок):

double *dM;
cudaMalloc((void **)&dM, sizeof_M);
cudaMemcpy(dM, hM, sizeof_M, cudaMemcpyHostToDevice);

int dimx = 32;
int dimy = 32;
dim3 block(dimx, dimy);
dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y);

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
    cudaDeviceSynchronize();
}

cudaMemcpy(hM, dM, sizeof_M, cudaMemcpyDeviceToHost);

[Для понимания dM относится к матрице расстояний, хранящейся на стороне устройства, и hM на стороне хоста, а n относится к числу узлов.]

* 1010Ядра * в цикле k должны выполняться последовательно, это объясняет, почему я пишу инструкцию cudaDeviceSynchronize() после каждого выполнения ядра.Тем не менее, я заметил, что помещение этой инструкции синхронизации за пределы цикла приводит к тому же результату.

Теперь мой вопрос.Два следующих фрагмента кода

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
    cudaDeviceSynchronize();
}

и

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
}
cudaDeviceSynchronize();

эквивалентны?

Answer 1

Они не эквивалентны, но дадут те же результаты. Первая заставит хост ждать после каждого вызова ядра, пока ядро ​​не вернется, а другая заставит его ждать только один раз. Может быть, запутанная часть почему это работает; в CUDA два последовательных вызова ядра в одном и том же потоке (в вашем случае поток по умолчанию) гарантированно будут выполняться последовательно.

По производительности рекомендуется использовать вторую версию, поскольку синхронизация с хостом увеличивает накладные расходы.

Редактировать: в этом конкретном случае вам даже не нужно звонить cudaDeviceSynchronize(), потому что cudaMemcpy будет синхронизироваться.