Cuda Streams的概述（三）-- 并发内存拷贝

并发内存拷贝

首先，回顾一下CUDA的内存。

• 设备内存：
  • 通过cudaMalloc申请内存。
  • 不能分页
• 主机端可分页内存：
  • 默认申请方式（malloc， calloc，new）
  • 可以在OS中调入和调出。
• 主机端页锁内存：
  • 通过特殊的申请方式
  • 不能由OS调出

ALLOCATING PINNED MEMORY

cudaMallocHost(...) / cudaHostAlloc(...) //在主机端申请和释放pinned 内存。
cudaFreeHost(...) // 释放由cudaMallocHost or cudaHostAlloc 申请的内存。
cudaHostRegister(...) / cudaHostUnregister(...) //固定/取消pinned可分页内存（使其成为pinned内存），但是由于很慢所以不考虑。

为什么用pinned内存

• 可分页内存通过主机端CPU进行传输。
• pinned内存通过DMA传输。
  • 异步地由CPU释放
  • 实现更高百分比的峰值带宽

cudaMemcpy()
//在默认流里传输，
//同步：必须在返回之前完成

cudaMemcpyAsync(..., &stream) //在流中传输，同时立刻返回

为了实现并发

• 传输要在非默认流里
• 必须用异步拷贝
• 同时只能有一个方向的一个传输
• 在主机端的内存必须是pinned

PAGED MEMORY EXAMPLE

int *h_ptr, *d_ptr;

h_ptr=malloc(bytes);
cudaMalloc(&d_ptr,bytes);

cudaMemcpy(d_ptr,h_ptr,bytes,cudaMemcpyHostToDevice);

free(h_ptr); 
cudaFree(d_ptr);

PINNED MEMORY: EXAMPLE 1

int *h_ptr, *d_ptr;

cudaMallocHost(&h_ptr,bytes); 
cudaMalloc(&d_ptr,bytes);

cudaMemcpy(d_ptr,h_ptr,bytes,cudaMemcpyHostToDevice);

cudaFreeHost(h_ptr); 
cudaFree(d_ptr);

PINNED MEMORY: EXAMPLE 2

int *h_ptr, *d_ptr;

h_ptr=malloc(bytes); 
cudaHostRegister(h_ptr,bytes,0); 
cudaMalloc(&d_ptr,bytes);

cudaMemcpy(d_ptr,h_ptr,bytes,cudaMemcpyHostToDevice);

cudaHostUnregister(h_ptr); 
free(h_ptr); 
cudaFree(d_ptr);

CONCURRENCY EXAMPLES

同步：

cudaMemcpy(...);
foo<<<...>>>();

同一个流里的异步

cudaMemcpyAsync(...,stream1); 
foo<<<...,stream1>>>();

不同流里的异步

cudaMemcpyAsync(...,stream1); 
foo<<<...,stream2>>>();

总结

并发地执行内存拷贝只能：

• 内存的拷贝要在不同的非默认流里
• 拷贝使用主机端的pinned内存
• 调用异步API
• 在同一时刻，同一方向不能有另外一个内存拷贝运行。