TensorRT笔记四（ShareMemory）

一、共享内存介绍

• 共享内存更靠近计算单元，访问速度更快（ShareMemory > GlobalMemory > PinnedMemory > Pageable Memory）
• 共享内存通常作为访问全局内存（GlobalMemory）的缓存使用
• 利用共享内存实现线程间的通信

二、静态共享内存和动态共享内存

1. 动态共享内存dynamic_shared_memory

• 定义时必须使用extern __shared__开头
• 无论定义多少个，动态共享内存，地址都一样
• 定义时为不确定大小的数组，例如：extern __shared__ char dynamic_shared_memory[];

2. 静态共享内存static_shared_memory

• 定义时不加extern，以__shared__开头
• 定义多个，地址随之叠加
• 定义时为确定大小的数组，例如__shared__ char static_shared_memory[2];

3. 静态共享变量和动态共享变量在kernel函数内/外定义均可，没有限制

三、共享内存的使用

• 共享内存通常与__syncthreads同时出现，该函数同步block内所有线程，全部执行到这一行再继续往下走
• 使用方式：在线程id为0时候，从GlobalMemory中取值，然后syncthreads，然后再使用
• 对一个block内的5个thread来说，如果他们返回结果都是a，要想更快得对这5个a求和，可以在ShareMemory里新建一个共享变量sum，通过共享变量读写数据更快。

例：

//同一个block内的线程相互通信
__shared__ int shared_value1;
__global__ void kernel()
{
	__shared__ int shared_value2;
	if(threadIdx.x == 0)
	{
		if(blockIdx.x == 0){
			shared_value1 = 123;
			shared_value2 = 55;}
		else
		{
			shared_value1 = 321;
			shared_value2 = 66;
		}
	}
	//等待block内所有线程执行
	__syncthreads();
}