|
|
发表于 5 天前
|
查看: 20 |
回复: 0
一、CUDA中的事件
在众多编程语言或框架中，你可能都接触过“事件”这个概念。在 CUDA 编程中，事件同样扮演着核心角色，但其含义更贴近字面本身。它本质上是一种标志，用于密切监视设备的执行进度，是一个强大的同步工具。
同时，CUDA 事件也是执行精确计时的利器。它允许应用程序在代码的任何位置异步地记录事件，并查询这些事件何时完成。当事件之前的所有任务（或者可选地，特定流中的所有命令）都已完成时，该事件即被标记为完成。对于在流零（即默认流）中记录的事件，则是在所有流的所有先前任务和命令完成后才会完成。
因此，在 CUDA 编程中，事件身兼两职：既是高效的同步工具，也是测量执行时间的精密仪器。
二、CUDA事件主要应用场景
理解了事件的基本定义后，它在实际编程中有哪些用武之地呢？
-
性能测量
事件可以精确测量内核函数的执行耗时，也常被用于测量涉及多个流、计算与数据传输重叠等复杂操作的整体性能。
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);
-
流的同步
事件更重要的用途是控制流之间的依赖关系。这包括多流之间的同步、不同计算阶段的协调、多GPU间的数据依赖，以及在动态工作流中进行控制。
__global__ void foo1(char *A){
*A = 0x1;
}
__global__ void foo2(char *B){
printf("%d\n", *B); // *B == *A == 0x1 assuming foo2 waits for foo1
// to complete before launching
}
cudaMemcpyAsync(B, input, size, stream1); // Aliases are allowed at
// operation boundaries
foo1<<<1,1,0,stream1>>>(A); // allowing foo1 to access A.
cudaEventRecord(event, stream1);
cudaStreamWaitEvent(stream2, event);
foo2<<<1,1,0,stream2>>>(B);
cudaStreamWaitEvent(stream3, event);
cudaMemcpyAsync(output, B, size, stream3); // Both launches of foo2 andcudaMemcpy (which both read)
// wait for foo1 (which writes) to complete before proceeding
注：上面代码来自CUDA官网
三、流和流同步
虽然事件常用于流同步，但它与常规的流同步操作（如 cudaStreamSynchronize）有显著区别：
- 同步机制更灵活：事件可以实现跨多个流的精细同步控制，而流同步通常作用于单个流的内部。
- 控制精度更高：事件可以标记和等待一个具体的“点”（如某个内核完成），而流同步则需等待整个流的所有任务结束。
- 开销更低：在某些场景下，事件同步的开销低于完全同步一个流。
- 应用更广泛：事件除了同步，还具备计时功能，这是流同步不具备的。
通过这样的对比，能加深对两者适用场景的理解。
四、例程
结合以上理论，我们来看一个综合性的实战例程。它演示了如何利用事件进行多流间的同步，并对整个操作流程进行精确计时。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
__global__ void kernelFunc(float* data, int num, float factor){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < num) {
data[idx] = data[idx] * factor + idx * 0.001f;
}
}
int main(){
const int N = 1 << 20; // 100w
const int numStreams = 4;
const int chunkSize = N / numStreams;
const size_t chunkBytes = chunkSize * sizeof(float);
const size_t totalBytes = N * sizeof(float);
printf("mul stream sync test...\n");
float *hData = NULL;
cudaMallocHost(&hData, totalBytes);
for (int i = 0; i < N; i++) {
hData[i] = (float)rand() / RAND_MAX;
}
float *dData=NULL;
cudaMalloc(&dData, totalBytes);
cudaStream_t streams[numStreams];
for (int i = 0; i < numStreams; i++) {
cudaStreamCreate(&streams[i]);
}
//1 cacl time
cudaEvent_t startEvent, stopEvent;
cudaEvent_t kernelEvents[numStreams];
cudaEventCreate(&startEvent);
cudaEventCreate(&stopEvent);
for (int i = 0; i < numStreams; i++) {
cudaEventCreateWithFlags(&kernelEvents[i], cudaEventDisableTiming);
}
cudaEventRecord(startEvent, 0);
int threadsPerBlock = 256;
int blocksPerChunk = (chunkSize + threadsPerBlock - 1) / threadsPerBlock;
for (int i = 0; i < numStreams; i++) {
int offset = i * chunkSize;
cudaMemcpyAsync(&dData[offset], &hData[offset],
chunkBytes, cudaMemcpyHostToDevice,
streams[i]);
kernelFunc << <blocksPerChunk, threadsPerBlock, 0, streams[i] >> >(
&dData[offset], chunkSize, (float)(i + 1) * 0.5f);
cudaGetLastError();
cudaEventRecord(kernelEvents[i], streams[i]);
cudaMemcpyAsync(&hData[offset], &dData[offset],
chunkBytes, cudaMemcpyDeviceToHost,
streams[i]);
}
for (int i = 0; i < numStreams; i++) {
cudaStreamSynchronize(streams[i]);
}
// record finish timepoint
cudaEventRecord(stopEvent, 0);
cudaEventSynchronize(stopEvent);
float totalTime = 0.f;
cudaEventElapsedTime(&totalTime, startEvent, stopEvent);
printf("sum time: %.3f ms\n", totalTime);
//2 stream depended
printf("\n stream sync :\n");
cudaEvent_t syncEvent;
cudaEventCreate(&syncEvent);
kernelFunc << <blocksPerChunk, threadsPerBlock, 0, streams[0] >> >(
dData, chunkSize, 2.0f);
cudaEventRecord(syncEvent, streams[0]);
for (int i = 1; i < numStreams; i++) {
cudaStreamWaitEvent(streams[i], syncEvent, 0);
kernelFunc << <blocksPerChunk, threadsPerBlock, 0, streams[i] >> >(
&dData[i * chunkSize], chunkSize, 1.5f);
}
for (int i = 0; i < numStreams; i++) {
cudaStreamSynchronize(streams[i]);
}
printf("sync finish \n");
printf("\n event query:\n");
cudaEvent_t queryEvent;
cudaEventCreate(&queryEvent);
kernelFunc << <blocksPerChunk, threadsPerBlock >> >(dData, chunkSize, 1.0f);
cudaEventRecord(queryEvent, 0);
int maxChecks = 100;
int checkCount = 0;
while (cudaEventQuery(queryEvent) == cudaErrorNotReady) {
checkCount++;
if (checkCount < maxChecks) {
int dummy = 0;
for (int j = 0; j < 1000; j++) {
dummy += j;
}
}
else {
cudaEventSynchronize(queryEvent);
break;
}
}
printf("event query count: %d\n", checkCount);
cudaEventDestroy(syncEvent);
cudaEventDestroy(queryEvent);
cudaEventDestroy(startEvent);
cudaEventDestroy(stopEvent);
for (int i = 0; i < numStreams; i++) {
cudaEventDestroy(kernelEvents[i]);
cudaStreamDestroy(streams[i]);
}
cudaFree(dData);
cudaFreeHost(hData);
cudaDeviceReset();
printf("\n all finish！\n");
return 0;
}
这段代码主要实现了多流间的同步，并计算了整体操作时间。如果你有前面的 [C/C++](https://yunpan.plus/f/25-1) 和 CUDA 基础，理解起来应该不难。如果遇到不清楚的地方，最好的方法是上机运行，通过添加打印日志来观察程序的执行流程。
五、总结
在之前学习 CUDA 流时，我们简要提到了事件。本文则对事件进行了深入展开，详细阐述了其工作原理、核心应用场景，并通过实际例程让大家更直观地理解事件的运行机制。掌握好事件这一工具，对于编写高效、协调的 [高性能计算](https://yunpan.plus/f/15-1) 程序至关重要。
希望这篇关于 CUDA 事件的实战解析对你有帮助。想了解更多 并行编程 或 GPU计算 相关的深度内容，欢迎持续关注 云栈社区 的更新。 |
上一篇：跨界融合与个人品牌：Dan Koe对一人企业未来的核心洞察下一篇：NotebookLM进阶教程：利用AI深度解析YouTube视频内容作为个人知识库
|
