CUDA零拷贝内存（zerocopy memory）

为了实现CPU与GPU内存的共享，cuda采用了零拷贝内存，它值固定内存的一种，当然，也就是实际存储空间实在cpu上。

零拷贝内存的延迟高，在进行频繁的读写操作时尽量少用，否则会大大降低性能。

/**创建固定内存映射** flags:   cudaHostAllocDefault:      make cudaHostAlloc same as "cudaMallocHost"*          cudaHostAllocPortable:     函数返回能被所有cuda上下文使用的固定内存，而不仅仅是执行内存分配的那个*          cudaHostAllocWriteCombined:返回写结合的内存，该内存可以在某个系统配置上通过PCIe总线上更快的传输*          cudaHostAllocMapped:       该标志返回映射到的设备内存地址空间的主机内存*/cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

/*** 获取映射得到的固定内存的设备指针*/cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

整体的使用过程：

    size_t nBytes = 2048;float *h_A;h_A = (float *)malloc(nBytes);initialData(h_A, nBytes);cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0);cudaFreeHost(h_A);

给个完整程序，源码网址：点击打开链接

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>/** This example demonstrates the use of zero-copy memory to remove the need to* explicitly issue a memcpy operation between the host and device. By mapping* host, page-locked memory into the device's address space, the address can* directly reference a host array and transfer its contents over the PCIe bus.** This example compares performing a vector addition with and without zero-copy* memory.*/void checkResult(float *hostRef, float *gpuRef, const int N)
{double epsilon = 1.0E-8;for (int i = 0; i < N; i++){if (abs(hostRef[i] - gpuRef[i]) > epsilon){printf("Arrays do not match!\n");printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],gpuRef[i], i);break;}}return;
}void initialData(float *ip, int size)
{int i;for (i = 0; i < size; i++){ip[i] = (float)( rand() & 0xFF ) / 10.0f;}return;
}void sumArraysOnHost(float *A, float *B, float *C, const int N)
{for (int idx = 0; idx < N; idx++){C[idx] = A[idx] + B[idx];}
}__global__ void sumArrays(float *A, float *B, float *C, const int N)
{int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < N) C[i] = A[i] + B[i];
}__global__ void sumArraysZeroCopy(float *A, float *B, float *C, const int N)
{int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < N) C[i] = A[i] + B[i];
}int main(int argc, char **argv)
{// set up deviceint dev = 0;CHECK(cudaSetDevice(dev));// get device propertiescudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(&deviceProp, dev));// check if support mapped memoryif (!deviceProp.canMapHostMemory){printf("Device %d does not support mapping CPU host memory!\n", dev);CHECK(cudaDeviceReset());exit(EXIT_SUCCESS);}printf("Using Device %d: %s ", dev, deviceProp.name);// set up data size of vectorsint ipower = 10;if (argc > 1) ipower = atoi(argv[1]);int nElem = 1 << ipower;size_t nBytes = nElem * sizeof(float);if (ipower < 18){printf("Vector size %d power %d  nbytes  %3.0f KB\n", nElem, ipower,(float)nBytes / (1024.0f));}else{printf("Vector size %d power %d  nbytes  %3.0f MB\n", nElem, ipower,(float)nBytes / (1024.0f * 1024.0f));}// part 1: using device memory// malloc host memoryfloat *h_A, *h_B, *hostRef, *gpuRef;h_A     = (float *)malloc(nBytes);h_B     = (float *)malloc(nBytes);hostRef = (float *)malloc(nBytes);gpuRef  = (float *)malloc(nBytes);// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef,  0, nBytes);// add vector at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// malloc device global memoryfloat *d_A, *d_B, *d_C;CHECK(cudaMalloc((float**)&d_A, nBytes));CHECK(cudaMalloc((float**)&d_B, nBytes));CHECK(cudaMalloc((float**)&d_C, nBytes));// transfer data from host to deviceCHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));// set up execution configurationint iLen = 512;dim3 block (iLen);dim3 grid  ((nElem + block.x - 1) / block.x);sumArrays<<<grid, block>>>(d_A, d_B, d_C, nElem);// copy kernel result back to host sideCHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));// check device resultscheckResult(hostRef, gpuRef, nElem);// free device global memoryCHECK(cudaFree(d_A));CHECK(cudaFree(d_B));// free host memoryfree(h_A);free(h_B);// part 2: using zerocopy memory for array A and B// allocate zerocpy memoryCHECK(cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped));CHECK(cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped));// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef,  0, nBytes);// pass the pointer to deviceCHECK(cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0));CHECK(cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0));// add at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// execute kernel with zero copy memorysumArraysZeroCopy<<<grid, block>>>(d_A, d_B, d_C, nElem);// copy kernel result back to host sideCHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));// check device resultscheckResult(hostRef, gpuRef, nElem);// free  memoryCHECK(cudaFree(d_C));CHECK(cudaFreeHost(h_A));CHECK(cudaFreeHost(h_B));free(hostRef);free(gpuRef);// reset deviceCHECK(cudaDeviceReset());return EXIT_SUCCESS;
}

CUDA零拷贝内存（zerocopy memory）相关推荐

CUDA Samples: dot product(使用零拷贝内存)
以下CUDA sample是分别用C++和CUDA实现的点积运算code,CUDA包括普通实现和采用零拷贝内存实现两种,并对其中使用到的CUDA函数进行了解说,code参考了<GPU高性能编程C ...
理解Netty中的零拷贝（Zero-Copy）机制
理解Netty中的零拷贝(Zero-Copy)机制发表于2年前(2014-01-13 15:11) 阅读(10209) | 评论(12) 164人收藏此文章,我要收藏赞29 12月12日北京O ...
【Netty】零拷贝（zero-copy）
目录 1. 零拷贝技术实现 2. 传统读取IO流的操作 2.1 读操作 2.2 写操作 2.3 MMAP+write 2.4 Sendfile 3. 零拷贝应用场景很多更新的技术在宣传的时候,都会提 ...
Linux Zero-copy零拷贝技术：源码示例
<Linux Zero-copy零拷贝技术:源码示例> <Linux Zero-copy零拷贝技术全面揭秘> <什么是mmap?零拷贝?DMA?> <Linu ...
Linux I/O原理和零拷贝Zero-copy技术全面揭秘
目录导言计算机存储器物理内存虚拟内存静态重定位存储器抽象交换(swapping)技术虚拟内存技术用户态和内核态 Linux I/O I/O 缓冲区 I/O 模式程序控制 I/O 中 ...
Kafka和RocketMQ底层存储：零拷贝技术
零拷本相关 <[转]零拷贝的实现原理> <[转]零拷贝的实现原理> <搞懂Linux零拷贝,DMA> <通过零拷贝进行有效的数据传输(java.c)> ...
什么是mmap？零拷贝？DMA？
<Linux Zero-copy零拷贝技术全面揭秘> <什么是mmap?零拷贝?DMA?> <Linux C语言:用零拷贝技术实现TCP代理(源代码+测试服务端客户端代码 ...
续说零拷贝(Zero-Copy) - DMA技术
如果想理解Kafaka为什么这么快,得先看DMA是什么. DMA: 无论 I/O 速度如何提升,比起 CPU,总还是太慢.SSD 硬盘的 IOPS 可以到 2 万.4 万,但是我们 CPU 的主频有 ...
linux dma 拷贝内存数据_原来 8 张图，就可以搞懂「零拷贝」了
前言磁盘可以说是计算机系统最慢的硬件之一,读写速度相差内存 10 倍以上,所以针对优化磁盘的技术非常的多,比如零拷贝.直接 I/O.异步 I/O 等等,这些优化的目的就是为了提高系统的吞吐量,另外操 ...

CUDA零拷贝内存（zerocopy memory）

CUDA零拷贝内存（zerocopy memory）相关推荐

最新文章

热门文章