GPU CUDA 使用shared memory 运行速度不升反降原因与解决方案

news2024/10/7 4:22:57

写了两张图像相加，以及图像滤波的的几个算子，分别采用shared memory 进行优化。

#include <stdio.h>
#include <cuda_runtime.h>

#include "helper_cuda.h"
#include "helper_timer.h"

#define BLOCKX 32
#define BLOCKY 32
#define BLOCK_SIZE 1024
#define PADDING 2

__global__ void filter5x5(float* in, float* out, int nW, int nH)
{
	// 线程索引 ---> 全局内存索引
	unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
	unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;

	if (row_idx >= PADDING && col_idx >= PADDING && row_idx < nH - PADDING && col_idx < nW - PADDING)
	{
		int sum = 0;
		for (int i = -PADDING; i <= PADDING; i++)
		{
			for (int j = -PADDING; j <= -PADDING; j++)
			{
				sum += in[(row_idx + i) * nW + col_idx + j];
			}
		}
		out[row_idx * nW + col_idx] = sum / ((PADDING * 2 + 1) * (PADDING * 2 + 1));
	}
}

__global__ void filter5x5shared(float* in, float* out, int nW, int nH)
{
	// 声明动态共享内存
	__shared__ int tile[BLOCKY + 4][BLOCKX + 4];

	// 线程索引 ---> 全局内存索引
	const int ig = blockIdx.x * blockDim.x + threadIdx.x;
	const int jg = blockIdx.y * blockDim.y + threadIdx.y;
	const int ijg = jg * nW + ig;

	const int is = threadIdx.x;
	const int js = threadIdx.y;
	//position within smem arrays
	const int ijs = (js + PADDING) * BLOCKX + is + PADDING;

	// 共享内存 存储 操作
	tile[js + PADDING][is + PADDING] = in[ijg];
	//if (is == 0 && js == 0)  //left top
	//{
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		for (int x = 0; x < PADDING; x++)
	//		{
	//			if (jg - (PADDING - y) >= 0 && ig - (PADDING - x) >= 0)
	//				tile[js + y][is + x] = in[(jg - (PADDING - y)) * nW + ig - (PADDING - x)];
	//			else
	//				tile[js + y][is + x] = 0;
	//		}
	//	}
	//}
	//else if (is == 0 && js == blockDim.y - 1)  //left bottom
	//{
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		for (int x = 0; x < PADDING; x++)
	//		{
	//			if (jg + y < nH && ig - (PADDING - x) >= 0)
	//				tile[js + y][is + x] = in[(jg + y) * nW + ig - (PADDING - x)];
	//			else
	//				tile[js + y][is + x] = 0;
	//		}
	//	}
	//}
	//else if (is == blockDim.x - 1 && js == 0)  //right top
	//{
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		for (int x = 0; x < PADDING; x++)
	//		{
	//			if (jg - (PADDING - y) >= 0 && ig + x < nW)
	//				tile[js + y][is + x] = in[(jg - (PADDING - y)) * nW + ig + x];
	//			else
	//				tile[js + y][is + x] = 0;
	//		}
	//	}
	//}
	//else if (is == blockDim.x - 1 && js == blockDim.y - 1)  //right bottom
	//{
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		for (int x = 0; x < PADDING; x++)
	//		{
	//			if (jg + y < nH && ig + x < nW)
	//				tile[js + y][is + x] = in[(jg + y) * nW + ig + x];
	//			else
	//				tile[js + y][is + x] = 0;
	//		}
	//	}
	//}

 //   if(is == 0) //left
 //   {
 //       for (int x = 0; x < PADDING; x++)
 //       {
 //           if (ig - (PADDING - x) >= 0)
 //               tile[js][is + x] = in[jg * nW + ig - (PADDING - x)];
 //           else
 //               tile[js][is + x] = 0;
 //       }
 //   }
 //   else if (js == 0) //top
 //   {
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		if (jg - (PADDING - y) >= 0)
	//			tile[js + y][is] = in[(jg - (PADDING - y)) * nW + ig];
	//		else
	//			tile[js + y][is] = 0;
 //       }
 //   }
 //   else if (is == blockDim.x - 1) //right
 //   {
 //       for (int x = 0; x < PADDING; x++)
 //       {
 //           if (ig + x < nW)
 //               tile[js][is + x] = in[jg * nW + ig + x];
 //           else
 //               tile[js][is + x] = 0;
 //       }
	//}
	//else if (js == blockDim.y - 1) //bottom
	//{
	//	for (int y = 0; y < PADDING; y++)
	//	{
	//		if (jg + y < nH)
	//			tile[js + y][is] = in[(jg + y) * nW + ig];
	//		else
	//			tile[js + y][is] = 0;
	//	}
	//}

	// 等待所有线程完成
	__syncthreads();

	if (jg >= 2 && ig >= 2 && jg < nH - 2 && ig < nW - 2)
	{
		int sum = 0;
		for (int i = -2; i <= 2; i++)
		{
			for (int j = -2; j <= -2; j++)
			{
				sum += tile[js + i][is + j];
			}
		}
		out[jg * nW + ig] = sum / 25;
	}
}

__global__ void vectorAddition(const float* a, const float* b, float* result, int nW, int nH)
{
	const int x = blockIdx.x * blockDim.x + threadIdx.x;
	const int y = blockIdx.y * blockDim.y + threadIdx.y;

	if (x < nW && y < nH)
	{
		result[x + y * nW] = a[x + y * nW] + b[x + y * nW];
	}
}

__global__ void vectorAdditionOptimized(const float* a, const float* b, float* result, int nW, int nH) {
	__shared__ float sharedA[BLOCKX * BLOCKY];
	__shared__ float sharedB[BLOCKX * BLOCKY];

	const int x = blockIdx.x * blockDim.x + threadIdx.x;
	const int y = blockIdx.y * blockDim.y + threadIdx.y;

	// Load data into shared memory
	if (x < nW && y < nH)
	{
		sharedA[threadIdx.x + blockDim.x * threadIdx.y] = a[x + y * nW];
		sharedB[threadIdx.x + blockDim.x * threadIdx.y] = b[x + y * nW];
	}
	else {
		sharedA[threadIdx.x] = 0.0f;
		sharedB[threadIdx.x] = 0.0f;
	}

	// Synchronize to make sure all threads have loaded their data
	__syncthreads();

	// Perform vector addition using data from shared memory
	if (x < nW && y < nH)
	{
		result[x + y * nW] = sharedA[threadIdx.x + blockDim.x * threadIdx.y] +
			sharedB[threadIdx.x + blockDim.x * threadIdx.y];
	}
}

__global__ void vectorAddition(const float* a, const float* b, float* result, int size) {
	int tid = blockIdx.x * blockDim.x + threadIdx.x;

	if (tid < size) {
		result[tid] = a[tid] + b[tid];
	}
}

__global__ void vectorAdditionOptimized(const float* a, const float* b, float* result, int size) {
	__shared__ float sharedA[BLOCK_SIZE];
	__shared__ float sharedB[BLOCK_SIZE];

	int tid = blockIdx.x * blockDim.x + threadIdx.x;

	// Load data into shared memory
	if (tid < size) {
		sharedA[threadIdx.x] = a[tid];
		sharedB[threadIdx.x] = b[tid];
	}
	else {
		sharedA[threadIdx.x] = 0.0f;
		sharedB[threadIdx.x] = 0.0f;
	}

	// Synchronize to make sure all threads have loaded their data
	__syncthreads();

	// Perform vector addition using data from shared memory
	if (threadIdx.x < size)
	{
		result[tid] = sharedA[threadIdx.x] + sharedB[threadIdx.x];
	}
}

int main()
{
	int nW = 1024;
	int nH = 1024;

	float* pHIn, * pHIn2, * pHOut, * pHOut2;
	pHIn = new float[nW * nH];
	pHIn2 = new float[nW * nH];
	pHOut = new float[nW * nH];
	pHOut2 = new float[nW * nH];
	for (int y = 0; y < nW; y++)
	{
		for (int x = 0; x < nH; x++)
		{
			pHIn[x + y * nW] = rand() % 1000;
			pHIn2[x + y * nW] = rand() % 1000;
		}
	}

	float* pIn, * pIn2, * pOut, * pOut2;
	cudaMalloc(&pIn, nW * nH * sizeof(float));
	cudaMalloc(&pIn2, nW * nH * sizeof(float));
	cudaMalloc(&pOut, nW * nH * sizeof(float));
	cudaMalloc(&pOut2, nW * nH * sizeof(float));

	cudaMemcpy(pIn, pHIn, nW * nH * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(pIn2, pHIn2, nW * nH * sizeof(float), cudaMemcpyHostToDevice);

	dim3 threadsPerBlock(BLOCKX, BLOCKY);
	dim3 blocksPerGrid((nW + BLOCKX - 1) / BLOCKX, (nH + BLOCKY - 1) / BLOCKY);

	dim3 threadsPerBlock2(BLOCK_SIZE, 1);
	dim3 blocksPerGrid2((nW * nH + BLOCK_SIZE - 1) / BLOCK_SIZE, 1);

	StopWatchInterface* sw;
	sdkCreateTimer(&sw);
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);


	//vectorAddition << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut, nW, nH);
	//vectorAdditionOptimized << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut, nW, nH);

	sdkStartTimer(&sw);
	cudaEventRecord(start);
	for (int i = 0; i < 100; i++)
	{
		//filter5x5 << <blocksPerGrid, threadsPerBlock >> > (pIn, pOut, nW, nH);
		//vectorAdditionOptimized << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut, nW, nH);
		vectorAddition << <blocksPerGrid2, threadsPerBlock2 >> > (pIn, pIn2, pOut, nW * nH);

	}
	cudaEventRecord(stop);
	cudaDeviceSynchronize();
	sdkStopTimer(&sw);
	float ms1 = sdkGetTimerValue(&sw);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, start, stop);
	std::cout << "Elapsed Time: " << elapsedTime << " milliseconds\n";


	sdkResetTimer(&sw);

	//vectorAddition << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut2, nW, nH);
	//vectorAdditionOptimized << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut, nW, nH);

	sdkStartTimer(&sw);
	cudaEvent_t start1, stop1;
	cudaEventRecord(start);
	for (int i = 0; i < 100; i++)
	{
		//filter5x5shared << <blocksPerGrid, threadsPerBlock >> > (pIn, pOut2, nW, nH);
		//vectorAddition << <blocksPerGrid, threadsPerBlock >> > (pIn, pIn2, pOut2, nW , nH);
		vectorAdditionOptimized << <blocksPerGrid2, threadsPerBlock2 >> > (pIn, pIn2, pOut2, nW * nH);
	}
	cudaEventRecord(stop);
	cudaDeviceSynchronize();
	sdkStopTimer(&sw);
	float ms2 = sdkGetTimerValue(&sw);

	cudaEventElapsedTime(&elapsedTime, start, stop);
	std::cout << "Elapsed Time: " << elapsedTime << " milliseconds\n";

	std::cout << "ms1:" << ms1 << ",ms2:" << ms2 << std::endl;

	cudaMemcpy(pHOut, pOut, nW * nH * sizeof(float), cudaMemcpyDeviceToHost);
	cudaMemcpy(pHOut2, pOut2, nW * nH * sizeof(float), cudaMemcpyDeviceToHost);

	for (int y = 0; y < nW; y++)
	{
		for (int x = 0; x < nH; x++)
		{
			if (abs(pHOut[x + y * nW] - pHOut2[x + y * nW]) > 0.01)
				std::cout << "error" << std::endl;
		}
	}

	cudaFree(pIn);
	cudaFree(pOut);

	return 0;
}

实际运行发现
在这里插入图片描述
使用shared memory的效率反而下降了，实现结果是一致的。

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.coloradmin.cn/o/1196417.html

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

【C语言 | 预处理】C语言预处理详解(三)——内存对齐、手把手带你计算结构体大小

【C语言 | 预处理】C语言预处理详解(三)——内存对齐、手把手带你计算结构体大小

&#x1f601;博客主页&#x1f601;&#xff1a;&#x1f680;https://blog.csdn.net/wkd_007&#x1f680; &#x1f911;博客内容&#x1f911;&#xff1a;&#x1f36d;嵌入式开发、Linux、C语言、C、数据结构、音视频&#x1f36d; &#x1f923;本文内容&#x1f923;&a…
阅读更多...
2023年加氢工艺证考试题库及加氢工艺试题解析

2023年加氢工艺证考试题库及加氢工艺试题解析

题库来源&#xff1a;安全生产模拟考试一点通公众号小程序 2023年加氢工艺证考试题库及加氢工艺试题解析是安全生产模拟考试一点通结合&#xff08;安监局&#xff09;特种作业人员操作证考试大纲和&#xff08;质检局&#xff09;特种设备作业人员上岗证考试大纲随机出的加氢…
阅读更多...
DDU框架学习之路

DDU框架学习之路

目录 MVVM对比 DDU 数据消费者UI 数据的转换者&#xff1a;Domain Layer 数据图生产者/提供者 DataLayer 遵循原理&#xff1a; 单一数据流&#xff1a; Android官方推荐架构&#xff1a;DDU MVVM对比 M&#xff1a;Model 网络层 用于获取远端数据 VM:ViewModel 中间转…
阅读更多...
【Shell脚本9】Shell test 命令

【Shell脚本9】Shell test 命令

Shell test 命令 Shell中的 test 命令用于检查某个条件是否成立&#xff0c;它可以进行数值、字符和文件三个方面的测试。 数值测试 num1100 num2100 if test $[num1] -eq $[num2] thenecho 两个数相等&#xff01; elseecho 两个数不相等&#xff01; fi输出结果&#xff1a…
阅读更多...
关于Android Studio 同步Gradle失败的解决方案

关于Android Studio 同步Gradle失败的解决方案

&#xff08;1&#xff09;打开Android Studio的Settings找到Gradle的目录 &#xff08;2&#xff09;打开本地文件目录&#xff0c;找到对应的gradle版本&#xff0c;可以通过Index of /gradle/ 下载gradle压缩包。把目录中gradle-7.0.2-bin\一堆字符\ 下 的.lck 和.part文…
阅读更多...
Oracle（16）Managing Privileges

Oracle（16）Managing Privileges

目录 一、基础知识 1、Managing Privileges管理权限 2、System Privileges 系统特权 3、System Privileges : Example系统权限&#xff1a;示例 4、Who Can Grant or Revoke? 谁可以授予或撤销权限&#xff1f; 5、The PUBLIC 6、SYSDBA and SYSOPER 7、Revoke with A…
阅读更多...
KT6368A蓝牙芯片的出现部分芯片距离短换芯片就好是什么问题呢

KT6368A蓝牙芯片的出现部分芯片距离短换芯片就好是什么问题呢

一、简介 KT6368A蓝牙芯片的出现部分芯片距离短&#xff0c;换一个芯片距离就好了&#xff0c;是什么问题呢&#xff1f;生产2K的样子 详细说明 按照我们出货客户的跟踪情况&#xff0c;这种问题&#xff0c;可能性极低因为芯片本身的不良率&#xff0c;目前是控制在千分之三…
阅读更多...
js 求数组中的对象某个属性和

js 求数组中的对象某个属性和

可以直接看下效果 代码&#xff1a; <script>let list [{num: 1,price: 10,},{num: 2,price: 10,},{num: 3,price: 10,},{num: 4,price: 10,},]// for循环 求总数和 num的和let num 0for (let i 0; i < list.length; i) {num list[i].num}console.log(第一种&am…
阅读更多...
深度学习 python opencv 火焰检测识别 计算机竞赛

深度学习 python opencv 火焰检测识别 计算机竞赛

文章目录 0 前言1 基于YOLO的火焰检测与识别2 课题背景3 卷积神经网络3.1 卷积层3.2 池化层3.3 激活函数&#xff1a;3.4 全连接层3.5 使用tensorflow中keras模块实现卷积神经网络 4 YOLOV54.1 网络架构图4.2 输入端4.3 基准网络4.4 Neck网络4.5 Head输出层 5 数据集准备5.1 数…
阅读更多...
《红蓝攻防对抗实战》九.内网穿透之利用GRE协议进行隧道穿透

《红蓝攻防对抗实战》九.内网穿透之利用GRE协议进行隧道穿透

​ 前文推荐&#xff1a; 《红蓝攻防对抗实战》一. 隧道穿透技术详解 《红蓝攻防对抗实战》二.内网探测协议出网之TCP/UDP协议探测出网 《红蓝攻防对抗实战》三.内网探测协议出网之HTTP/HTTPS协议探测出网 《红蓝攻防对抗实战》四.内网探测协议出网之ICMP协议探测出网 《红蓝…
阅读更多...
AI：86-基于深度学习的人体姿态估计与运动分析

AI：86-基于深度学习的人体姿态估计与运动分析

🚀 本文选自专栏：人工智能领域200例教程专栏 从基础到实践，深入学习。无论你是初学者还是经验丰富的老手，对于本专栏案例和项目实践都有参考学习意义。 ✨✨✨ 每一个案例都附带有在本地跑过的代码，详细讲解供大家学习，希望可以帮到大家。欢迎订阅支持，正在不断更新中，…
阅读更多...
数据结构线性表——带头双向循环链表

数据结构线性表——带头双向循环链表

前言&#xff1a;小伙伴们好久不见啦&#xff0c;上篇文章我们一起学习了数据结构线性表其一的单链表&#xff0c;了解了单链表的不少好处&#xff0c;但是不可能有完美的数据结构&#xff0c;就算是单链表&#xff0c;也会有很多缺点。 那么今天这篇文章&#xff0c;我们就来…
阅读更多...
全网最细，Apipost接口自动化测试-关联配置，老鸟带你上高速...

全网最细，Apipost接口自动化测试-关联配置，老鸟带你上高速...

目录&#xff1a;导读 前言一、Python编程入门到精通二、接口自动化项目实战三、Web自动化项目实战四、App自动化项目实战五、一线大厂简历六、测试开发DevOps体系七、常用自动化测试工具八、JMeter性能测试九、总结&#xff08;尾部小惊喜&#xff09; 前言 在接口自动化测试…
阅读更多...
Arduino到底适不适合做产品

Arduino到底适不适合做产品

文章目录 一、Arduino性能很低&#xff0c;不如树莓派等开发板&#xff0c;所以不要用Arduino做开发二、Arduino程序效率很低&#xff0c;所以不要用Arduino做开发三、Arduino只能开发玩具&#xff0c;不能做产品四、Arduino开发板成本太高&#xff0c;不适合做产品总结个人见解…
阅读更多...
iPhone或在2024开放第三方应用商店。

iPhone或在2024开放第三方应用商店。

iPhone或开放第三方应用商店&#xff0c;可以说这是一个老生常谈的话题。对于像是iOS这样封闭的系统来说&#xff0c;此前传出苹果可能开放侧载消息的时候&#xff0c;又有谁能信&#xff0c;谁会信&#xff1f; 如果是按照苹果自身的意愿&#xff0c;这种事情自然是不可能发生…
阅读更多...
【LeetCode笔试题】88.合并两个有序数组

【LeetCode笔试题】88.合并两个有序数组

问题描述 给你两个按 非递减顺序 排列的整数数组 nums1 和 nums2&#xff0c;另有两个整数 m 和 n &#xff0c;分别表示 nums1 和 nums2 中的元素数目。 请你 合并 nums2 到 nums1 中&#xff0c;使合并后的数组同样按 非递减顺序 排列。 注意&#xff1a;最终&#xff0c;合…
阅读更多...
王学岗visibility改变后调用onLayout()

王学岗visibility改变后调用onLayout()

自定义控件的时候发现了一个bug。 Button位移动画执行结束后我设置了一个不相关的TextView的可见性由gone变为visible.令人郁闷的是&#xff0c;只要我注释的地方放开。动画执行结束后button都会重新绘制在位移动画开始的位置。注释掉这段代码就正常。 经过分析后得知 View的Vi…
阅读更多...
python OrderedDict类（有序字典）

python OrderedDict类（有序字典）

嗨喽~大家好呀&#xff0c;这里是魔王呐 ❤ ~! python更多源码/资料/解答/教程等 点击此处跳转文末名片免费获取 创建有序字典 import collectionsdic collections.OrderedDict() dic[k1] v1 dic[k2] v2 dic[k3] v3 print(dic)#输出&#xff1a;OrderedDict([(k1, v1), (…
阅读更多...
Vatee万腾科技决策力的未来展望：开创数字化创新的新高度

Vatee万腾科技决策力的未来展望：开创数字化创新的新高度

随着科技不断演进&#xff0c;Vatee万腾的科技决策力在数字化创新领域展现出了强大的潜力和前瞻性。 Vatee万腾的科技决策力被视为数字化创新的引擎&#xff0c;为未来创新注入了新的动力。通过深刻的市场洞察和科学决策&#xff0c;Vatee万腾致力于推动数字化创新走向新的高度…
阅读更多...
图论11-欧拉回路与欧拉路径+Hierholzer算法实现

图论11-欧拉回路与欧拉路径+Hierholzer算法实现

文章目录 1 欧拉回路的概念2 欧拉回路的算法实现3 Hierholzer算法详解4 Hierholzer算法实现4.1 修改Graph&#xff0c;增加API4.2 Graph.java4.3 联通分量类4.4 欧拉回路类 1 欧拉回路的概念 2 欧拉回路的算法实现 private boolean hasEulerLoop(){CC cc new CC(G);if(cc.cou…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023