做视频的软件模板下载网站有哪些内容百度广告收费
常量内存
__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);
 
-  
cudaStream_t是 CUDA 中的流对象类型,表示一个指令队列。 -  
默认情况下,所有核函数、内存操作都在 默认流(stream 0) 上执行,所有操作是顺序同步执行的。
 -  
使用多个流,可以让不同核函数或操作并发执行,前提是它们之间没有依赖关系。
 
cudaEventRecord(event, stream1);
 
 在 stream1 上记录事件
-  
表示:当 stream1 中所有之前的操作完成时,该事件被“标记为完成”。
 -  
也就是说,
event表示kernel1执行完成的时间点。 
cudaStreamWaitEvent(stream2, event, 0);
 
stream2 等待 event 完成
-  
这句的含义是:stream2 中后续的所有操作都要等到
event完成之后才能执行。 -  
所以此处,stream2 中的 kernel 将会等待
kernel1执行完。 
虽然 stream1 和 stream2 原本是可以并发执行的,
但因为用 cudaStreamWaitEvent 明确加了依赖关系,所以 kernel2 会等 kernel1 完成。
