CUDA中下面的一些执行步骤能够互相异步并发执行:
主机上的计算
设备上的计算
主机到设备的数据传输
设备到主机的数据传输
一个设备中内存的数据传输
不同设备之间的数据传输
注意,上面的6个操作是互相之间能够并发,在每个单独项上是不能够并发的(除了第二项设备上可以通过流并发来进行设备上计算的并发),但并不是每个英伟达的显卡都支持上面所有的并发,因此使用之前需要查询显卡是否支持。
主机和设备之间的并发:
核函数的启动
一个设备内,内存的复制
从主机到设备复制的内存块大小小于等于64kb时
函数后缀上带有Async的内存复制
内存设置函数调用(比如将某块内存全部置零)
最后,可以通过将变量CUDA_LAUNCH_BLOCKING置1来将核函数的启动变成同步,这个特性只用于调试
核函数之间的并发:
在设备计算能力大于2.1时支持核函数的并发,这个特性可以通过查询变量concurrentKernels来知道,当这个变量为1时说明设备支持,另外需要注意以下几点:
一个cuda上下文中的核函数不能够和另一个cuda上下文中的核函数并行
当一个核函数中大量使用纹理或局部内存时,核函数之间能够并行的可能性就在降低
数据传输与核函数执行的并发:
通过查询变量asyncEngineCount来检查设备是否支持,当设备支持时其值应当大于0
数据传输并发(指数据从主机到设备和从设备到主机的并发):
查询变量asyncEngineCount,支持时,其值应当等于2
流
前面的并发操作都是通过流来管理的:流与流之间可以并发,但流内部是有顺序的
流的同步:
//等待所有的流都执行完
cudaDeviceSynchronize()
//等待指定流中的任务全部执行完
cudaStreamSynchronize(cudaStream_t stream)
//某个流等待某个事件执行完再执行,下面的例子中使用了这个函数:需要注意的是这个函数也是非阻塞的
cudaStreamWaitEvent(cudaStream_t stream,cudaEvent_t event,unsigned int flags)
//查询某个流中的任务是否完成
cudaTreamQuery(cudaStream_t stream)
上面是流的显式同步,而隐式同步有以下几点:
如果主机线程中涉及到以下几点的话,那么两个流中的命令是无法同步的:
主机分配锁页内存
设备分配内存
将设备内存全部置成指定的值(原话是:a device memory set)
将两个地址的数据复制到同一个设备内存(原话是:a memory copy between two addresses to the same device memory)
任何附加到默认流中的命令
流中添加回调函数:
cudaStreamAddCallback(cudaStream_t stream,cudaStreamCallback_t callback,void *userdata,unsigned int flags)
需要注意的是第三个参数,是用户传递给回调函数的参数,cudaStreamCallback_t的定义:
typedef void (CUDART_CB cudaStreamCallback_t) (cudaStream_t stream,cudaError_t status,void userData)
流的优先级,流在创建时可以指定优先级:涉及的函数如下:
cudaStreamCreateWaitPriority()
cudaDeviceGetStreamPriorityRange()
事件:事件有两个作用:
第一个是记录某个流中某个操作的运行事件,创建时使用cudaEventCreate(cudaEvent_t *event),记录时间点:
cudaEventRecord(cudaEvent_t event,cudaStream_t stream),记录某个流下的时间点,由于是在主机上打印事件信息,因此主机需要等待事件完成,
因此还需要加上事件同步:cudaEventSynchronize(cudaEvent_t event),一个标准的事件记录:
cudaEvent_t start,end;
cudaEventCreate(&start);
cudaEventCreate(&end);
cudaEventRecord(start,stream)
//some stream operation
cudaEventRecord(end event,stream)
cudaEventSynchronize(end)
float elaspsedTime;
cudaEventElapsedTime(&elapsedTime,start,end);
第二个是作为一个标记,放在某个流操作的后面,标记这个流的该操作完成:
创建时使用:cudaEventCreateWithFlags(cudaEvent_t *event,unsigned int flags)
flags应该指定为cudaEventDisableTiming;flags还有其他值,请参考具体的手册
同步调用;相关函数:
cudaSetDeviceFlags(unsigned int flags):相见手册
核函数的并发执行,在下面需要注意的是,流与流之间通过事件来同步,即一个流等待另一个流执行完某个操作后再进行下一步操作,这时,创建事件时,应该使用:
checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming));
使用时应该是:当前流等待上一个流运行后当前流再执行:
for (int i=0; i<nkernels; ++i)
{
clock_block<<<1,1,0,streams[i]>>>(&d_a[i], time_clocks);
total_clocks += time_clocks;
checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i]));
//make the last stream wait for the kernel event to be recorded
checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0));
}
/*
* Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
*
* Please refer to the NVIDIA end user license agreement (EULA) associated
* with this source code for terms and conditions that govern your use of
* this software. Any use, reproduction, disclosure, or distribution of
* this software and related documentation outside the terms of the EULA
* is strictly prohibited.
*
*/
//
// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to
// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced
// in CUDA 3.2.
//
// Devices of compute capability 1.x will run the kernels one after another
// Devices of compute capability 2.0 or higher can overlap the kernels
//
#include <stdio.h>
#include <helper_functions.h>
#include <helper_cuda.h>
// This is a kernel that does no real work but runs at least for a specified number of clocks
__global__ void clock_block(clock_t *d_o, clock_t clock_count)
{
unsigned int start_clock = (unsigned int) clock();
clock_t clock_offset = 0;
while (clock_offset < clock_count)
{
unsigned int end_clock = (unsigned int) clock();
// The code below should work like
// this (thanks to modular arithmetics):
//
// clock_offset = (clock_t) (end_clock > start_clock ?
// end_clock - start_clock :
// end_clock + (0xffffffffu - start_clock));
//
// Indeed, let m = 2^32 then
// end - start = end + m - start (mod m).
clock_offset = (clock_t)(end_clock - start_clock);
}
d_o[0] = clock_offset;
}
// Single warp reduction kernel
__global__ void sum(clock_t *d_clocks, int N)
{
__shared__ clock_t s_clocks[32];
clock_t my_sum = 0;
for (int i = threadIdx.x; i < N; i+= blockDim.x)
{
my_sum += d_clocks[i];
}
s_clocks[threadIdx.x] = my_sum;
syncthreads();
for (int i=16; i>0; i/=2)
{
if (threadIdx.x < i)
{
s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i];
}
syncthreads();
}
d_clocks[0] = s_clocks[0];
}
int main(int argc, char **argv)
{
int nkernels = 8; // number of concurrent kernels
int nstreams = nkernels + 1; // use one more stream than concurrent kernel
int nbytes = nkernels * sizeof(clock_t); // number of data bytes
float kernel_time = 10; // time the kernel should run in ms
float elapsed_time; // timing variables
int cuda_device = 0;
printf("[%s] - Starting...\n", argv[0]);
// get number of kernels if overridden on the command line
if (checkCmdLineFlag(argc, (const char **)argv, "nkernels"))
{
nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels");
nstreams = nkernels + 1;
}
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
cuda_device = findCudaDevice(argc, (const char **)argv);
cudaDeviceProp deviceProp;
checkCudaErrors(cudaGetDevice(&cuda_device));
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device));
if ((deviceProp.concurrentKernels == 0))
{
printf("> GPU does not support concurrent kernel execution\n");
printf(" CUDA kernel runs will be serialized\n");
}
printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n",
deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
// allocate host memory
clock_t *a = 0; // pointer to the array data in host memory
checkCudaErrors(cudaMallocHost((void **)&a, nbytes));
// allocate device memory
clock_t *d_a = 0; // pointers to data and init value in the device memory
checkCudaErrors(cudaMalloc((void **)&d_a, nbytes));
// allocate and initialize an array of stream handles
cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t));
for (int i = 0; i < nstreams; i++)
{
checkCudaErrors(cudaStreamCreate(&(streams[i])));
}
// create CUDA event handles
cudaEvent_t start_event, stop_event;
checkCudaErrors(cudaEventCreate(&start_event));
checkCudaErrors(cudaEventCreate(&stop_event));
// the events are used for synchronization only and hence do not need to record timings
// this also makes events not introduce global sync points when recorded which is critical to get overlap
cudaEvent_t *kernelEvent;
kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t));
for (int i = 0; i < nkernels; i++)
{
checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming));
}
//////////////////////////////////////////////////////////////////////
// time execution with nkernels streams
clock_t total_clocks = 0;
#if defined(__arm__) || defined(__aarch64__)
// the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks.
clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000));
#else
clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate);
#endif
cudaEventRecord(start_event, 0);
// queue nkernels in separate streams and record when they are done
for (int i=0; i<nkernels; ++i)
{
clock_block<<<1,1,0,streams[i]>>>(&d_a[i], time_clocks);
total_clocks += time_clocks;
checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i]));
// make the last stream wait for the kernel event to be recorded
checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0));
}
// queue a sum kernel and a copy back to host in the last stream.
// the commands in this stream get dispatched as soon as all the kernel events have been recorded
sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels);
checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1]));
// at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel
// in this sample we just wait until the GPU is done
checkCudaErrors(cudaEventRecord(stop_event, 0));
checkCudaErrors(cudaEventSynchronize(stop_event));
checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event));
printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f);
printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f);
printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f);
bool bTestResult = (a[0] > total_clocks);
// release resources
for (int i = 0; i < nkernels; i++)
{
cudaStreamDestroy(streams[i]);
cudaEventDestroy(kernelEvent[i]);
}
free(streams);
free(kernelEvent);
cudaEventDestroy(start_event);
cudaEventDestroy(stop_event);
cudaFreeHost(a);
cudaFree(d_a);
if (!bTestResult)
{
printf("Test failed!\n");
exit(EXIT_FAILURE);
}
printf("Test passed\n");
exit(EXIT_SUCCESS);
}