Question

Есть ли функция, которая предоставляет эти функции одновременно?Я ищу функцию, которая распределяет память, которая имеет черты как «сопоставленные с памятью» (например, выделенные с mmap), так и UVM (доступная как с хоста, так и с устройства с графическим процессором).Я вижу, что cudaHostAlloc выделяет память в памяти хоста, которая доступна для устройств, но нет очевидного способа объявить выделенные диапазоны памяти как отображенные в памяти!

У меня такой вопрос: есть ли APIфункция для выделения памяти с вышеупомянутыми признаками?

Если ответ на поставленный выше вопрос - «нет», то есть ли набор функций API, которые я могу вызвать, что приводит к тому же поведению?

Например, сначала мы используем cudaMallocManaged для выделения памяти на основе UVM, а затем используем определенный API (либо POSIX, либо CUDA API), чтобы объявить ранее выделенную память как «отображенную в память» (точно так же, какmmap)?Или, наоборот, (выделите с помощью mmap, а затем объявите диапазон как UVM для драйвера CUDA)?

Любые другие предложения также будут оценены!

ОБНОВЛЕНИЕ 13 декабря 2018 года:

К сожалению, предложение @tera, похоже, не работает, как ожидалось.Когда код выполняется на устройстве, кажется, что устройство не может видеть память на хосте!

Ниже приведен код, который я использую с командой компиляции.

#include <stdio.h>
#include <stdlib.h>
#include <sys/mman.h>
#include <sys/types.h>
#include <fcntl.h>
#include <unistd.h>
#include <sys/stat.h>
#include <assert.h>


__global__
void touchKernel(char *d, char init, int n) {
    int index =  blockIdx.x *blockDim.x + threadIdx.x;
    if(index >= n)
        return;
    d[index] = init;
}


void process_file(char* filename, int n) {
    if(n < 0) {
        printf("Error in n: %d\n", n);
        exit(1);
    }
    size_t filesize = n*sizeof(char);
    size_t pagesize = (size_t) sysconf (_SC_PAGESIZE);

    //Open file
    int fd = open(filename, O_RDWR|O_CREAT, 0666);
    // assert(fd != -1);
    if(fd == -1) {
        perror("Open API");
        exit(1);
    }
    ftruncate(fd, filesize);

    //Execute mmap
    char* mmappedData = (char*) mmap(0, filesize, PROT_READ|PROT_WRITE, MAP_SHARED|MAP_LOCKED, fd, 0);
    assert(mmappedData != MAP_FAILED);
    printf("mmappedData: %p\n", mmappedData);

    for(int i=0;i<n;i++)
        mmappedData[i] = 'z';

    if(cudaSuccess != cudaHostRegister(mmappedData, filesize, cudaHostRegisterDefault)) {
        printf("Unable to register with CUDA!\n");
        exit(1);
    }

    int vec = 256;
    int gang = (n) / vec + 1;
    printf("gang: %d - vec: %d\n", gang, vec);
    touchKernel<<<gang, vec>>>((char*) mmappedData, 'a', n);
    cudaDeviceSynchronize();

    //Cleanup
    int rc = munmap(mmappedData, filesize);
    assert(rc == 0);


    close(fd);
}


int main(int argc, char const *argv[])
{
    process_file("buffer.obj", 10);

    return 0;
}

И для компиляции, вот он:

nvcc -g -O0 f1.cu && cuda-memcheck ./a.out

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

========= Invalid __global__ write of size 1
=========     at 0x000000b0 in touchKernel(char*, char, int)
=========     by thread (2,0,0) in block (0,0,0)
=========     Address 0x7fdc8e137002 is out of bounds
=========     Device Frame:touchKernel(char*, char, int) (touchKernel(char*, char, int) : 0xb0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x24d9dd]
=========     Host Frame:./a.out [0x22b22]
=========     Host Frame:./a.out [0x22d17]
=========     Host Frame:./a.out [0x570d5]
=========     Host Frame:./a.out [0x6db8]
=========     Host Frame:./a.out [0x6c76]
=========     Host Frame:./a.out [0x6cc3]
=========     Host Frame:./a.out [0x6a4c]
=========     Host Frame:./a.out [0x6ade]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./a.out [0x673a]
=========
========= Invalid __global__ write of size 1
=========     at 0x000000b0 in touchKernel(char*, char, int)
=========     by thread (1,0,0) in block (0,0,0)
=========     Address 0x7fdc8e137001 is out of bounds
=========     Device Frame:touchKernel(char*, char, int) (touchKernel(char*, char, int) : 0xb0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x24d9dd]
=========     Host Frame:./a.out [0x22b22]
=========     Host Frame:./a.out [0x22d17]
=========     Host Frame:./a.out [0x570d5]
=========     Host Frame:./a.out [0x6db8]
=========     Host Frame:./a.out [0x6c76]
=========     Host Frame:./a.out [0x6cc3]
=========     Host Frame:./a.out [0x6a4c]
=========     Host Frame:./a.out [0x6ade]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./a.out [0x673a]
=========
========= Invalid __global__ write of size 1
=========     at 0x000000b0 in touchKernel(char*, char, int)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7fdc8e137000 is out of bounds
=========     Device Frame:touchKernel(char*, char, int) (touchKernel(char*, char, int) : 0xb0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x24d9dd]
=========     Host Frame:./a.out [0x22b22]
=========     Host Frame:./a.out [0x22d17]
=========     Host Frame:./a.out [0x570d5]
=========     Host Frame:./a.out [0x6db8]
=========     Host Frame:./a.out [0x6c76]
=========     Host Frame:./a.out [0x6cc3]
=========     Host Frame:./a.out [0x6a4c]
=========     Host Frame:./a.out [0x6ade]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./a.out [0x673a]
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize. 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x351c13]
=========     Host Frame:./a.out [0x40a16]
=========     Host Frame:./a.out [0x6a51]
=========     Host Frame:./a.out [0x6ade]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./a.out [0x673a]
=========

Вышеприведенный вывод означает, что код не был успешно выполнен на устройстве.

Есть предложения?

ОБНОВЛЕНИЕ 14 декабря 2018 года

Я изменил код на следующий:

__global__
void touchKernel(char *d, char init, int n) {
    int index =  blockIdx.x *blockDim.x + threadIdx.x;
    if(index >= n || index < 0)
        return;
    printf("index %d\n", index);
    d[index] = init + (index%20);
    printf("index %d - Done\n", index);
}

Если приведенный выше код заменить старым, можно увидеть вывод обеих printf команд.Если проверить файл buffer.obj, они увидят, что файл содержит правильный вывод!

ОБНОВЛЕНИЕ 14 декабря 2018

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

tera · Answer 1 · 13 декабря 2018

По совпадению я только что ответил на похожий вопрос на форуме Nvidia.

Вы можете cudaHostRegister() отобразить память, если передать флаг MAP_LOCKED в mmap().

Возможно, вам придется увеличитьограничение для заблокированной памяти (ulimit -m в bash) при этом.

Обновление: Оказывается, MAP_LOCKED flag до mmap() даже не требуется.Однако в документации к cudaHostRegister() перечислены некоторые другие ограничения:

В системах без унифицированной виртуальной адресации флаг cudaHostRegisterMapped необходимо передать в cudaHostRegister() илипамять не будет отображена.Если устройство не имеет ненулевого значения для атрибута cudaDevAttrCanUseHostPointerForRegisteredMem, это также означает, что вам нужно запросить адрес устройства для диапазона отображаемой памяти через cudaHostGetDevicePointer().
Контекст CUDA должен быть создан с флагом cudaMapHost, чтобы было возможно отображение.Поскольку контекст создается лениво API-интерфейсом среды выполнения, вам необходимо создать контекст самостоятельно с помощью API-драйвера перед любым вызовом API-интерфейса среды выполнения, чтобы иметь возможность влиять на флаги, с которыми создается контекст.

Сочетание функций mmap и UVM

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

Сочетание функций mmap и UVM

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

Похожие темы