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;
}

在这里插入图片描述

关注

笔者 - 东旭

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

信必诺

嗨,支持下哥们呗。

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值