在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;
}