是否可以直接从gpu访问硬盘？

Question

是否可以直接从GPU(CUDA/openCL)访问硬盘/闪存盘并直接从GPU的内存加载/存储内容？

我试图避免将内容从磁盘复制到内存,然后将其复制到GPU的内存中.

我读到了有关Nvidia GPUDirect但不确定它是否符合我上面解释的内容.它谈到远程GPU内存和磁盘,但我的情况下的磁盘是GPU的本地磁盘.

基本思想是加载内容(类似于dma) - >执行一些操作 - >将内容存储回磁盘(再次以dma方式).

我试图尽可能少地涉及CPU和RAM.

请随时提供有关设计的任何建议.

Answer 1

对于其他任何寻找这个问题的人来说,"懒惰的解开"或多或少地做了我想要的事情.

请仔细阅读以下内容,看看这对您是否有帮助.

使用RDMA for GPUDirect的最简单实现将在每次传输之前固定内存,并在传输完成后立即取消固定.不幸的是,这通常表现不佳,因为固定和取消固定内存是昂贵的操作.但是,执行RDMA传输所需的其余步骤可以在不进入内核的情况下快速执行(可以使用MMIO寄存器/命令列表缓存和重放DMA列表).

因此,延迟取消内存是高性能RDMA实现的关键.这意味着,即使在传输完成后也要保持内存固定.这利用了以下事实:可能将相同的存储器区域用于将来的DMA传输,因此延迟取消固定可以节省pin/unpin操作.

惰性取消固定的示例实现将保留一组固定的存储器区域并且仅取消其中一些(例如,最近最少使用的一个),如果区域的总大小达到某个阈值,或者如果由于BAR而固定新区域失败空间耗尽(参见PCI BAR尺寸).

这是指向应用程序指南和nvidia文档的链接.

Answer 2

尝试使用此功能，我在 Windows x64 上编写了一个小示例来实现此功能。在这个例子中，内核“直接”访问磁盘空间。实际上，正如@RobertCrovella 之前提到的，操作系统正在完成这项工作，可能还有一些 CPU 工作；但没有补充编码。

__global__ void kernel(int4* ptr)
{
    int4 val ; val.x = threadIdx.x ; val.y = blockDim.x ; val.z = blockIdx.x ; val.w = gridDim.x ;
    ptr[threadIdx.x + blockDim.x * blockIdx.x] = val ;
    ptr[160*1024*1024 + threadIdx.x + blockDim.x * blockIdx.x] = val ;
}

#include "Windows.h"

int main()
{
    // 4GB - larger than installed GPU memory
    size_t size = 256 * 1024 * 1024 * sizeof(int4) ;

    HANDLE hFile = ::CreateFile ("GPU.dump", (GENERIC_READ | GENERIC_WRITE), 0, 0, CREATE_ALWAYS, FILE_ATTRIBUTE_NORMAL, NULL) ;

    HANDLE hFileMapping = ::CreateFileMapping (hFile, 0, PAGE_READWRITE, (size >> 32), (int)size, 0) ;

    void* ptr = ::MapViewOfFile (hFileMapping, FILE_MAP_ALL_ACCESS, 0, 0, size) ;

    ::cudaSetDeviceFlags (cudaDeviceMapHost) ;

    cudaError_t er = ::cudaHostRegister (ptr, size, cudaHostRegisterMapped) ;
    if (cudaSuccess != er)
    {
        printf ("could not register\n") ;
        return 1 ;
    }

    void* d_ptr ;
    er = ::cudaHostGetDevicePointer (&d_ptr, ptr, 0) ;
    if (cudaSuccess != er)
    {
        printf ("could not get device pointer\n") ;
        return 1 ;
    }

    kernel<<<256,256>>> ((int4*)d_ptr) ;

    if (cudaSuccess != ::cudaDeviceSynchronize())
    {
        printf ("error in kernel\n") ;
        return 1 ;
    }

    if (cudaSuccess != ::cudaHostUnregister (ptr))
    {
        printf ("could not unregister\n") ;
        return 1 ;
    }

    ::UnmapViewOfFile (ptr) ;

    ::CloseHandle (hFileMapping) ;
    ::CloseHandle (hFile) ; 

    ::cudaDeviceReset() ;

    printf ("DONE\n");

    return 0 ;
}