当前位置：首页 > news >正文

cuda编程笔记（3）--常量内存与事件

news 来源：原创 2025/6/13 9:54:57

常量内存

__constant__ 是 CUDA 中用于声明常量内存变量的关键字，其作用是在设备的常量内存（constant memory）中分配内存空间。这种内存适用于少量、在核函数中频繁读取但不改变的全局常量数据。

特性	描述
存储位置	GPU 的常量内存（显存中的一个小区域，通常为 64KB）
访问方式	所有线程共享，可高效广播
修改方式	主机端用 `cudaMemcpyToSymbol` 修改，设备端不能修改
使用场景	所有线程频繁访问同一组常量，例如卷积核、查找表等

1.声明常量变量（在全局作用域）

__constant__ int constData[256];

注意：这声明了一个存放在 GPU 常量内存上的全局整型数组 constData，不可在设备代码中修改。

2.在主机端初始化常量内存

int hostData[256] = { /* 初始化数据 */ };
cudaMemcpyToSymbol(constData, hostData, sizeof(int) * 256);

⚠️ 使用 cudaMemcpyToSymbol 是唯一向 __constant__ 内存赋值的方式。不能直接写 constData[i] = x;。

对比其他类型

类型	可见性	访问速度	可否修改	使用方式
`__device__`	所有设备函数	较快	✅（设备端）	通常用来存储全局变量
`__shared__`	一个线程块内共享	非常快	✅	线程块内共享存储
`__constant__`	所有线程共享	非常快（广播）	❌（只读）	频繁读取、所有线程访问的全局只读变量

完整示例代码

#include <iostream>
#include <cuda_runtime.h>__constant__ float coeffs[4];__global__ void kernel(float* data) {int i = threadIdx.x;data[i] = data[i] * coeffs[i % 4];  // 使用常量内存做乘法
}int main() {float hostCoeffs[4] = { 1.0f, 2.0f, 3.0f, 4.0f };float data[8] = { 10, 20, 30, 40, 50, 60, 70, 80 };float* devData;cudaMalloc(&devData, sizeof(data));cudaMemcpy(devData, data, sizeof(data), cudaMemcpyHostToDevice);// 把常量数据复制到 GPUcudaMemcpyToSymbol(coeffs, hostCoeffs, sizeof(hostCoeffs));kernel<<<1, 8>>>(devData);cudaMemcpy(data, devData, sizeof(data), cudaMemcpyDeviceToHost);for (int i = 0; i < 8; ++i) {std::cout << data[i] << " ";}std::cout << std::endl;cudaFree(devData);return 0;
}

事件

CUDA 的 事件（Event） 是一种用于 性能测量、同步和流控制 的机制，它可以记录 GPU 中某个时间点的状态，用于：

⏱️ 精确测量 GPU 上操作（如 kernel 执行、内存拷贝）耗时
🔁 在不同 CUDA 流（stream）中进行同步控制
🧵 检查某些异步操作是否完成

CUDA 中的事件对象由 cudaEvent_t 表示，可以用来：

作用	描述
记录时间点	通过 `cudaEventRecord()` 在 GPU 某个位置插入事件
测量时间间隔	用 `cudaEventElapsedTime()` 计算两个事件之间的时间（以毫秒为单位）
异步检测完成	使用 `cudaEventQuery()` 检查事件是否已完成
流同步	用 `cudaEventSynchronize()` 等待某个事件完成
跨流依赖控制	使用事件实现不同 stream 间的同步

常用 API

函数	说明
`cudaEventCreate(&event)`	创建事件对象
`cudaEventRecord(event, stream)`	在某个 stream 中记录事件（stream 可为 0，表示默认流）
`cudaEventSynchronize(event)`	等待事件完成（阻塞主机直到事件触发）
`cudaEventElapsedTime(&ms, start, stop)`	计算两个事件之间的时间间隔（单位：毫秒）
`cudaEventDestroy(event)`	销毁事件对象，释放资源

典型使用场景：测量核函数运行时间

#include <cuda_runtime.h>
#include <iostream>__global__ void myKernel() {int tid = threadIdx.x + blockIdx.x * blockDim.x;if (tid < 10000) {// 模拟计算for (int i = 0; i < 1000; ++i) {}}
}int main() {cudaEvent_t start, stop;float elapsedTime;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);       // 在默认流中记录开始时间myKernel<<<100, 128>>>();        // 启动核函数cudaEventRecord(stop, 0);        // 记录结束时间cudaEventSynchronize(stop);      // 等待 stop 事件代表的所有前序操作完成cudaEventElapsedTime(&elapsedTime, start, stop);std::cout << "Kernel execution time: " << elapsedTime << " ms" << std::endl;cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}

cudaEventRecord(start, 0) 只是记录了一个事件插入点

它 插入到当前设备上的默认流（stream 0）中，表示 “从这时开始有事件要记录”。

stop 为什么必须同步？

因为：

核函数是 异步执行 的，也就是：
- 主机代码调用 kernel 时不会等待其执行完成；
- 继续往下执行 cudaEventRecord(stop, 0)；
- 但是这行本身也只是把事件插入到流中，不意味着 kernel 已完成。

为了让我们能正确读取 kernel 执行完的时间，必须同步 stop 事件：

cudaEventSynchronize(stop); // 等待 stop 事件代表的所有前序操作完成

跨流同步示例：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);cudaEvent_t event;
cudaEventCreate(&event);kernel1<<<blocks, threads, 0, stream1>>>();  // 在 stream1 中执行
cudaEventRecord(event, stream1);             // 在 stream1 的尾部记录事件cudaStreamWaitEvent(stream2, event, 0);      // 让 stream2 等待事件完成
kernel2<<<blocks, threads, 0, stream2>>>();  // stream2 中的 kernel 等待 eventcudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);
cudaEventDestroy(event);