网站导航
> 文章列表 > 第七天 CUDA Stream

第七天 CUDA Stream

网友: 文章列表 2024-03-21 23:48:06

第七天 CUDA Stream

CUDA Stream的理解

CUDA的程序一般需要处理海量的数据，内存带宽经常会成为主要的瓶颈

在Stream的帮助下，CUDA程序可以有效地将内存读取和数值运算并行，从而提升数据的吞吐量

由于GPU和CPU不能直接读取对方的内存，CUDA程序一般会有一下三个步骤：

  1. 将数据从CPU内存转移到GPU内存
  2. GPU进行运算并将结果保存在GPU内存
  3. 将结果从GPU内存拷贝到CPU内存

如果不做特别处理，那么CUDA会默认只使用一个Stream（Default Stream）

在这种情况下，必须等一步完成了才能进行下一步， 为此可以看出Stream的作用。

CUDA流表示一个GPU操作队列，该队列中的操作将以添加到流中的先后顺序而依次执行

可以将一个流看做是GPU上的一个任务，不同任务可以并行执行

使用CUDA流，首先要选择一个支持设备重叠（Device Overlap）功能的设备

支持设备重叠功能的GPU能够在执行一个CUDA核函数的同时，还能在主机和设备之间执行复制数据操作

支持重叠功能的设备的这一特性很重要，可以在一定程度上提升GPU程序的执行效率

一般情况下，CPU内存远大于GPU内存，对于数据量比较大的情况，不可能把CPU缓冲区中的数据一次性传输给GPU，需要分块传输

如果能够在分块传输的同时，GPU也在执行核函数运算，这样的异步操作，就用到设备的重叠功能，能够提高运算性能

CUDA Stream执行流程

  1. stream是一个流句柄，可以当做是一个队列
    • cuda执行器从stream中一条条的读取并执行指令
    • 例如cudaMemcpyAsync函数等同于向stream这个队列中加入一个cudaMemcpy指令并排队
    • 使用到了stream的函数，便立即向stream中加入指令后立即返回，并不会等待指令执行结束
    • 通过cudaStreamSynchronize函数，等待stream中所有指令执行完毕，也就是队列为空
  2. 当使用stream时，要注意
    • 由于异步函数会立即返回，因此传递进入的参数要考虑其生命周期，应确认函数调用结束后再做释放
  3. 还可以向stream中加入Event，用以监控是否到达了某个检查点
    • cudaEventCreate，创建事件
    • cudaEventRecord，记录事件，即在stream中加入某个事件，当队列执行到该事件后，修改其状态
    • cudaEventQuery，查询事件当前状态
    • cudaEventElapsedTime，计算两个事件之间经历的时间间隔，若要统计某些核函数执行时间，请使用这个函数，能够得到最准确的统计
    • cudaEventSynchronize，同步某个事件，等待事件到达
    • cudaStreamWaitEvent，等待流中的某个事件
  4. 默认流，对于cudaMemcpy等同步函数，其等价于执行了
    • cudaMemcpyAsync(… 默认流) 加入队列
    • cudaStreamSynchronize(默认流) 等待执行完成
    • 默认流与当前设备上下文类似，是与当前设备进行的关联
    • 因此，如果大量使用默认流，会导致性能低下

代码示例


// CUDA运行时头文件
#include <cuda_runtime.h>#include <stdio.h>
#include <string.h>#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){if(code != cudaSuccess){    const char* err_name = cudaGetErrorName(code);    const char* err_message = cudaGetErrorString(code);  printf("runtime error %s:%d  %s failed. \\n  code = %s, message = %s\\n", file, line, op, err_name, err_message);   return false;}return true;
}int main(){int device_id = 0;checkRuntime(cudaSetDevice(device_id));cudaStream_t stream = nullptr;checkRuntime(cudaStreamCreate(&stream));// 在GPU上开辟空间float* memory_device = nullptr;checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float)));// 在CPU上开辟空间并且放数据进去，将数据复制到GPUfloat* memory_host = new float[100];memory_host[2] = 520.25;checkRuntime(cudaMemcpyAsync(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice, stream)); // 异步复制操作，主线程不需要等待复制结束才继续// 在CPU上开辟pin memory,并将GPU上的数据复制回来 float* memory_page_locked = nullptr;checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float)));checkRuntime(cudaMemcpyAsync(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost, stream)); // 异步复制操作，主线程不需要等待复制结束才继续checkRuntime(cudaStreamSynchronize(stream));printf("%f\\n", memory_page_locked[2]);// 释放内存checkRuntime(cudaFreeHost(memory_page_locked));checkRuntime(cudaFree(memory_device));checkRuntime(cudaStreamDestroy(stream));delete [] memory_host;return 0;
}
网络标签：

相关问题