CUDA编程之Stream

‌1、基本概念‌

• ‌定义‌：CUDA流（Stream）是一组在GPU上按顺序执行的操作序列，不同流中的操作可并行或交错执行，以提升计算效率‌。

• ‌作用‌：通过异步操作和并发执行，实现数据传输与核函数执行的重叠，减少整体执行时间‌。

• ‌条件‌：使用CUDA流，首先要选择一个支持设备重叠（Device Overlap）功能的设备，支持设备重叠功能的GPU能够在执行一个CUDA核函数的同时，还能在主机和设备之间执行复制数据操作。

2、核心操作‌

• ‌创建与销毁‌：
  使用 cudaStreamCreate() 创建流，cudaStreamDestroy() 销毁流。支持创建多个流以实现并发‌。

  cudaStream_t stream‌[2];
  for (int i=0; i<2; ++i) 
      cudaStreamCreate(&stream[i]);
  // 销毁
  for (int i=0; i<2; ++i)
      cudaStreamDestroy(stream[i]);
  

3、同步机制

• ‌流内同步‌：cudaStreamSynchronize(stream) 等待指定流内所有操作完成‌。
• ‌全局同步‌：cudaDeviceSynchronize() 等待设备上所有流完成‌。

4. ‌性能优化‌

• ‌数据传输与计算重叠‌：
  使用 cudaMemcpyAsync() 异步传输数据，结合不同流中的核函数执行，隐藏内存复制延迟‌。

  // 示例：异步数据传输与核函数并发执行 
  cudaMemcpyAsync(devPtr, hostPtr, size, cudaMemcpyHostToDevice, stream); 
  kernel<<<grid, block, 0, stream>>>(devPtr);

• ‌多流并发‌：
  将独立任务分配到不同流中，利用GPU多任务处理能力提升吞吐量‌。
• ‌事件同步‌：通过 cudaEventRecord() 和 cudaStreamWaitEvent() 实现流间依赖控制‌。

  cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags = 0);

• ‌stream‌：需等待事件的目标流。
• ‌event‌：被等待的事件，需已通过 cudaEventRecord() 记录到另一个流中。
• ‌flags‌：保留参数，通常设为 0。

•  cudaStreamWaitEvent强制指定一个流（stream）等待某个事件（event）完成后再继续执行后续任务，从而实现流间同步‌。
• 示例：
  流 A 的核函数需等待流 B 的数据传输完成后才能执行。

  cudaEvent_t event;
  cudaEventCreate(&event);
  // 流 B 记录事件
  cudaMemcpyAsync(..., stream_B);
  cudaEventRecord(event, stream_B);
  // 流 A 等待事件
  cudaStreamWaitEvent(stream_A, event);
  kernel_A<<<..., stream_A>>>();

  此操作确保 kernel_A 仅在流 B 的数据传输完成后启动‌ 。

5. 代码实例

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void vector_add(float *a, float *b, float *c, int size){
	int tid = blockIdx.x*blockDim.x + threadIdx.x;
	if(tid < size){
		c[tid] = a[tid] + b[tid];
	}
}

int main() {

	//获取设备属性  
	cudaDeviceProp prop;
	int deviceID;
	cudaGetDevice(&deviceID);
	cudaGetDeviceProperties(&prop, deviceID);
 
	//检查设备是否支持重叠功能  
	if (!prop.deviceOverlap)
	{
		printf("No device will handle overlaps. so no speed up from stream.\n");
		return 0;
	}

	//创建一个CUDA流
	cudaStream_t stream;
	cudaStreamCreate(&stream);

	int size = 1000;
	float *a, *b, *c, *c_host;
	c_host = (float*)malloc(sizeof(float)*size);
	//设备内存分配
	cudaMalloc(&a, sizeof(float)*size);
	cudaMalloc(&b, sizeof(float)*size);
	cudaMalloc(&c, sizeof(float)*size);

	//初始化及数据传输
	float a_host[1000];
	float b_host[1000];
	int n = 0;
	for(n=0; n<size; n++){
		a_host[n] = 1.0;
		b_host[n] = 99.0;
	}
	cudaMemcpyAsync(a,a_host,sizeof(float)*size,cudaMemcpyHostToDevice, stream);
	cudaMemcpyAsync(b,b_host,sizeof(float)*size,cudaMemcpyHostToDevice, stream);
	//////////////////////////////////

	dim3 block(256); //每Block 256线程
	dim3 grid((size+256-1)/256);//计算所需Block数(这里根据size计算)
	vector_add<<<grid,block,0,stream>>>(a,b,c,size);
	
	//结果回传到主机内存
	cudaMemcpyAsync(c_host,c,sizeof(float)*size,cudaMemcpyDeviceToHost, stream);

  	//等待指定流内所有操作完成‌
	cudaStreamSynchronize(stream);

	printf("host data head:%.2f tail:%.2f\n", c_host[0], c_host[size-1]);

	cudaFree(a);
	cudaFree(b);
	cudaFree(c);
	free(c_host);

	cudaStreamDestroy(stream);

    return 0;
}