CUDA —— 2.4、cuda共享内存__shared__使用介绍（附：多个动态共享内存变量手动分配办法、及完整代码）

限定符__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;
}

 

关注

笔者 - 东旭