共享内存作用
- 共享内存在片上(on-chip),与本地内存和全局内存相比具有更高的带宽和更低的延迟;
- 共享内存中的数据在线程块内所有线程可见,可用线程间通信,共享内存的生命周期也与所属线程块一致;
- 使用__shared__修饰的变量存放于共享内存中
- 共享内存可定义动态与静态两种;
- 每个SM的共享内存数量是一定的,也就是说,如果在单个线程块中分配过度的共享内存,将会限制活跃线程束的数量;
- 访问共享内存必须加入同步机制:线程块内同步 void_syncthreads();
- 不同计算能力的架构,每个SM中拥有的共享内存大小是不同的
- 每个线程块使用共享内存的最大数量不同架构是不同的
-
此功能(与线程同步结合)有许多用途,例如用户管理的数据缓存、高性能协作并行算法(例如并行缩减),以及在原本不可能的情况下促进全局内存合并。
静态共享内存
定义在核函数中,定义时确定大小。
示例:CUDAExample: 点积运算 Dot Product-CSDN博客
动态共享内存
在声明的时候没有指定数组的大小
静态和动态共享内存的示例:实现数组的逆向交给交互
#include <stdio.h>
#include "cpu_anim.h"
#include "cuda_runtime.h"
#include <device_launch_parameters.h>
#include <device_functions.h>// 静态共享内存
__global__ void staticReverse(int* d, int n)
{__shared__ int s[64];int t = threadIdx.x;int tr = n - t - 1;s[t] = d[t];__syncthreads();d[t] = s[tr];
}// 动态共享内存的第一种调用方法
__global__ void dynamicReverse(int* d, int n)
{extern __shared__ int s[];int t = threadIdx.x;int tr = n - t - 1;s[t] = d[t];__syncthreads();d[t] = s[tr];
}// 动态共享内存的第一种调用方法
extern __shared__ int d_s[]; // 固定格式
__global__ void dynamicReverseEx(int* d, int n)
{int t = threadIdx.x;int tr = n - t - 1;d_s[t] = d[t];__syncthreads();d[t] = d_s[tr];
}int main(void)
{const int n = 64;int a[n], r[n], d[n];for (int i = 0; i < n; i++) {a[i] = i;r[i] = n - i - 1;d[i] = 0;}int* d_d;cudaMalloc(&d_d, n * sizeof(int));// run version with static shared memorycudaMemcpy(d_d, a, n * sizeof(int), cudaMemcpyHostToDevice);staticReverse << <1, n >> > (d_d, n);cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);for (int i = 0; i < n; i++)if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);// run dynamic shared memory versioncudaMemcpy(d_d, a, n * sizeof(int), cudaMemcpyHostToDevice);// 第三个参数指定的是共享内存的大小dynamicReverse << <1, n, n * sizeof(int) >> > (d_d, n);cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);for (int i = 0; i < n; i++)if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);// run dynamic shared memory version cudaMemcpy(d_d, a, n * sizeof(int), cudaMemcpyHostToDevice);// 第三个参数指定的是共享内存的大小dynamicReverseEx << <1, n, n * sizeof(int) >> > (d_d, n);cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);for (int i = 0; i < n; i++)if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);return 0;
}