在核函数中，不同变量的生命周期

好记性不如烂笔头

---

前言

各变量生命周期

---

一、核函数参数（函数参数）

声明方式：

__global__ void kernel(float *dA, int lda) {
    // 核函数的代码
}

存储位置：全局内存（global memory）或寄存器
生命周期：
核函数参数的生命周期从核函数启动时开始，直到核函数执行完毕为止。
指针类型参数（如 float *dA）指向的是全局内存，这些指针在核函数中通过全局内存访问数据。指针本身可能被存储在寄存器中，但它指向的实际数据通常位于全局内存中。
标量参数（如 int lda）可以被存储在寄存器中，它们的生命周期只在核函数执行期间有效。

二、局部变量（Local Variables）

局部变量是声明在核函数中的变量，它们的生命周期和存储位置依赖于其具体类型和大小。

局部标量变量
声明方式：

__global__ void kernel() {
    int tmp = 0;  // 局部标量变量
    // 执行过程中的计算
}

存储位置：寄存器（register）或本地内存（local memory）
生命周期：
寄存器存储：对于小型的局部变量，通常会存储在寄存器中。寄存器中的变量生命周期只在当前线程的执行过程中有效。每个线程都有自己的私有寄存器空间，因此局部变量在每个线程中是独立的，且无法跨线程访问。
本地内存存储：如果局部变量过大或寄存器不足，编译器会将这些变量存储在本地内存中（本地内存实际上是位于全局内存的一部分）。虽然这些变量在每个线程内都是私有的，但它们的生命周期依然只在核函数执行期间。

三、共享内存变量（Shared Memory Variables）

共享内存是在同一个线程块中的线程之间共享的数据。它们可以由所有属于同一个线程块的线程访问。

声明方式：

__global__ void kernel() {
    __shared__ float sharedData[32];  // 共享内存变量
    // 线程块内的所有线程可以访问 sharedData
}

存储位置：共享内存（shared memory）
生命周期：
共享内存变量的生命周期从线程块开始执行时创建，直到线程块执行完成时销毁。
这些变量只能在声明它们的线程块内访问，线程块外的其他线程块无法访问它们。
当核函数启动时，CUDA 将为每个线程块分配共享内存，并在所有线程块中的线程都执行完毕后释放共享内存。

四、全局内存变量（Global Memory Variables）

全局内存是最常见的存储空间，适合存储大规模的数据结构（如矩阵、数组等），并且可以在核函数间共享。

声明方式：

float *dA;
cudaMalloc(&dA, size);  // 在全局内存中分配空间

存储位置：全局内存（global memory）
生命周期：
全局内存变量的生命周期通常从通过 cudaMalloc 分配时开始，直到使用 cudaFree 释放为止。这意味着它们可以在多个核函数调用之间存在。
全局内存可以被任何核函数访问，无论哪个线程块、哪个网格，甚至可以被 CPU 端（host）代码访问。
这些变量在 GPU 上分配时，与其声明的核函数执行无关，除非显式释放它们，否则会一直占用全局内存资源。

五、常量内存变量（Constant Memory Variables）

常量内存是用于存储只读数据的内存空间，适合存储小型的只读数据，并且在整个核函数执行期间可以被所有线程访问。

声明方式：

__constant__ float constData[256];  // 声明在常量内存中的数据

存储位置：常量内存（constant memory）
生命周期：
常量内存变量的生命周期从它们被初始化时开始，直到程序执行结束或被释放为止。
常量内存是只读的，所有线程都可以读取这些变量，但不能修改它们。常量内存的生命周期通常较长，可以在多个核函数间共享。

六、本地内存变量（Local Memory Variables）

本地内存是每个线程私有的存储空间，但实际上是位于全局内存中。通常用于存储超出寄存器容量的大型局部变量或编译时无法确定大小的局部变量。

声明方式：

__global__ void kernel() {
    int large_array[1000];  // 局部数组，可能存储在本地内存中
}

存储位置：本地内存（local memory，实际上是全局内存的一部分）
生命周期：
本地内存变量的生命周期仅限于它们声明的核函数执行期间，类似于寄存器变量。但它们存储在全局内存中，因此每个线程的本地内存变量是独立的，无法被其他线程访问。
它们在核函数执行结束后会自动销毁。

七、纹理内存变量（Texture Memory Variables）

纹理内存是一种只读内存，适合处理二维或三维图像数据。纹理内存提供了硬件支持的过滤和缓存功能，适用于需要高效读取二维、三维数据的场景。

声明方式：

texture<float, 2, cudaReadModeElementType> tex;  // 声明一个纹理变量

存储位置：纹理内存（texture memory）
生命周期：
纹理内存变量的生命周期从它们绑定到内存时开始，直到解除绑定时结束。通常在 cudaBindTexture 函数调用时开始，在 cudaUnbindTexture 调用时结束。
纹理内存的生命周期可以跨越多个核函数调用，适用于多次需要读取同样数据的情况。

---

总结

记一下，方便随时回来看