广东网站seo策划,苏州高端网页设计建设,国际最好的摄影作品网站,微信商城网站开发常量内存
__constant__ 是 CUDA 中用于声明常量内存变量的关键字#xff0c;其作用是在设备的常量内存#xff08;constant memory#xff09;中分配内存空间。这种内存适用于少量、在核函数中频繁读取但不改变的全局常量数据。
特性描述存储位置GPU 的常量内存#xff0…常量内存
__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));kernel1, 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); // 在默认流中记录开始时间myKernel100, 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);kernel1blocks, threads, 0, stream1(); // 在 stream1 中执行
cudaEventRecord(event, stream1); // 在 stream1 的尾部记录事件cudaStreamWaitEvent(stream2, event, 0); // 让 stream2 等待事件完成
kernel2blocks, 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 完成。
