CUDA优化方案—stream的使用

---

一、 什么是CUDA stream

CUDA的stream流，类似我们经常使用CPU时开多线程。

• 当我们使用GPU进行计算时，如果我们没有主动开启stream流，GPU会自动创建默认流来执行核函数，默认流和CPU端的计算是同步的。（也即在CPU执行任务过程中，必须等GPU执行完核函数后，才能继续往下执行）
• 当我们使用GPU进行计算时，我们可以主动开启多个stream流，类似CPU开启多线程。我们可以将大批量文件读写分给多个流去执行，或者用不同的流分别计算不同的核函数。开启的多个流之间是异步的，流与CPU端的计算也是异步的。所以我们需要注意加上同步操作。
  值得注意的是，受PCIe总线带宽的限制，当一个流在进行读写操作时，另外一个流不能同时进行读写操作，但是其他流可以进行数值计算任务。这个有点类似与CPU中的流水线机制。
  

二、CUDA stream API介绍

• 创建一个stream
cudaStream_t stream;
cudaStreamCreate(&stream);
• 将host数据拷贝到device
cudaMemcpyAsync(dst, src, size, type, stream)
• kernel在流中执行
kernel_name<<<grid, block, stream>>>(praments);
• 同步和查询
cudaError_t cudaStreamSynchronize(cudaStream_t stream)
cudaError_t cudaStreamQuery(cudaStream_t stream);
• 销毁流
cudaError_t cudaStreamDestroy(cudaStream_t stream)

三、stream使用示例

//创建多个流
cudaStream_t stream[3];
for(int i=0; i<3; i++)
	cudaStreamCreat(&stream[i]);

float* hostPtr;
cudaMallocHost(&hostPtr, 3*size);
......
//开启了三个流，将数据分成三份分别交给了三个流进行计算
for(int i=0; i<3; i++){
	//从CPU内存复制数据到GPU显存
	cudaMemcpyAsync(DevPtr + i*size, hostPtr + i*size, size, cudaMemcpyHostToDevice, stream[i]);
	//执行Kernel函数
	Mykernel<<<grid, block, 0, stream[i]>>>(outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
	//将结果拷贝回CPU内存
	cudaMemcpyAsync(hostPtr + i*size, outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
}
//同步流
for(int i=0; i<3; i++)
	cudaStreamSynchronise(stream[i]);
//销毁流
for(int i=0; k<3; i++)
	cudaStreamDestroy(stream[i]);
	
}

四、CUDA 多个stream测试—CUDA Event API

主要用于标记stream执行过程中特定的点，我们可以插入event到stream中，查询该event是否完成。只有当该event标记的stream位置的所有操作都被执行完毕，该event才算完成。关联到默认stream上的event则对所有的stream有效

API
• 创建event
cudaEvent_t event
cudaError_t cudaEventCreate(cudaEvent_t* event);
• 将event插入流中(stream=0为默认流)
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);
• 销毁event
cudaError_t cudaEventDestroy(cudaEvent_t event);
• 阻塞host,等待event执行完成
cudaError_t cudaEventSynchronize(cudaEvent_t event);
• 查询stream是否结束，该函数不会阻塞host
cudaError_t cudaEventQuery(cudaEvent_t event);
• 进阶同步函数
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event)

// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// record start event（如果没指明stream,则为默认流）
cudaEventRecord(start);
// 执行 kernel
kernel<<<grid, block>>>(argulist);
// record stop event
cudaEventRecord(stop);

// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);

单线程内，编译加上–default-stream per-thread 后默认流的执行是异步的，显式流的执行是异步的
nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread
多线程下，正常编译，多个线程共享一个默认流
nvcc ./pthread_test.cu -o pthreads_legacy
加上–default-stream per-thread后，每个线程都有一个默认流
nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread