CUDA学习(五)之使用共享内存(shared memory)进行归约求和(一个包含N个线程的线程块)
2024-09-07 02:19:46
共享内存(shared memory)是位于SM上的on-chip(片上)一块内存,每个SM都有,就是内存比较小,早期的GPU只有16K(16384),现在生产的GPU一般都是48K(49152)。
共享内存由于是片上内存,因而带宽高,延迟小(较全局内存而言),合理使用共享内存对程序效率具有很大提升。
下面是使用共享内存对一个数组进行求和,使用全局内存进行归约求和可以浏览https://www.cnblogs.com/xiaoxiaoyibu/p/11397205.html
#pragma once
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h" #include <iostream> using namespace std; const int N = ; //数组长度 __global__ void d_ParallelTest(double *Para)
{
int tid = threadIdx.x; //----使用shared memory-------------------------------------------------------------- __shared__ double s_Para[N]; //定义长度为N的共享内存数组
if (tid < N) //循环整个数组,每个线程负责将一个元素从全局内存载入共享内存
s_Para[tid] = Para[tid];
__syncthreads(); //(红色下波浪线提示由于VS不识别,不影响运行)同步,等待所有线程把自己负责的元素载入到共享内存再执行下面代码
for (int index = ; index < blockDim.x; index *= )
{
__syncthreads(); //同步,以防止归约过程中某个线程运行速度过快导致计算错误(后面线程计算使用错误的前面线程结果值)
if (tid % ( * index) == )
{
s_Para[tid] += s_Para[tid + index];
}
} if (tid == ) //整个数组相加完成后,将共享内存数组0号元素的值赋给全局内存数组0号元素,最后返回CPU端
Para[tid] = s_Para[tid];
} void ParallelTest()
{ double *Para;
cudaMallocManaged((void **)&Para, sizeof(double) * N); //统一内存寻址,CPU和GPU都可以使用 double ParaSum = ; for (int i=; i<N; i++)
{
Para[i] = (i + ) * 0.1; //数组赋值
ParaSum += Para[i]; //CPU端数组累加
}
cout << " CPU result = " << ParaSum << endl; //显示CPU端结果 double d_ParaSum;
d_ParallelTest << < , N>> > (Para); //调用核函数(一个包含N个线程的线程块) cudaDeviceSynchronize(); //等待设备端同步
d_ParaSum = Para[]; //从累加过后数组的0号元素得出结果
cout << " GPU result = " << d_ParaSum << endl; //显示GPU端结果
} int main()
{ //并行归约
ParallelTest(); system("pause"); return ;
}
结果如下(CPU和GPU结果一致):
最新文章
- Atitit zxing二维码qr码识别解析
- 初识c#
- 亚马逊EC2弹性IP收费
- Base64正反编码
- Java之构造器的作用
- [51nod1685]第k大区间
- mysql导入数据出错
- 织梦(dedecms) 5.7 /plus/car.php sql注入0day
- 20145305 《Java程序设计》实验一
- Mssql显错和不显错模式下的注入
- WebBrowser中取对应的图片资源
- Node.js学习 - Global Object
- 实战-CentOS6.8配置nfs服务
- matplotlib学习笔记
- C#发布和调试WebService
- 笔记:MYSQL四种事务隔离级。
- 13.5.SolrCloud集群使用手册之数据导入
- linux Shell 脚本编写
- 【Unity】第6章 Unity脚本开发基础
- 用emoji表情包来可视化北京市历史天气状况!