Эффект использования страничной памяти для асинхронного копирования памяти?

В руководстве CUDA C Best Practices Guide версии 5.0, раздел 6.1.2, написано, что:

В отличие от cudaMemcpy (), асинхронная версия передачи требуется закрепленная память хоста (см. закрепленная память), и она содержит дополнительный аргумент-идентификатор потока.

Это означает, что функция cudaMemcpyAsync должна завершиться неудачей, если я использую простую память.

Но это не то, что произошло.

Просто для тестирования я попробовал следующее программа:

Ядро: 

__global__ void kernel_increment(float* src, float* dst, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if(tid<n)   
        dst[tid] = src[tid] + 1.0f;
}

Главная: 

int main()
{
    float *hPtr1, *hPtr2, *dPtr1, *dPtr2;

    const int n = 1000;

    size_t bytes = n * sizeof(float);

    cudaStream_t str1, str2;

    hPtr1 = new float[n];
    hPtr2 = new float[n];

    for(int i=0; i<n; i++)
        hPtr1[i] = static_cast<float>(i);

    cudaMalloc<float>(&dPtr1,bytes);
    cudaMalloc<float>(&dPtr2,bytes);

    dim3 block(16);
    dim3 grid((n + block.x - 1)/block.x);

    cudaStreamCreate(&str1);
    cudaStreamCreate(&str2);

    cudaMemcpyAsync(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice,str1);
    kernel_increment<<<grid,block,0,str2>>>(dPtr1,dPtr2,n);
    cudaMemcpyAsync(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost,str1);

    printf("Status: %sn",cudaGetErrorString(cudaGetLastError()));

    cudaDeviceSynchronize();

    printf("Status: %sn",cudaGetErrorString(cudaGetLastError()));

    cudaStreamDestroy(str1);
    cudaStreamDestroy(str2);

    cudaFree(dPtr1);
    cudaFree(dPtr2);

    for(int i=0; i<n; i++)
        std::cout<<hPtr2[i]<<std::endl;

    delete[] hPtr1;
    delete[] hPtr2;

    return 0;
}

Программа выдала правильный вывод. Массив успешно увеличивался.

Как cudaMemcpyAsync выполняется без блокировки страницы памяти? Я что-то упустил?

1 7
cuda asynchronous

1 ответ:

cudaMemcpyAsync является принципиально асинхронной версией cudaMemcpy. Это означает, что он не блокирует вызывающий поток хоста при выполнении вызова копирования. Это основное поведение вызова.

опционально , если вызов запускается в поток, не являющийся потоком по умолчанию, и если память хоста является закрепленным выделением, а устройство имеет свободный механизм копирования DMA, операция копирования может выполняться, пока GPU одновременно выполняет другую операцию: либо выполнение ядра, либо другую копию (в случай графического процессора с двумя DMA-копирайтерами). Если все эти условия не выполняются, то операция на GPU функционально идентична стандартному вызову cudaMemcpy, т. е. он сериализует операции на GPU, и одновременное выполнение ядра копирования или одновременное многократное копирование не может произойти. Единственное отличие заключается в том, что операция не блокирует вызывающий поток хоста.

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

cudaMemcpy(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice);
kernel_increment<<<grid,block>>>(dPtr1,dPtr2,n);
cudaMemcpy(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost);

За исключением того, что все вызовы асинхронны на хосте, поэтому хост-поток блокируется при вызове cudaDeviceSynchronize(), а не при каждом вызове передачи памяти.

Это абсолютно ожидаемое поведение.
13