GPU-CUDA编程学习（四）-共享内存

Shared memory

共享内存在芯片上可用，因此它比全局内存快得多。共享内存延迟大约比未调用的全局内存延迟低100倍。来自同一块的所有线程都可以访问共享内存。这在许多需要与其他线程共享结果的应用程序中非常有用。但是，如果没有同步，也会产生混乱或错误的结果。如果一个线程在其他线程写入数据之前从内存中读取数据，那么可能会导致错误的结果。因此，应该适当地控制或管理内存访问。这是由剩余的**syncthreads()**指令完成的，该指令确保在程序继续执行之前，对内存的所有写操作都已完成。这也叫做阻塞。阻塞的意思是所有线程都将到达这一行并等待其他线程完成。在所有线程都达到这个阻塞之后，它们可以进一步移动。为了演示共享内存和线程同步的使用，本文给出了一个移动平均的示例。其核函数如下:

#include <stdio.h>
__global__ void gpu_shared_memory(float *d_a)
{
   
int i, index = threadIdx.x;
float average, sum = 0.0f;
//Defining shared memory
__shared__ float sh_arr[10];
sh_arr[index] = d_a[index];
// This directive ensure all the writes to shared memory have completed
__syncthreads();
for (i = 0; i<= index; i++)
{
   
sum += sh_arr[i];
}
average = sum / (index + 1.0f);
d_a[index] = average;
//This statement is redundant and will have no effect on overall code execution
sh_arr[index] = average;
}

移动平均操作只是找到一个数组中直到当前元素的所有元素的平均值。许多线程的计算都需要相同的数组数据。这是使用共享内存的理想情况，它将比全局内存提供更快的数据。这将减少每个线程的全局内存访问次数，从而减少程序的延迟。共享内存的位置是用__shared__指令定义的。在本例中，定义了十个浮点数元素的共享内存。通常，共享内存的大小应该等于每个块的线程数。在这里，我们处理的是数组长度10，因此我们获得了这个大小的共享内存。下一步是将数据从全局内存复制到这个共享内存。所有线程都将按其线程ID索引的元素复制到共享数组中。现在，这是一个共享内存的写操作，在下一行中，我们将从这个共享数组中读取。因此，在继续之前，我们应该确保所有共享内存写操作都已完成。因此，让我们引入**__synchronizethreads()**阻塞。

接下来，for循环使用共享内存中的值计算到当前元素的所有元素的平均值，并将结果存储在由当前线程ID索引的全局内存中。现在，我们将尝试为这段代码编写如下的main函数：

int main(int argc, char **argv)
{
   
float h_a[10];
float *d_a;
//Initialize host Array
for (int i = 0; i < 10; i++)
{
   
h_a[i] = i;
}
// allocate global memory on the device
cudaMalloc((void **)&d_a, sizeof(float) * 10);
// 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
cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);
printf("Use of Shared Memory on GPU: \n");
for (int i = 0