限定符__shared__
共享内存是CUDA中最重要的性能优化工具之一,它位于每个流多处理器(SM)上,提供比全局内存快100倍以上的访问速度。__shared__限定符用于声明共享内存变量。
核心特性
基本声明方式
静态分配(编译时确定大小)
1、生命周期与线程块绑定:静态分配的共享内存与线程块的生命周期一致,线程块启动时分配,结束时释放。
2、编译时确定大小:需要在核函数编译时确定共享内存的大小(如__shared__ int smem),运行时无法动态调整。
3、线程块内可见:同一线程块内的所有线程均可访问该内存,但不同线程块无法共享数据。
__shared__ float s_data[1024]; // 静态共享内存数组
动态分配(运行时确定大小)
1、动态共享内存通过extern __shared__声明,在核函数调用时通过第三个执行配置参数指定大小
2、所有动态声明的共享内存变量共享同一块内存区域,编译器不会为每个变量分配独立空间
3、动态共享内存的地址空间是连续的,多个变量会按声明顺序占用同一内存块的不同偏移位置
4、若两个变量总大小超过执行配置中分配的空间,会导致运行时错误。使用建议: 手动管理变量在共享内存中的偏移量、或通过指针算术或结构体方式组织数据
extern __shared__ int s_dynamic[]; // 声明动态共享内存
// 内核调用时指定大小(参数3)
myKernel<<<blocks, threads, sharedMemSize>>>(...);
避免存储体冲突
共享内存被分为32个存储体(32位架构)。当同一个warp中的多个线程访问同一个存储体的不同地址时,会发生串行访问。可使用 __syncthreads() 所有线程同步。
运行效果及代码
静态分配代码示例
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define N 256 // 数组大小
#define THREADS_PER_BLOCK 16 // 每个线程块的线程数
__global__ void addArraysWithSharedMemory(int *A, int *B, int *C)
{
// 线程的唯-TD
int index = threadIdx.x + blockIdx.x*blockDim.x;
//printf("%d + %d * %d = %d\n", threadIdx.x, blockIdx.x, blockDim.x, index);
// 在共享内存中分配一个大小为THREADS_PER_BLOCK 的空间
__shared__ int shared_A[THREADS_PER_BLOCK];
__shared__ int shared_B[THREADS_PER_BLOCK];
//每个线程将A和B的元素加载到共享内存中
if (index < N)
{
shared_A[threadIdx.x] = A[index];
shared_B[threadIdx.x] = B[index];
}
// 确保所有线程都完成了数据加载到共享内存中
__syncthreads();
// 每个线程计算 A[i]+B[i],并将结果存储到
if (index < N)
{
C[index] = shared_A[threadIdx.x] + shared_B[threadIdx.x];
}
}
int main()
{
int size = N * sizeof(int);
int *h_A = (int*)malloc(size);
int *h_B = (int*)malloc(size);
int *h_C = (int*)malloc(size);
// 初始化数据
for (int i = 0; i < N; i++)
{
h_A[i] = i;
h_B[i] = i * 2;
}
int *d_A,*d_B,*d_C;
cudaMalloc((void**)&d_A, size);
cudaMalloc((void**)&d_B, size);
cudaMalloc((void**)&d_C, size);
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 启动核函数,计算数组的和
int numBlocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
addArraysWithSharedMemory << <numBlocks, THREADS_PER_BLOCK >> > (d_A, d_B, d_C);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 打印前10个结果
for (int i= 0; i < 10; i++)
{
std::cout << "C[" << i << "]=" << h_C[i] << std::endl;
}
// 释放内存
free(h_A);
free(h_B);
free(h_C);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
return 0;
}
动态分配代码示例
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define N 256 // 数组大小
#define THREADS_PER_BLOCK 16 // 每个线程块的线程数
extern __shared__ int shared[]; // 声明动态共享内存
__global__ void addArraysWithSharedMemory(int *A, int *B, int *C)
{
int* shared_A = shared;
int* shared_B = (int*)&shared[16];
// 线程的唯-TD
int index = threadIdx.x + blockIdx.x*blockDim.x;
//printf("%d + %d * %d = %d\n", threadIdx.x, blockIdx.x, blockDim.x, index);
//每个线程将A和B的元素加载到共享内存中
if (index < N)
{
shared_A[threadIdx.x] = A[index];
shared_B[threadIdx.x] = B[index];
printf("%d -> A(%d),B(%d)\n", index, shared_A[threadIdx.x], shared_B[threadIdx.x]);
}
// 确保所有线程都完成了数据加载到共享内存中
__syncthreads();
// 每个线程计算 A[i]+B[i],并将结果存储到
if (index < N)
{
C[index] = shared_A[threadIdx.x] + shared_B[threadIdx.x];
}
}
int main()
{
int size = N * sizeof(int);
int *h_A = (int*)malloc(size);
int *h_B = (int*)malloc(size);
int *h_C = (int*)malloc(size);
// 初始化数据
for (int i = 0; i < N; i++)
{
h_A[i] = i;
h_B[i] = i*2;
}
int *d_A, *d_B, *d_C;
cudaMalloc((void**)&d_A, size);
cudaMalloc((void**)&d_B, size);
cudaMalloc((void**)&d_C, size);
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 启动核函数,计算数组的和。
//(由于核函数使用2个动态共享内存变量,那么两个变量内存占用内存一致,所以需要分配2份内存并在核函数内手动管理变量在共享内存中的偏移量,使其计算正确)
int numBlocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
addArraysWithSharedMemory << <numBlocks, THREADS_PER_BLOCK , THREADS_PER_BLOCK*2 >> > (d_A, d_B, d_C);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 打印前10个结果
for (int i = 0; i < 10; i++)
{
std::cout << "C[" << i << "]=" << h_C[i] << std::endl;
}
// 释放内存
free(h_A);
free(h_B);
free(h_C);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
return 0;
}
关注
笔者 - 东旭