GPU存储器架构-- 全局内存本地内存寄存器堆共享内存常量内存纹理内存

Posted 2023-04-02 给算法爸爸上香

tags:

篇首语：本文由小常识网(cha138.com)小编为大家整理，主要介绍了GPU存储器架构-- 全局内存本地内存寄存器堆共享内存常量内存纹理内存相关的知识，希望对你有一定的参考价值。

上表表述了各种存储器的各种特性。作用范围栏定义了程序的哪个部分能使用该存储器。而生存期定义了该存储器中的数据对程序可见的时间。除此之外，Ll和L2缓存也可以用于GPU程序以便更快地访问存储器。
总之，所有线程都有一个寄存器堆，它是最快的。共享内存只能被块中的线程访问,但比全局内存块。全局内存是最慢的，但可以被所有的块访问。

全局内存

所有的块都可以对全局内存进行读写。该存储器较慢，但是可以从代码的任何地方进行读写。缓存可加速对全局内存的访问。所有通过cudaMalloc分配的存储器都是全局内存。下面的简单代码演示了如何从程序中使用全局内存:

#include <stdio.h>
#define N 5

__global__ void gpu_global_memory(int *d_a)

	// "array" is a pointer into global memory on the device
	d_a[threadIdx.x] = threadIdx.x;


int main()

	// Define Host Array
	int h_a[N];
	//Define device pointer	
	int *d_a;       
						
	cudaMalloc((void **)&d_a, sizeof(int) *N);
	// now copy data from host memory to device memory 
	cudaMemcpy((void *)d_a, (void *)h_a, sizeof(int) *N, cudaMemcpyHostToDevice);
	// launch the kernel 
	gpu_global_memory << <1, N >> >(d_a);  
	// copy the modified array back to the host memory
	cudaMemcpy((void *)h_a, (void *)d_a, sizeof(int) *N, cudaMemcpyDeviceToHost);
	printf("Array in Global Memory is: \\n");
	//Printing result on console
	for (int i = 0; i < N; i++) 
	
		printf("At Index: %d --> %d \\n", i, h_a[i]);
	

	return 0;

本地内存和寄存器堆

本地内存和寄存器堆对每个线程都是唯一的。寄存器是每个线程可用的最快存储器。当内核中使用的变量在寄存器堆中装不下的时候，将会使用本地内存存储它们，这叫寄存器溢出。请注意使用本地内存有两种情况:一种是寄存器不够了，-种是某些情况根本就不能放在寄存器中，例如对一个局部数组的下标进行不定索引的时候。基本上可以将本地内存看成是每个线程的唯一的全局内存部分。相比寄存器堆，本地内存要慢很多。虽然本地内存通过Ll缓存和L2缓存进行了缓冲，但寄存器溢出可能会影响你的程序的性能。
下面演示一个简单的程序:

#include <stdio.h>
#define N 5

__global__ void gpu_local_memory(int d_in)

	int t_local;    
	t_local = d_in * threadIdx.x;     
	printf("Value of Local variable in current thread is: %d \\n", t_local);


int main()


	printf("Use of Local Memory on GPU:\\n");
	gpu_local_memory << <1, N >> >(5);  
	cudaDeviceSynchronize();
	return 0;

代码中的t_local变量是每个线程局部唯一的，将被存储在寄存器堆中。用这种变量计算的时候，计算速度将是最快速的。

共享内存

共享内存位于芯片内部，因此它比全局内存快得多。(CUDA里面存储器的快慢有两方面，一个是延迟低，一个是带宽大。这里特指延迟低)，相比没有经过缓存的全局内存访问,共享内存大约在延迟上低100倍。同一个块中的线程可以访问相同的一段共享内存(注意:不同块中的线程所见到的共享内存中的内容是不相同的)，这在许多线程需要与其他线程共享它们的结果的应用程序中非常有用。但是如果不同步，也可能会造成混乱或错误的结果。如果某线程的计算结果在写入到共享内存完成之前被其他线程读取，那么将会导致错误。因此，应该正确地控制或管理内存访问。这是由_syncthreads()指令完成的，该指令确保在继续执行程序之前完成对内存的所有写入操作。这也被称为barrier。barrier 的含义是块中的所有线程都将到达该代码行，然后在此等待其他线程完成。当所有线程都到达了这里之后，它们可以一起继续往下执行。为了演示共享内存和线程同步的使用，我们这里给出一个计算MA的例子:

#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void gpu_shared_memory(float *d_a)

	// Defining local variables which are private to each thread
	int i, index = threadIdx.x;
	float average, sum = 0.0f;

	//Define shared memory
	__shared__ float sh_arr[10];
	sh_arr[index] = d_a[index];
	__syncthreads();    // This ensures all the writes to shared memory have completed

	for (i = 0; i<= index; i++) 
	 
		sum += sh_arr[i]; 
	
	average = sum / (index + 1.0f);
	d_a[index] = average; 


int main(int argc, char **argv)

	//Define Host Array
	float h_a[10];   
	//Define Device Pointer
	float *d_a;       
	
	for (int i = 0; i < 10; i++)
	
		h_a[i] = i;
	
	// allocate global memory on the device
	cudaMalloc((void **)&d_a, sizeof(float) * 10);
	// now copy data from host memory  to device memory 
	cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10, cudaMemcpyHostToDevice);
	
	gpu_shared_memory << <1, 10 >> >(d_a);
	// copy the modified array back to the host memory
	cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);
	printf("Use of Shared Memory on GPU:  \\n");
	//Printing result on console
	for (int i = 0; i < 10; i++) 
	
		printf("The running average after %d element is %f \\n", i, h_a[i]);
	
	return 0;

MA操作很简单，就是计算数组中当前元素之前所有元素的平均值，很多线程计算的时候将会使用数组中的同样的数据。这就是一种理想的使用共享内存的用例，这样将会得到比全局内存更快的数据访问。这将减少每个线程的全局内存访问次数，从而减少程序的延迟。共享内存上的数字或者变量是通过__shared__修饰符定义的。我们在本例中，定义了具有10个float元素的共享内存上的数组。通常，共享内存的大小应该等于每个块的线程数。因为我们要处理10个(元素)的数组，所以我们也将共享内存的大小定义成这么大。
下一步就是将数据从全局内存复制到共享内存。每个线程通过自己的索引复制一个元素，这样块整体完成了数据的复制操作，这样数据写到了共享内存中。在下一行，我们开始读取使用这个共享内存中的数组，但是在继续之前，我们应当保证所有(线程）都已经完成了它们的写入操作。所以，让我们使用__syncthreads()进行一次同步。
接着就是（每个线程)通过for循环，利用这些存储在共享内存中的值(读取后）计算(从第一个元素）到当前元素的平均值，并且将对应每个线程的结果存放到全局内存中的相应位置。

常量内存

CUDA程序员会经常用到另外一种存储器——常量内存，NVIDIA GPU 卡从逻辑上对用户提供了64KB的常量内存空间，可以用来存储内核执行期间所需要的恒定数据。常量内存对一些特定情况下的小数据量的访问具有相比全局内存的额外优势。使用常量内存也一定程度上减少了对全局内存的带宽占用。在本小节中，我们将看看如何在CUDA中使用常量内存。我们将用一个简单的程序进行a * x + b的数学运算，其中a,b都是常数，程序代码如下:

#include "stdio.h"
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining two constants
__constant__ int constant_f;
__constant__ int constant_g;
#define N	5

//Kernel function for using constant memory
__global__ void gpu_constant_memory(float *d_in, float *d_out) 

	//Thread index for current kernel
	int tid = threadIdx.x;	
	d_out[tid] = constant_f*d_in[tid] + constant_g;


int main() 

	//Defining Arrays for host
	float h_in[N], h_out[N];
	//Defining Pointers for device
	float *d_in, *d_out;
	int h_f = 2;
	int h_g = 20;
	// allocate the memory on the cpu
	cudaMalloc((void**)&d_in, N * sizeof(float));
	cudaMalloc((void**)&d_out, N * sizeof(float));
	//Initializing Array
	for (int i = 0; i < N; i++) 
	
		h_in[i] = i;
	
	//Copy Array from host to device
	cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
	//Copy constants to constant memory
	cudaMemcpyToSymbol(constant_f, &h_f, sizeof(int), 0, cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(constant_g, &h_g, sizeof(int));

	//Calling kernel with one block and N threads per block
	gpu_constant_memory << <1, N >> >(d_in, d_out);
	//Coping result back to host from device memory
	cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
	//Printing result on console
	printf("Use of Constant memory on GPU \\n");
	for (int i = 0; i < N; i++) 
	
		printf("The expression for input %f is %f\\n", h_in[i], h_out[i]);
	
	//Free up memory
	cudaFree(d_in);
	cudaFree(d_out);
	return 0;

常量内存中的变量使用__constant__ 关键字修饰。在之前的代码中，两个浮点数constant_f，constant_g 被定义成在内核执行期间不会改变的常量。需要注意的第二点是，使用__constant__ (在内核外面)定义好了它们后，它们不应该再次在内核内部定义。内核函数将用这两个常量进行一个简单的数学运算，在main 函数中，我们用一个特殊的方式将这两个常量的值传递到常量内存中。
在main 函数中，h_f， h_g两个常量在主机上被定义并初始化，然后将被复制到设备上的常量内存中。我们将用cudaMemcpyToSymbol函数把这些常量复制到内核执行所需要的常量内存中。该函数具有五个参数:第一个参数是(要写入的)目标，也就是我们刚才用__constant__ 定义过的h_f或者h_g常量;第二个参数是源主机地址;第三个参数是传输大小;第四个参数是写人目标的偏移量，这里是0;第五个参数是设备到主机的数据传输方向;最后两个参数是可选的，因此后面我们第二次cudaMemcpyToSymbol函数调用的时候省略掉了它们。

纹理内存

纹理内存是另外一种当数据的访问具有特定的模式的时候能够加速程序执行，并减少显存带宽的只读存储器。像常量内存一样，它也在芯片内部被cache缓冲。该存储器最初是为了图形绘制而设计的，但也可以被用于通用计算。当程序进行具有很大程度上的空间邻近性的访存的时候，这种存储器变得非常高效。空间邻近性的意思是，每个线程的读取位置都和其他线程的读取位置邻近。这对那些需要处理4个邻近的相关点或者8个邻近的点的图像处理应用非常有用。
通用的全局内存的cache将不能有效处理这种空间邻近性，可能会导致进行大量的显存读取传输。纹理存储被设计成能够利用这种访存模型，这样它只会从显存读取1次，然后缓冲掉，所以执行速度将会快得多。纹理内存支持2D和3D的纹理读取操作，在你的CUDA程序里面使用纹理内存可没有那么轻易，特别是对那些并非编程专家的人来说。我们将在本小节中为你解释一个如何通过纹理存储进行数组赋值的例子:

#include "stdio.h"
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>

#define NUM_THREADS 10
#define N 10
texture <float, 1, cudaReadModeElementType> textureRef;

__global__ void gpu_texture_memory(int n, float *d_out)

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	if (idx < n)
	
		float temp = tex1D(textureRef, float(idx));
		d_out[idx] = temp;
	


int main()

	//Calculate number of blocks to launch
	int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);
	//Declare device pointer
	float *d_out;
	// allocate space on the device for the result
	cudaMalloc((void**)&d_out, sizeof(float) * N);
	// allocate space on the host for the results
	float *h_out = (float*)malloc(sizeof(float) * N);
	//Declare and initialize host array
	float h_in[N];
	for (int i = 0; i < N; i++) 
	
		h_in[i] = float(i);
	
	//Define CUDA Array
	cudaArray *cu_Array;
	cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);
	//Copy data to CUDA Array
	cudaMemcpyToArray(cu_Array, 0, 0, h_in, sizeof(float)*N, cudaMemcpyHostToDevice);
	
	// bind a texture to the CUDA array
	cudaBindTextureToArray(textureRef, cu_Array);
	//Call Kernel	
  	gpu_texture_memory << <num_blocks, NUM_THREADS >> >(N, d_out);
	
	// copy result back to host
	cudaMemcpy(h_out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);
	printf("Use of Texture memory on GPU: \\n");
	for (int i = 0; i < N; i++)
	
		printf("Texture element at %d is : %f\\n",i, h_out[i]);
	
	free(h_out);
	cudaFree(d_out);
	cudaFreeArray(cu_Array);
	cudaUnbindTexture(textureRef);

通过“纹理引用”来定义一段能进行纹理拾取的纹理内存。纹理引用是通过texture<>类型的变量进行定义的。定义的时候，它具有3个参数:第一个是texture<>类型的变量定义时候的参数，用来说明纹理元素的类型。在本例中，是float类型;第二个参数说明了纹理引用的类型，可以是1D的，2D的，3D的。在本例中，是1D的纹理引用;第三个参数则是读取模式，这是一个可选参数，用来说明是否要执行读取时候的自动类型转换。请一定要确保纹理引用被定义成全局静态变量，同时还要确保它不能作为参数传递给任何其他函数。在这个内核函数中，每个线程通过纹理引用读取自己线程ID作为索引位置的数据,然后复制到d_out 指针指向的全局内存中。
在main函数中，定义并分配了内存和显存上的数组后，主机上的数组(中的元素）被初始化为0-9的值。本例中，你会第一次看到CUDA数组的使用。它们类似于普通的数组，但是却是纹理专用的。CUDA数组对于内核函数来说是只读的。但可以在主机上通过cudaMemcpyToArray函数写入，如同你在之前的代码中看到的那样。在cudaMemcpyToArray函数中，第二个和第三个参数中的0代表传输到的目标CUDA数组横向和纵向上的偏移量。两个方向上的偏移量都是О代表我们的这次传输将从目标CUDA数组的左上角（0，0）开始。CUDA数组中的存储器布局对用户来说是不透明的，这种布局对纹理拾取进行过特别优化。
cudaBindTextureToArray函数，将纹理引用和CUDA数组进行绑定。我们之前写入内容的CUDA数组将成为该纹理引用的后备存储。纹理引用绑定完成后我们调用内核，该内核将进行纹理拾取，同时将结果数据写入到显存中的目标数组。注意:CUDA对于显存中常见的大数据量的存储方式有两种，一种是普通的线性存储，可以直接用指针访问。另外一种则是CUDA数组，对用户不透明，不能在内核里直接用指针访问，需要通过texture或者surface的相应函数进行访问。本例的内核中，从texture reference进行的读取使用了相应的纹理拾取函数，而写入直接用普通的指针(d_out[])进行。当内核执行完成后，结果数组被复制回到主机上的内存中，然后在控制台窗口中显示出来。当使用完纹理存储后，我们需要执行解除绑定的代码，这是通过调用cudaUnbindTexture函数进行的。然后使用cudaFreeArray()函数释放刚才分配的CUDA数组空间。

GPU结构与CUDA系列4GPU存储资源：寄存器，本地内存，共享内存，缓存，显存等存储器细节

0 软件抽象和硬件结构对应关系的例子

把GPU跟一个学校对应起来，学校里有教学楼、操场、食堂，还有老师和学生们；很快有领导（CPU）来检查卫生（需要执行的任务Host程序），因此这个学校的学生们要完成打扫除的工作（Device程序）。

软件抽象资源包括Thread、Warp、Block和Grid

硬件资源包括SP和SM

0.1 软件抽象

Grid对应的是年级

是抽象的划分组织方式

根据年级划分任务，Grid可以分为多个不同的班级

Block对应的是班级

是抽象的划分组织方式

每个班级有若干的同学（线程），可能一个两个不同的年级会出现在同一层楼（SM），或者一层楼只有一个班级，或者没有班级，但是每一层楼的班级最大数量是固定的

Warp对应的是兴趣小组

每个小组有32个学生；（同一时间他们一定是一个班级下的小组）

并且数量固定，即使凑不满这么多学生需要加进来不干活的学生，凑够一个小组

只要求他们有着一样的兴趣爱好（能执行相同的任务）

Thread对应的是学生

一个Thread对应一个SP

每个学生都有个课桌，放自己的物品，不能让别人用，表示每个Thread在软件上都有自己的空间（寄存器等）

0.2 硬件资源

SM对应的是教学楼的一个楼层

是实际存在的资源

一个楼层上可以有多个班级，年级和楼层并没有确定的对应关系，一个楼层中可以有很多来自不同的年级的Block

SM中的SP会被分成兴趣小组，承接不同的任务

SP对应的是学生

一个SP对应一个Thread

是实际存在的资源

每个学生都有个课桌，放自己的物品，不能让别人用，表示每个SP在硬件上都有自己的空间（local memory + registers)；

在楼层中，有公共的空间（走廊、厕所等），这一层楼的所有同学都可以停留，表示一个SM中有shared memory，这个SM上的Block都可以访问；（shared memory是不是所有的block都可以访问）

学校里的公共区域，比如操场、食堂等，所有同学都可以去运动、吃饭，表示GPU中有一些公共的存储空间供所有的Grid访问。

0.3 执行任务

虽然GPU是并行运行，但也并不是我们理想中所有的Thread一起工作，在打扫卫生时，并不是所有学生一起干活，学生经过老师（这里我们理解为Wrap Scheduler）安排后，分为一组一组的小组，每一个小组都只会做一件一样的事情，如果有人先做完了或者不需要做，那么他也会在旁边等他的组员，处于等待状态idle。

1 GPU不同存储的辨析

1.1 总述

这一点跟CPU比较像，就是存储空间越大，访问速度越慢。

GPU越靠近SM的内存就越快。内存的访问速度从快到慢依次为：

Registers->Caches->Shared Memory->Gloabl Memory（Local Memory）。

1.1.1 GPU存储与CPU存储

`CPU`的典型存储结构如下：

一般来说，CPU和内存之间的带宽只有数十GB/s。比如对于Intel Xeon E5-2699 v3，内存带宽达到68GB/s（(2133 * 64 / 8)*4 MB/s）:

内存规格
最大内存大小（取决于内存类型）	768 GB
内存类型	DDR4 1600/1866/2133
最大内存通道数	4
最大内存带宽	68 GB/s

`GPU`的存储结构一般如下:

GPU的高速缓存较小，上图的Memory实际上是指GPU卡内部的显存。但是与显存之间的带宽可以达到数百GB/s，比如P40的显存带宽为346GB/s，远远大于CPU的内存带宽，但是，相对于GPU的计算能力，显存仍然是瓶颈所在。

1.1.3 CPU与GPU交互

在现代的异构计算系统中，GPU是以PCIe卡作为CPU的外部设备存在，两者之间通过PCIe总线通信:

 ----------           ------------
|___DRAM___|         |___GDRAM____|
      |                    |
 ----------           ------------
|   CPU    |         |    GPU     |
|__________|         |____________|
      |                    |
  ---------            --------
 |___IO____|---PCIe---|___IO___|

对于PCIe Gen3 x1理论带宽约为1000MB/s，所以对于Gen3 x32的最大带宽为32GB/s，而受限于本身的实现机制，有效带宽往往只有理论值的2/3还低。所以，CPU与GPU之间的通信开销是比较大的。

1.2 Registers

寄存器是访问速度最快的空间。
当我们在核函数中不加修饰的声明一个变量，那该变量就是寄存器变量，如果在核函数中定义了常数长度的数组，那也会被分配到Registers中；寄存器变量是每个线程私有的，当这个线程的核函数执行完成后，寄存器变量也就不能访问了。
寄存器是比较稀缺的资源，空间很小，Fermi架构中每个线程最多63个寄存器，Kepler架构每个线程最多255个寄存器；一个线程中如果使用了比较少的寄存器，那么SM中就会有更多的线程块，GPU并行计算速度也就越快。
如果一个线程中变量太多，超出了Registers的空间，这时寄存器就会发生溢出，就需要其他内存（Local Memory）来存储，当然程序的运行速度也会降低。
因此，在程序中，对于那种循环操作的变量，我们可以放到寄存器中；同时要尽量减少寄存器的使用数量，这样线程块的数量才能增多，整个程序的运行速度才能更快。

1.3 Local Memory

Local Memory也是每个线程私有的，但却是存储在于Global Memory中的。在核函数中符合存储在寄存器中但不能进入核函数分配的寄存器空间中的变量将被存储在Local Memory中，Local Memory中可能存放的变量有以下几种：

使用未知索引的本地数组
较大的本地数组或结构体
任何不满足核函数寄存器限定条件的变量

1.4 Shared Memory

每个SM中都有共享内存，使用__shared__关键字（CUDA关键字的下划线一般都是两个）定义，共享内存在核函数中声明，生命周期和线程块一致。

同样需要注意的是，SM中共享内存使用太多，会导致SM上活跃的线程数量减少，也会影响程序的运行效率。

数据的共享肯定会导致线程间的竞争，可以通过同步语句来避免内存竞争，同步语句为：

void __syncthreads();

当所有线程都执行到这一步时，才能继续向下执行；频繁调用__syncthreads()也会影响核函数的执行效率。

共享内存因为需要分配给不同的线程所以被分成了不同个Bank，一个Warp中有32个线程，在比较老的GPU中，16个Bank可以同时互相访问，即一条指令就可以让半个Warp同时访问16个Bank，这种并行访问的效率可以极大的提高GPU的效率。比较新的GPU中，一个Warp即32个SP可以同时访问32个Bank，效率又提升了一倍。

下面这个图中：

左边的图每个线程访问一个Bank，不存在内存冲突，通过一个指令即可完成访问所有的访问操作；
中间的图虽然看起来有些乱，但还是一个线程对应一个Bank，也不存在冲突，一个指令即可完成。
右边的图中，存在多个Thread访问一个Bank的情况，如果是读操作，那么GPU底层可以通过广播的方式将数据传给各个Thread，延迟不会很大，但如果是写操作，就必须要等上一个线程写完成后才能进行下一个线程的写操作，延时会比较大。

1.5 Constant Memroy

常量内存驻留在设备内存中，每个SM都有专用的常量内存空间，使用__constant__关键字来声明，可以用来声明一些滤波系数等常量。

常量内存存在于核函数之外，在kernel函数外声明，即常量内存存在于内存中，并不在片上，常量内容的访问速度也是很快的，这是因为每个SM都有专用的常量内存缓存，会把片外的常量读取到缓存中；对所有的核函数都可见，在Host端进行初始化后，核函数不能再修改。

1.6 Texture Memory

纹理内存的使用并不多，它是为了GPU的显示而设计的，这里不多讲了。纹理内存也是存在于片外。

1.7 Global Memory

全局内存，就是我们常说的显存，就是GDDR的空间，全局内存中的变量，只要不销毁，生命周期和应用程序是一样的。

在访问全局内存时，要求是对齐的，也就是一次要读取指定大小（32、64、128）整数倍字节的内存，数据对齐就意味着传输效率降低，比如我们想读33个字节，但实际操作中，需要读取64字节的空间。

对于Global和Constant，Host可以通过下面的函数访问：

cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()

1.8 GPU缓存

每个SM都有一个一级缓存，所有SM公用一个二级缓存，GPU读操作是可以使用缓存的，但写操作不能被缓存。

L1 Cache：Pascal架构上，L1 Cache和Texture已经合为一体（Unified L1/Texture Cache），作为一个连续缓存供给warp使用。

L2 Cache：用来做Global Memory的缓存，容量大，给整个GPU使用。

每个SM有一个只读常量缓存，只读纹理缓存，它们用于设备内存中提高来自于各个内存空间内的读取性能。

讲到缓存，我们就必须要提一点，CPU和GPU在缓存上的一个重要差别就是“缓存一致性”（cache coherency）问题。缓存一致是指一个内存的写操作需要通知所有核的各个级别的缓存，因此，无论何时，所有处理器核看到的内存视图是完全一样的。随着处理器中核数量的增多，这个“通知”的开销迅速增大，使得“缓存一致性”成为限制一个处理器中核数不能太多的一个重要因素。“缓存一致”系统中最坏的情况是，一个内存写操作会强迫每个核的缓存都进行更新，进而每个核都要对相邻的内存单元进行写操作。

CPU遵循“缓存一致”原则，而GPU不是。在GPU中系统不会自动的更新其他核的缓存。所以GPU能扩展到一个芯片内具有大数量的核心。它需要由程序员写清楚每个处理器核输出的各自不同的目标区域。从程序的视角看，这支持一个核仅负责一个输出或者一个小的输出集。

1.9 总结如下：

存储器	作用域	声明期
Register	Thread	Kernel
Local Memory	Thread	Kernel
Shared Memory	Block	Kernel
Global Memory	Grid	Application
Constant	Grid	Application

2 内存与软件硬件的一一对应

2.1 Thread

每一个Thread都有自己的local memory和Registers

即每个同学都可以把自己的东西放到自己的课桌上，别的同学不可以使用；

Local Memory，它是每个线程专有的线程，但却是存在于Global Memory中的，结合我们在第0节例子中拿学校和学生举的例子，可以理解为：学生的课桌都放满了，只能在操场里给他再找个地方放东西，所以访问速度是很慢的，但是这部分还是属于他的local memory，别的线程应该是访问不了的。

2.2 Block

每一个Block有自己的shared memory，构成Block的所有Thread都可以访问。可以被线程中所有的线程共享，其生命周期与线程块一致

即每个班所在的教室里的走道、讲台等，是这个班里同学们的公共区域，别的班级的同学不能进入；

2.3 Grid

Grid之间会有Global memory和Cache

所有的Grid都可以访问，即学校里的操场、餐厅等，是全校同学的公共区域，所有年级的同学都可以共享。

所有的thread（包括不同block的thread）都共享一份 global memory、constant memory、和texture memory。所有的线程都可以访问全局内存（Global Memory）

2.4 Warp

每一个时钟周期内，Warp（一个block里面一起运行的thread，其中各个线程对应的数据资源不同，因为指令相同但是数据不同）现在规定的thread数量是32个。一个block中最多含有16个warp。所以一个block中最多含有512个线程。

2.5 其他存储

只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）
每个SM有自己的L1 cache，SM通过L2 cache连接到Global Memory

3 SM中的存储实现细节

GPU基本处理单元是流多处理器，有关处理单元介绍，之前的文章中有叙述。这篇主要讲存储结构

3.1 寄存器

流处理器先与寄存器交换数据，寄存器负责存储指令，以及指令计算所需的变量。寄存器速度最高，与核心同频。可以无延迟的让Core执行指令运算。

3.2 L1、L2和SMEM

L1和SMEM（Shared Memory）其实算作一个东西，或者准确来讲，在一同一块芯片区域当中。只是根据程序需要，决定哪一部分划分成L1缓存，哪一部分划分成Shared Memory，这个操作是动态的。
L1 cache负责缓存内存地址，而SMEM负责缓存计算所需的变量（包括顶点数据，纹理数据，以及计算引入的中间变量等）。
L1/SMEM缓存频率非常高，延迟非常低，而且访问这个东西所需要的指令周期是可预测的。这个可预测十分重要，这种情况下算法可以优化其计算和访问SMEM里面数据的策略，比如，遇到数据访问操作，可以切换执行其他计算指令，待SMEM数据被取回到寄存器后，就切回刚刚执行数据访问操作的位置，继续执行后续的指令。
L1/SMEM每个SM独享一份，不与其他SM共享。如果有共享需求，则这时候应该L2出场。L2相比L1，其频率没那么高，相应的访问其中的数据需要付出一定的延迟代价。当然了，GPU执行当中为了最大化执行效率，都有动态缓存执行状态和切换执行其他指令的策略。
L2缓存保存了显存的一部分拷贝。在必要情况下（GPU没有独立显存，或者系统也在内存当中划分了一部分区域提供给显存使用），它也可以保存一部分内存的拷贝。所谓拷贝，意味着这是一份连续的显存/内存拷贝（所以指望tex2DLod这样的操作可以减少显存读写以提升性能的可以省省了）。
对于CPU而言，多线程访问L2可能是不安全的。所以CPU一般是提供了额外的指令集去确保L2读-改-写操作是安全的。但是现在某些GPU（比如NVidia Fermi系列）就可以保证这套流程是安全的。GPU相比CPU还有一个好处，它往往对一套数据（比如数组，纹理等）访问是并行的，所以在GPU侧进行这个访问操作就可以安全很多。