CUDA内存模型

核函数性能并不只与线程束的执行有关。

CUDA内存模型概述

GPU和CPU内存模型的主要区别是，CUDA编程模型能将内存层次结构更好地呈现给用户，能让我们显示的控制它的行为。

对程序员来说，一般有两种类型的存储器：

可编程的：需要显示的控制哪些数据存放在可编程内存中
不可编程的：不能决定数据的存放位置，程序自动生成存放位置

CUDA内存模型提供了多种可编程的内存类型：

寄存器
共享内存
本地内存
常量内存
纹理内存
全局内存

图中所示为这些内存空间的层次结构，每个都有不同的作用域、生命周期和缓存行为。一个核函数中的线程都有自己私有的本地内存。一个线程块有自己的共享内存，对同一线程块内的所有线程可见，其内容持续线程块的整个生命周期。所有线程都可以访问全局内存。所有线程都能访问的只读空间有：常量内存和纹理内存。对一个应用程序来说，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

全局内存（Global Memory）：

特性：全局内存是GPU中所有线程共享的内存空间，可以被所有线程访问。
优点：全局内存容量较大，适合存储大量数据。
缺点：全局内存访问延迟较高，需要通过内存控制器访问，因此访问速度相对较慢。
适用场景：适用于需要在不同线程之间共享数据的情况，但对访问速度要求不高的情况。

常量内存（Constant Memory）：

特性：常量内存是只读的内存空间，所有线程都可以读取其中的数据，但不能写入。
优点：常量内存具有高速缓存，访问速度较快。
缺点：常量内存容量有限（通常为64KB），且只能在程序启动时初始化。
适用场景：适用于存储只读数据，且需要高速访问的情况，如常数表、查找表等。

纹理内存（Texture Memory）：

特性：纹理内存是一种特殊的只读内存，具有缓存和插值功能，适合于图像处理等需要纹理特性的应用。
优点：纹理内存具有缓存和插值功能，可以提高某些类型数据的访问速度。
缺点：容量相对较小，不能直接写入。
适用场景：适用于需要纹理特性的数据访问，如图像处理、光线追踪等。

寄存器

寄存器是GPU上运行速度最快的内存空间。核函数中声明的变量通常储存在寄存器中。寄存器对每个线程来说是私有的。与核函数生命周期相同。

寄存器和线程不是一对一的，费米架构一个线程可以有63个寄存器，开普勒架构可以有255个寄存器。

如果，一个核函数使用了超过硬件限制数量的寄存器，则会用本地内存替代。这种寄存器溢出的情况会降低性能。

本地内存

核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。

在编译时，使用未知索引引用的本地数组
可能会占用大量寄存器空间的较大本地结构体或数组
任何不满足核函数寄存器限定条件的变量

共享内存

在核函数中使用如下修饰符修饰的变量存放在共享内存中

__shared__

每个SM都有一定数量的由线程块分配的共享内存。共享内存在核函数的范围内声明，其生命周期伴随着整个线程块。

共享内存是线程之间相互通信的基本方式。访问共享内存必须同步使用如下调用，

void __syncthreads();

SM中的一级缓存和共享内存都使用64KB的片上内存，通过静态划分，在运行时也可以进行动态配置。

片上内存（On-chip Memory）是指集成在芯片（例如GPU、CPU等处理器）内部的内存单元。与传统的外部内存（如RAM）相比，片上内存具有更低的访问延迟和更高的带宽，因为它们与处理器核心更加接近，减少了数据传输的距离和延迟。

常量内存

常量内存保存在设备内存中，每个SM都有自己专门用于缓存常量内存的缓存。

常量变量用以下修饰符修饰：

__constant__

常量变量必须在全局空间内和所有核函数之外进行声明。常量内存是静态声明的，并对同一编译单元内的所有核函数可见。

核函数只能从常量内存中读取数据，因此，常量内存必须在主机端由以下函数初始化

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);

将count个字节从src指向的内存复制到symbol指向的内存中，这个变量存放在设备的全局内存或常量内存中。

线程束中的所有线程从相同的地址读取数据时，常量内存表现最好。比如，所有线程都需要读取一个公式的系数，对不同的数据根据同样的系数做相同的计算，这时，系数存放在常量内存中最好。

纹理内存

特性：纹理内存是一种特殊的只读内存，具有缓存和插值功能，适合于图像处理等需要纹理特性的应用。
优点：纹理内存具有缓存和插值功能，可以提高某些类型数据的访问速度。
缺点：容量相对较小，不能直接写入。
适用场景：适用于需要纹理特性的数据访问，如图像处理、光线追踪等。

纹理特性的数据访问指的是使用CUDA中的纹理内存（Texture Memory）来访问数据的一种特殊方式。纹理内存在访问数据时具有缓存和插值等特性，适用于某些类型的数据访问需求，例如图像处理、光线追踪等。纹理内存的特性包括以下几个方面：

缓存：纹理内存具有缓存功能，可以将最近访问的数据缓存在内存中，以提高后续访问相同数据的速度。这种缓存机制有助于减少对全局内存的访问次数，从而提高访问速度。
插值：纹理内存支持插值功能，可以对访问的数据进行插值计算，以获取平滑的插值结果。这在图像处理等需要对像素进行插值的应用中特别有用，可以减少图像处理过程中的锯齿状边缘或伪影现象。
边界处理：纹理内存支持边界处理功能，可以在访问超出边界的数据时进行特定的处理，例如将边界外的像素值设置为固定值或通过重复边界像素值来填充。
硬件优化：纹理内存的访问方式经过硬件优化，可以提高数据访问的效率，并充分利用GPU的并行计算能力。

全局内存

全局内存是GPU中最大、延迟最高且最常使用的内存。

可以使用如下修饰符静态声明一个全局内存变量

__device__

在主机端使用cudaMalloc函数分配全局内存，使用cudaFree函数释放全局内存。全局内存可以存在于应用程序的整个生命周期中，可以被所有核函数的所有线程访问。从多个线程访问全局变量时，由于线程之间不能跨线程块同步，不同线程块内的多个线程并发地修改全局内存变量可能会出现问题。

GPU缓存

与CPU缓存一样，GPU缓存是不可编程的内存。在GPU上有四种缓存

一级缓存
二级缓存
只读常量缓存
只读纹理缓存

每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级缓存和二级缓存都被用来存储本地内存和全局内存中的数据。

实践

静态声明一个全局变量，并且在一个核函数中修改值，在主机端获取修改后的值。

__device__ float num;

__global__ void test_num() {
    printf("the num is %f\n", num);
    num += 0.3;
}

int main(){
    float a = 3.14;

    cudaMemcpyToSymbol(num, &a, sizeof(float));
    test_num << <1, 1 >> > ();

    cudaMemcpyFromSymbol(&a, num, sizeof(float));
    cudaDeviceReset();

    cout << a;
    return 0;
}

这里面cudaMemcpyToSymbol函数在VS里可能会标红