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

CUDA —— 2.3、cuda静态全局变量device使用介绍（附：完整代码）

news 2025/7/12 11:16:54

限定符device

在CUDA中，我们可以在设备（GPU）上定义全局变量，这些变量在设备的整个生命周期内存在，并且可以被所有线程访问。这些变量通常使用__device__限定符声明，并且可以具有静态存储期（即它们在程序启动时初始化，在程序结束时销毁）。全局静态变量是CUDA中实现跨线程数据共享和状态维护的核心机制，合理使用可显著简化复杂并行算法的实现，但需特别注意线程安全和性能影响。
在CUDA中，从CUDA 6.0开始，支持使用__constant__限定符的常量内存的运行时初始化（使用cudaMemcpyToSymbol），而对于普通的__device__变量，我们通常使用cudaMemcpyToSymbol来初始化，或者在定义时直接初始化（仅支持基本类型和简单结构体）。全局设备变量在主机代码中不能直接访问，必须使用CUDA运行时API（如cudaMemcpyToSymbol和cudaMemcpyFromSymbol）进行读写。 多个文件定义同名的全局设备变量会导致链接错误，因此通常在一个.cu文件中定义，在其他文件中使用extern声明。

普通全局变量

__device__ int globalCounter;          // 整型全局变量
__device__ float deviceArray[1024];    // 全局数组

常量内存变量

__constant__ float constParams[16];    // 常量内存（只读，高速缓存）

核心特性与使用限制
在这里插入图片描述

运行效果及代码

在这里插入图片描述

#include "cuda_runtime.h"
#include "device_launch_parameters.h"#include <stdio.h>/* __device__定义为静态全局内存变量。定义时只能在函数外定义。 */
__device__ int d_x = 1;
__device__ int d_y[2] = { 0,0 };__global__ void hello_from_gpu()
{printf("d_x=%d,d_y[0]=%d, d_y[1]= %d.\n", d_x, d_y[0], d_y[1]);d_y[0] += d_x; d_y[1] += d_x;printf("d_x=%d,d_y[0]=%d, d_y[1]= %d.\n", d_x, d_y[0], d_y[1]);
}int main(void)
{// 选用gpu设备cudaSetDevice(0);// cuda创建启停事件对象cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);/* 将主机的数据传递给全局静态变量 */int h_y[2] = { 10,20 };cudaMemcpyToSymbol(d_y, h_y, sizeof(int) * 2);// 记录事件cudaEventRecord(start);//hello_from_gpu << <dim3(2,2), dim3(4,4) >> > ();	// 执行核函数hello_from_gpu << <1,1 >> > ();	// 执行核函数// 记录事件cudaEventRecord(stop);// 事件同步(直到设备完成该事件,等待核函数执行结束)cudaEventSynchronize(stop);/* 将全局静态变量数据传递给主机数据 */cudaMemcpyFromSymbol(h_y,d_y,sizeof(int) * 2);printf("hy{%d,%d}", h_y[0], h_y[1]);// 计算事件之间经过的时间float time;cudaEventElapsedTime(&time, start, stop);printf("\n耗时为: %g ms\n", time);// 销毁事件对象cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}