告訴你cuda共享記憶體的使用

想必大家都知道，cuda裡面每一個block上有一塊高速緩沖區，這就是提供給block裡面各個線程使用的shared memory，那怎麼使用這一塊記憶體呢？

首先，shared memory分為固定配置設定方式和動态配置設定方式，就是上圖的Static Shared Memory和Dynamic Shared Memory

1，固定配置設定

直接__shared__ int seme[5] ;這就是在每一個block裡面配置設定5個int（20B）

__global__ void addKernel(int *c, const int *a)
{
	int i = threadIdx.x;
	 __shared__ int smem[5];
	smem[i] = a[i];
	__syncthreads();
	if (i == 0)	//0号線程做平方和
	{
		c[0] = 0;
		for (int d = 0; d<5; d++)
		{
			c[0] += smem[d] * smem[d];
		}
	}
	if (i == 1)//1号線程做累加
	{
		c[1] = 0;
		for (int d = 0; d<5; d++)
		{
			c[1] += smem[d];
		}
	}
	if (i == 2)	//2号線程做累乘
	{
		c[2] = 1;
		for (int d = 0; d<5; d++)
		{
			c[2] *= smem[d];
		}
	}

}
           

調用，啟動的時候，block個數1，是以shared memory使用20B

通過nsight可以看出，使用了20B的共享記憶體，并且是Static的；

2，動态配置設定

沒錯，就是在block裡面聲明，前面加上extern；

__global__ void addKernel(int *c, const int *a)
{
	int i = threadIdx.x;
	 extern __shared__ int smem[];
	smem[i] = a[i];
	__syncthreads();
	if (i == 0)	//0号線程做平方和
	{
		c[0] = 0;
		for (int d = 0; d<5; d++)
		{
			c[0] += smem[d] * smem[d];
		}
	}
	if (i == 1)//1号線程做累加
	{
		c[1] = 0;
		for (int d = 0; d<5; d++)
		{
			c[1] += smem[d];
		}
	}
	if (i == 2)	//2号線程做累乘
	{
		c[2] = 1;
		for (int d = 0; d<5; d++)
		{
			c[2] *= smem[d];
		}
	}

}
           

那在哪裡指定大小呢？

原來是啟動核函數的時候指定的第三個參數，之前使用多個流的時候，第四個參數綁定流的序号，第三個參數總是設為0，現在終于明白它的含義了

這幾天正在準備寫一篇關于cuda流的使用，然後會加上一些自己的學習總結，年輕，幹就完了，奧利幹！