常量?jī)?nèi)存

__constant__ 是 CUDA 中用于聲明常量?jī)?nèi)存變量的關(guān)鍵字，其作用是在設(shè)備的常量?jī)?nèi)存（constant memory）中分配內(nèi)存空間。這種內(nèi)存適用于少量、在核函數(shù)中頻繁讀取但不改變的全局常量數(shù)據(jù)。

特性	描述
存儲(chǔ)位置	GPU 的常量?jī)?nèi)存（顯存中的一個(gè)小區(qū)域，通常為 64KB）
訪問(wèn)方式	所有線程共享，可高效廣播
修改方式	主機(jī)端用 cudaMemcpyToSymbol 修改，設(shè)備端不能修改
使用場(chǎng)景	所有線程頻繁訪問(wèn)同一組常量，例如卷積核、查找表等

1.聲明常量變量（在全局作用域）?

__constant__ int constData[256];

注意：這聲明了一個(gè)存放在 GPU 常量?jī)?nèi)存上的全局整型數(shù)組 constData，不可在設(shè)備代碼中修改。

2.在主機(jī)端初始化常量?jī)?nèi)存

int hostData[256] = { /* 初始化數(shù)據(jù) */ };
cudaMemcpyToSymbol(constData, hostData, sizeof(int) * 256);

?? 使用 cudaMemcpyToSymbol 是唯一向 __constant__ 內(nèi)存賦值的方式。不能直接寫(xiě) constData[i] = x;。

?對(duì)比其他類(lèi)型

類(lèi)型	可見(jiàn)性	訪問(wèn)速度	可否修改	使用方式
__device__	所有設(shè)備函數(shù)	較快	?（設(shè)備端）	通常用來(lái)存儲(chǔ)全局變量
__shared__	一個(gè)線程塊內(nèi)共享	非?？?/td>	?	線程塊內(nèi)共享存儲(chǔ)
__constant__	所有線程共享	非常快（廣播）	?（只讀）	頻繁讀取、所有線程訪問(wèn)的全局只讀變量

?完整示例代碼

#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];  // 使用常量?jī)?nèi)存做乘法
}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);// 把常量數(shù)據(jù)復(fù)制到 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） 是一種用于 性能測(cè)量、同步和流控制 的機(jī)制，它可以記錄 GPU 中某個(gè)時(shí)間點(diǎn)的狀態(tài)，用于：

• ?? 精確測(cè)量 GPU 上操作（如 kernel 執(zhí)行、內(nèi)存拷貝）耗時(shí)

• 🔁 在不同 CUDA 流（stream）中進(jìn)行同步控制

• 🧵 檢查某些異步操作是否完成

?CUDA 中的事件對(duì)象由 cudaEvent_t 表示，可以用來(lái)：

作用	描述
記錄時(shí)間點(diǎn)	通過(guò) cudaEventRecord() 在 GPU 某個(gè)位置插入事件
測(cè)量時(shí)間間隔	用 cudaEventElapsedTime() 計(jì)算兩個(gè)事件之間的時(shí)間（以毫秒為單位）
異步檢測(cè)完成	使用 cudaEventQuery() 檢查事件是否已完成
流同步	用 cudaEventSynchronize() 等待某個(gè)事件完成
跨流依賴(lài)控制	使用事件實(shí)現(xiàn)不同 stream 間的同步

常用 API?

函數(shù)	說(shuō)明
cudaEventCreate(&event)	創(chuàng)建事件對(duì)象
cudaEventRecord(event, stream)	在某個(gè) stream 中記錄事件（stream 可為 0，表示默認(rèn)流）
cudaEventSynchronize(event)	等待事件完成（阻塞主機(jī)直到事件觸發(fā)）
cudaEventElapsedTime(&ms, start, stop)	計(jì)算兩個(gè)事件之間的時(shí)間間隔（單位：毫秒）
cudaEventDestroy(event)	銷(xiāo)毀事件對(duì)象，釋放資源

典型使用場(chǎng)景：測(cè)量核函數(shù)運(yùn)行時(shí)間?

#include <cuda_runtime.h>
#include <iostream>__global__ void myKernel() {int tid = threadIdx.x + blockIdx.x * blockDim.x;if (tid < 10000) {// 模擬計(jì)算for (int i = 0; i < 1000; ++i) {}}
}int main() {cudaEvent_t start, stop;float elapsedTime;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);       // 在默認(rèn)流中記錄開(kāi)始時(shí)間myKernel<<<100, 128>>>();        // 啟動(dòng)核函數(shù)cudaEventRecord(stop, 0);        // 記錄結(jié)束時(shí)間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) 只是記錄了一個(gè)事件插入點(diǎn)

它 插入到當(dāng)前設(shè)備上的默認(rèn)流（stream 0）中，表示 “從這時(shí)開(kāi)始有事件要記錄”。

stop 為什么必須同步？

因?yàn)?#xff1a;

• 核函數(shù)是 異步執(zhí)行 的，也就是：

  • 主機(jī)代碼調(diào)用 kernel 時(shí)不會(huì)等待其執(zhí)行完成；

  • 繼續(xù)往下執(zhí)行 cudaEventRecord(stop, 0)；

  • 但是這行本身也只是把事件插入到流中，不意味著 kernel 已完成。

為了讓我們能正確讀取 kernel 執(zhí)行完的時(shí)間，必須同步 stop 事件：

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

跨流同步示例：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);cudaEvent_t event;
cudaEventCreate(&event);kernel1<<<blocks, threads, 0, stream1>>>();  // 在 stream1 中執(zhí)行
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 中的流對(duì)象類(lèi)型，表示一個(gè)指令隊(duì)列。

• 默認(rèn)情況下，所有核函數(shù)、內(nèi)存操作都在 默認(rèn)流（stream 0） 上執(zhí)行，所有操作是順序同步執(zhí)行的。

• 使用多個(gè)流，可以讓不同核函數(shù)或操作并發(fā)執(zhí)行，前提是它們之間沒(méi)有依賴(lài)關(guān)系。

cudaEventRecord(event, stream1);

?在 stream1 上記錄事件

• 表示：當(dāng) stream1 中所有之前的操作完成時(shí)，該事件被“標(biāo)記為完成”。

• 也就是說(shuō)，event 表示 kernel1 執(zhí)行完成的時(shí)間點(diǎn)。

cudaStreamWaitEvent(stream2, event, 0);

stream2 等待 event 完成

• 這句的含義是：stream2 中后續(xù)的所有操作都要等到 event 完成之后才能執(zhí)行。

• 所以此處，stream2 中的 kernel 將會(huì)等待 kernel1 執(zhí)行完。

雖然 stream1 和 stream2 原本是可以并發(fā)執(zhí)行的，

但因?yàn)橛?cudaStreamWaitEvent 明確加了依賴(lài)關(guān)系，所以 kernel2 會(huì)等 kernel1 完成。