cuda 统一内存模式

cuda基本编程步骤这一篇文章中,可以看到cuda编程中需要在两个地方分配内存(主机和设备之间),还需要手动拷贝数据,这样使用起来非常的不方便。从cuda 6.0开始,提供了统一内存的模式,可以方便cuda的编程。要利用统一内存模式,需要SM的架构大于3.0,而且操作系统必须时64位的。

在统一内存模式中,内存分配的方式是:

int *value;
cudaMallocManaged(&value, sizeof(int));

还可以结合std中的allocator进行高效的内存分配:

#include <cstdio>
#include <cstddef>
#include <vector>
#include <cuda_runtime.h>
// 
__global__ void kernel(int *arr, int n)
{
    for(int i=blockDim.x * blockIdx.x + threadIdx.x;
        i<n; i+=blockDim.x * gridDim.x)
    {
        arr[i] = i;
    }
}

template<class T>
struct cudaAllocator
{
    using value_type = T;
    
    T* allocate(size_t size)
    {
        T* ptr = nullptr;
        cudaError error = cudaMallocManaged(&ptr, sizeof(T) * size);
        return ptr;
    }

    void deallocate(T* ptr, size_t size = 0)
    {
        cudaFree(ptr);
    }

    template<class ...Args>
    void contruct(T*, Args && ... args)
    {
        if constexpr (!(sizeof...(Args) == 0 && std::is_pod_v<T>))
            ::new((void*)p) T(std::forward<Args>(args)...);
    }

    template<class U>
    cudaAllocator(const cudaAllocator<U>&) noexcept {}

    cudaAllocator() noexcept {}
};

int main()
{
    int n = 65535;
    std::vector<int, cudaAllocator<int>> arr(n);

    kernel<<<32,128>>>(arr.data(), n);
    cudaDeviceSynchronize();
    for(int i=0; i<n; i++)
    {
        printf("arr[%d]:%d\n", i, arr[i]);
    }
    
    return 0;
}