一、 什么是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