CUDA学习（十四）

回调：
运行时提供了一种方法，通过cudaStreamAddCallback（）将任何一个回调插入到一个流中。 回调是在回调完成之前，一旦所有发送到流的命令都在主机上执行的函数。 一旦在回调完成之前在所有流中发出的所有前面的任务和命令都执行流0中的回调。
以下代码示例在向每个流发出主机到设备内存副本，内核启动和设备到主机内存副本之后，将回调函数MyCallback添加到两个流中的每一个。 每个设备到主机内存副本完成后，回调将在主机上开始执行。

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data) {
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice,
        stream[i]);
    MyKernel << <100, 512, 0, stream[i] >> >(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost,
        stream[i]);
    cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0);
}

在回调完成之前，在回调之后，在流中发出的命令（或者如果回调被发送给流0，则发送到任何流的所有命令）不会开始执行。 cudaStreamAddCallback（）的最后一个参数保留给将来使用.
回调不能直接或间接地进行CUDA API调用，因为如果这样的调用导致死锁，它可能最终会等待自己.
流优先权:
可以使用cudaStreamCreateWithPriority（）在创建时指定流的相对优先级。 可以使用cudaDeviceGetStreamPriorityRange（）函数获取允许的优先级范围，以[最高优先级，最低优先级]排序。 在运行时，当低优先级方案中的块结束时，优先级较高的流中的等待块将被安排在其位置上。
以下代码示例获取当前设备允许的优先级范围，并创建具有最高和最低可用优先级的流：

// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

事件：
运行时还提供了一种方法来密切监视设备的进度，以及执行准确的计时，让应用程序在程序的任何时间点异步记录事件，并查询这些事件何时完成。 当事件之前的所有任务（或者可选地，给定流中的所有命令）已经完成时，事件已经完成。 在所有流中的所有前面的任务和命令完成之后，流Zero中的事件完成。
创造和销毁：
以下代码示例创建两个事件：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
他们的销毁方式：

cudaEventDestroy(start);
cudaEventDestroy(stop);

经过的时间：
创建和销毁中创建的事件可用于按以下方式计时创建和销毁的代码示例：

cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,
        size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel << <100, 512, 0, stream[i] >> >
        (outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,
        size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

同步呼叫：
当调用同步函数时，在设备完成所请求的任务之前，控制权不会返回到主机线程。 宿主线程随后会通过主机线程执行任何其他CUDA调用之前，通过调用cudaSetDeviceFlags（）来指定主机线程是否产生，阻塞或旋转，并指定一些特定的标志。