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

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

Очень мощной возможностью, впервые появившейся в CUDA 6, является так называемая unified-память. Она позволяет вам работать с единым адресным пространством, в который входит как вся системная (обычная) память, так и память всех установленных GPU (подерживающих CUDA).

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

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

unified memory diagram

Рис 1. Unified=память в системе

GPU класса Pascal (серия 10хх) и выше заметно расширяют возможности и удобство использоваиня этой памяти. Поэтому далее мы будем рассматривать все только для GPU Pascal и выше. С одной стороны используется единое 49-битовое адресное пространство. А с другой - в архитектуре GPU появился Page Migration engine.

Для GPU предыдущих поколений перед запуском ядра производится полное копирование всех данных, перенесенных ранее в память CPU, обратно на GPU (вне зависимости от того, будут они использоваться GPU или нет). наличие Page Migration engine позволяет избежать ненужного копирования - при обращении GPU к памяти, которая еще не пересена на данный GPU, происходит page fault. Далее останавливаются два SM (два в силу особеностей Page Migration engine), после чего производится копирование соответствующей памяти из CPU на GPU. После того, как копирование будет завершено, работа SM будет продолжена.

Для выделения unified-память используеться функция cudaMallocManaged, для освобождения - уже знакомая функция cudaFree.

cudaError_t cudaMallocManaged ( void **ptr, size_t size, unsigned int flags = cudaMemAttachGlobal );

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

__global__ void kernel ( int * data, int n )
{
    size_t index = threadIdx.x + blockIdx.x * blockDim.x;

    if ( index < n )
        data [index] = index;
}

int main ()
{
    int   * ptr;
    int n = 1024*1024;      // # of elements
    int blockSize = 1024;
    int numBlocks = (n + blockSize - 1) / blockSize;
                    // allocate managed memory
    cudaMallocManaged ( &ptr, n * sizeof ( int ) );

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

    cudaDeviceSynchronize ();

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

    cudaFree ( ptr );

    printf ( "Success\n" );

    return 0;
}

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

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

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

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

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

cudaError_t cudaMemAdvise ( const void * ptr, size_t size, cudaMemoryAdvice advice, int device );

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

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

cudaError_t cudaMemPrefetchAsync ( const void * ptr, size_t size, int device, cudaStream_t stream = 0 );
Эта функция служит для раннего копирования данных на заданное устройство (CPU или GPU). Параметр ptr задает адрес начала блока, который необходимо заранее скопирвать в память, и параметр size задает размер блока в байтах. Задаваемый блок памяти должен целиком лежать в области памяти, выделделнной через cudaMallocManaged. Параметр device задает устройство, в память которого необходимо произвести копирование. Если в качестве этого устройства выступает CPU, то этот параметр принимает значение cudaCpuDeviceId.

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), если соответствующbq поток (имеется в виду именно stream) pfyzn/ В этом случае, когда все операции к этом потоке, поданные на выполнения до запроса 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;
}

Важное замечение - на текущий момент (CUDA 9.1) полноцення поддержка копирования только используемых страниц под Windows не работает, под Linux все работает отлично :(((

По этой ссылке можно скачать весь исходный код к этой статье.