Работа с unified-памятью в CUDA

Обратите внимание, что можно использовать эту функцию при переопределении оператора new - в этом случае будет сразу выделяться unified-память. Это позволяет легко размещать различные объекты в такой памяти и делать их доступными для всех. Обратите внимание, что для GPU до Pascal нельзя одновременно работать с такой памятью и со стороны GPU и со стороны CPU. Для Pascal и выше такое уже становится возможно, но при этом надо соблюдать осторожность, чтобы не возникали race condition.

Кроме того, использование подобной памяти дает еще одно преимущество. Обычно объем памяти GPU ограничен (в отличии от объема легко наращиваемой оперативной памяти) и поэтому выделить десятки гигабайт памяти GPU обычным образом невозможно. Однако поскольку при выделении через cudaMallocManaged мы фактически имеем дело с виртуальной памятью, то возможно выделять большие объемы памяти. При этом при обращении к ней со стороны GPU (любого из) будет происходить автоматический перенос данных в память данного GPU. таким образом появляется возможность работы с очень большим объемом виртуальной памяти, которая при обращении к ней, автоматически копируется в нужную сторону, снимая тем самым ограничения на объем памяти, фактически установленной на GPU.

Использование подобной памяти сильно упрощает жизнь программиста, но также оно обладает определенной ценой, которую приходится платить на подобную схему копирования данных. В результате быстродействие обычно оказывается ниже, чем при явном копировании данных. Тем не менее, существую способы повысить скорость работы c unified-памятью. Для этого можно использовать функции cudaMemAdvise и cudaMemPrefetchAsync.

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

При помощи функции cudaMemAdvise можно сообщать о предполагаемом характере использования памяти, что может помочь системе организовать работу более эффективно.

В этом вызове параметр advice может принимать следующие значения:

• cudaMemAdviseSetReadMostly - данное значение задает, что доступ к данной области памяти в основном будет на чтение;
• cudaMemAdviseUnsetReadMostly - данное значение отменяет cudaMemAdviseSetReadMostly;
• cudaMemAdviseSetPreferredLocation - данное значение задает предпочтительное размещение данного блока в памяти устройства, задаваемого параметрами device;
• cudaMemAdviseUnsetPreferredLocation - отменяет cudaMemAdviseSetPreferredLocation;
• cudaMemAdviseSetAccessedBy - задает, что к данной области памяти будет идти со стороны устройства device;
• cudaMemAdviseUnsetAccessedBy - отменяет cudaMemAdviseUnsetAccessedBy

Однако большего ускорения можно добиться при помощи команды cudaMemPrefetchAsync, описание которой приводится ниже.

int main ()
{
    int   * ptr;
    int n = 160*1024*1024;      // # of elements
    int blockSize = 1024;
    int numBlocks = (n + blockSize - 1) / blockSize;
    int device;

    cudaGetDevice ( &device );

    printf ( "Device %d\n", device );

                    // allocate managed memory
    cudaMallocManaged ( &ptr, n * sizeof ( int ) );

    for ( int i = 0; i < n; i++ )
        ptr [i] = i;

    cudaMemPrefetchAsync ( ptr, n * sizeof(int), device );

    kernel<<<numBlocks, blockSize>>> ( ptr, n );

    cudaDeviceSynchronize ();

    for ( int i = 0; i < n; i++ )
        if ( ptr [i] != i + 1 )
        {
            printf ( "Error at %d - %d\n", i, ptr [i] );
            return 1;
        }

    cudaFree ( ptr );

    printf ( "Success\n" );

    return 0;
}

При запросе на ранее копирование (prefetch), этот запрос может быть отложен (defer), если соответствующий поток (имеется в виду именно stream) занят. В этом случае, когда все операции к этом потоке, поданные на выполнения до запроса cudaMemPrefetchAsync будут завершены, системная нить (на этот раз имеется в виду именно thread) в драйвере выполнит данный запрос. Если поток не занят, то запрос по решению драйвера может быть как отложен, так и сразу же выполнен.

При немедленном выполнении такого запроса его поведение зависит от того, идет ли копирование из CPU на GPU или наоборот. В случае копирования из памяти CPU в память GPU управление возвращается сразу же, после обновления таблиц страниц на стороне CPU, при этом само же копирование просто ставится в очередь.

Однако при копировании со стороны GPU на сторону CPU управление возвращается только после того, как будут завершены все операции копирования - это связано с тем, что таблицы страниц для CPU не могут обновляться асинхронно.

Для получения хорошего быстродействия при копирования данных в одну сторону, необходимо сначала запустить ядро (это асинхронная операция, она просто поставит запрос в очередь), а уже потом вызвать cudaMemPrefetchAsync.

Для получения хорошего быстродействия в случае копирования данных в обе стороны, получения оптимального высокого быстродействия несколько сложнее. Сайт developer.nvidia.com рекомендует следующий подход, основанный на том, чтобы избежать так называемой сериализации операций копирования - когда они будут выполняться не параллельно, а последовательно. Это снизит итоговое быстродействие программы.

Рассмотрим следующий пример:

__global__ void kernel ( int n, int * x )
{
    int index  = blockIdx.x * blockDim.x + threadIdx.x;
    
    if ( index < n )
        x [index] += index;
}

int main ()
{
    int device;

    cudaGetDevice ( &device );

    int   n = 1<<25;
    int * a;
    int   tileSize = 1024 * 128;        // 64K items per tile
    int   numTiles = (n + tileSize - 1) / tileSize;
    int   numRuns  = 10;

    printf ( "Device %d\nnumTiles %d\n", device, numTiles );

    
    cudaStream_t    s1, s2, s3;
    
    cudaStreamCreate ( &s1 );
    cudaStreamCreate ( &s2 );
    cudaStreamCreate ( &s3 );
    
            // Allocate Unified Memory -- accessible from CPU or GPU
    cudaMallocManaged ( &a, n * sizeof(int) );
 
            // initialize array a on the host
    for ( int i = 0; i < n; i++ ) 
        a [i] = 1;
 
            // do processing in tiles
    for ( int i = 0; i < numTiles; i++ ) 
    { 
            // offload previous tile to the cpu if (i > 0) 
        if ( i > 0 )
            cudaMemPrefetchAsync ( a + tileSize * (i-1), tileSize * sizeof(int), cudaCpuDeviceId, s1 ); 

            // run multiple kernels on current tile 
        for ( int j = 0; j < numRuns; j++ ) 
            kernel<<<1024, 1024, 0, s2>>>( tileSize, a + tileSize * i ); 

            // prefetch next tile to the gpu 
        if ( i < numTiles ) 
            cudaMemPrefetchAsync ( a + tileSize * (i+1), tileSize * sizeof(int), device, s3 ); 

            // sync all streams 
        cudaDeviceSynchronize(); 
    } 
            // Check for errors (all values should be 3.0f)
        // Free memory
    cudaFree          ( a );
    cudaStreamDestroy ( s1 );
    cudaStreamDestroy ( s2 );
    cudaStreamDestroy ( s3 );
    
    printf ( "Sucess\n" );
    
    return 0;
}

В данном случае мы используем три потока (stream) - s1, s2 и s3 - по одному для копирования в каждую сторону и выполнения ядра. Сама работа разбита на части, в ходе каждой из частей на своем потоке запускается ранее чтения данных в каждую из сторон и ядро. Согласно developer.nvidia.com для этого кода операции копирования будут происходиьт последовательно. На CUDA 9.1 я этого не наблюдал.

Однако ниже приводится рекомендованная версия, в которой мы переносим копирования GPU->CPU в "загруженный" поток (на котором выполняется ядро), в копирование CPU->GPU оставляем в своем потоке. Она дала небольшой прирост скорости, по сравнению с первоначальной.

int main ()
{
    int device;

    cudaGetDevice ( &device );

    int   n = 1<<25;
    int * a;
    int   tileSize = 1024 * 128;        // 64K items per tile
    int   numTiles = (n + tileSize - 1) / tileSize;
    int   numRuns  = 10;

    printf ( "Device %d\nnumTiles %d\n", device, numTiles );

    
    cudaStream_t    s1, s2, s3, st;
    cudaEvent_t     e1, e2, et;
    
    cudaStreamCreate ( &s1 );
    cudaStreamCreate ( &s2 );
    cudaStreamCreate ( &s3 );
    cudaEventCreate  ( &e1 );
    cudaEventCreate  ( &e2 );
    
            // Allocate Unified Memory -- accessible from CPU or GPU
    cudaMallocManaged ( &a, n * sizeof(int) );
 
            // initialize array a on the host
    for ( int i = 0; i < n; i++ ) 
        a [i] = 1;

            // prefetch first tile
    cudaMemPrefetchAsync ( a, tileSize * sizeof(int), 0, s2);
    cudaEventRecord      ( e1, s2 ); 

    for ( int i = 0; i < numTiles; i++ ) 
    { 
            // make sure previous kernel and current tile copy both completed 
        cudaEventSynchronize ( e1 );  
        cudaEventSynchronize ( e2 );

            // run multiple kernels on current tile 
        for (int j = 0; j < numRuns; j++)
            kernel<<<1024, 1024, 0, s1>>> ( tileSize, a + tileSize * i ); 

        cudaEventRecord ( e1, s1 );

            // prefetch next tile to the gpu in a separate stream 
        if ( i < numTiles - 1 ) 
        {
                // make sure the stream is idle to force non-deferred HtoD prefetches first 
            cudaStreamSynchronize ( s2 );       
            cudaMemPrefetchAsync  ( a + tileSize * (i+1), tileSize * sizeof(int), 0, s2 ); 
            cudaEventRecord       ( e2, s2 );
        } 

            // offload current tile to the cpu after the kernel is completed using the deferred path 
        cudaMemPrefetchAsync ( a + tileSize * i, tileSize * sizeof(int), cudaCpuDeviceId, s1 );

            // rotate streams and swap events 
        st = s1; s1 = s2; s2 = st; 
        st = s2; s2 = s3; s3 = st; 
        et = e1; e1 = e2; e2 = et; 
    }
 
        // Free memory
    cudaFree          ( a  );
    cudaEventDestroy  ( e1 );
    cudaEventDestroy  ( e2 );
    cudaStreamDestroy ( s1 );
    cudaStreamDestroy ( s2 );
    cudaStreamDestroy ( s3 );
    
    printf ( "Sucess\n" );

    return 0;
}