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

写了两张图像相加，以及图像滤波的的几个算子，分别采用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.mfbz.cn/a/132044.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

永达理简析：利用保险的“财务规划”功能维持退休后生活水平

永达理简析：利用保险的“财务规划”功能维持退休后生活水平

现代社会环境背景下&#xff0c;“自养自老”已经是一种未来养老趋势&#xff0c;很多人会为自己准备一份长期、比较周全的保障&#xff0c;这样财务规划不仅会分担子女的压力&#xff0c;也让自己有一个长远的保障。在各种财务储蓄工具中&#xff0c;商业保险占据着不可取代的…
阅读更多...
Redis实现分布式锁

Redis实现分布式锁

文章目录 前言一、概述为什么使用分布式锁基本原理分布式锁应该具备哪些条件常见的三种分布式锁 二、基于Redis实现分布式锁误删锁问题原子性问题最终代码实现 总结 前言 Redis实现简单分布式锁。 一、概述 为什么使用分布式锁 在多线程环境中&#xff0c;如果多个线程同时访…
阅读更多...
2023 年最好的 Android 系统修复/刷机应用程序和软件

2023 年最好的 Android 系统修复/刷机应用程序和软件

任何 Android 设备要顺利运行&#xff0c;其操作系统必须运行良好。幸运的是&#xff0c;对于大多数 Android 用户来说&#xff0c;这是不间断的。设备运行良好&#xff0c;打电话、共享文档等都没有问题。尽管如此&#xff0c;Android 操作系统可能会停止运行。这可能是由于特…
阅读更多...
HTTP-FLV详解及分析

HTTP-FLV详解及分析

文章目录 前言一、HTTP-FLV 简介1、市场上使用 http-flv 的商家2、http-flv、rtmp 和 hls 直播的优缺点3、http-flv 技术实现 二、Nginx 配置 http-flv1、Windows 安装 nginx&#xff0c;已经集成 nginx-http-flv-module2、nginx.conf 配置文件3、运行 nginx 服务器4、ffmpeg 推…
阅读更多...
Windows系统安装2个版本得的MySQL

Windows系统安装2个版本得的MySQL

一、MySQL官网下载对应版本的zip文件 最新版本8.0.34下载链接&#xff1a;https://dev.mysql.com/downloads/mysql/ MySQL 5.7下载链接&#xff1a;https://downloads.mysql.com/archives/community/ 二、将下载到的压缩包解压到指定目录 使用解压工具将下载到的压缩包解…
阅读更多...
开源Gimp动态压感笔刷设置方法

开源Gimp动态压感笔刷设置方法

一、问题描述 开源绘画工具的Gimp的笔刷压感在哪里控制和开启呢&#xff1f; 二、解决方法 1、Gimp有专用的笔刷集&#xff1a;如下图。开启需要在主窗口window下拉菜单开启&#xff0c;或在右侧面板里的左箭头按钮里打开。一般绘画够用了。比用自定义特殊笔刷。 2、如果要调…
阅读更多...
C++：this指针和构造与析构的运用

C++：this指针和构造与析构的运用

目录 一&#xff0c;this指针 二&#xff0c;构造函数 三&#xff0c;析构函数 四&#xff0c;析构与构造的调用 一&#xff0c;this指针 首先&#xff0c;我们先观察以下类&#xff1a; #include <iostream> using namespace std; class Date { public: void In…
阅读更多...
HCIA-hybrid经典小实验

HCIA-hybrid经典小实验

hybrid经典小实验 实验拓扑配置实现SW1SW2 配置验证PC1-PC3 不能通信PC1-PC2 正常通信其他自行测试 实验拓扑 配置实现 SW1 sysname SW1 # undo info-center enable # vlan batch 10 20 30 # interface Ethernet0/0/1 //接口发送该vlan-id的数据帧时&#xff0c;不剥离帧中的…
阅读更多...
JavaSE 类与对象

JavaSE 类与对象

前言 我们之前学的都是面向过程&#xff0c;面向过程研究的是对单个对象的一种方法实现过程&#xff0c;比如求一个数的阶乘&#xff0c;强调的是怎么实现这个方法的过程&#xff0c;但对我们以后来说&#xff0c;如果想要应用到更广的层面&#xff0c;不能只是学习一个方法的…
阅读更多...
JMeter 常见函数讲解

JMeter 常见函数讲解

&#x1f4e2;专注于分享软件测试干货内容&#xff0c;欢迎点赞 &#x1f44d; 收藏 ⭐留言 &#x1f4dd; 如有错误敬请指正&#xff01;&#x1f4e2;交流讨论&#xff1a;加入1000人软件测试技术学习交流群&#x1f4e2;资源分享&#xff1a;进了字节跳动之后&#xff0c;才…
阅读更多...
宠物社区系统宠物领养小程序，宠物救助小程序系统多少钱？

宠物社区系统宠物领养小程序，宠物救助小程序系统多少钱？

当前很多的宠物被抛弃和虐杀&#xff0c;它们没有选择权&#xff0c;我们强制性的把狗带进人类的生活中&#xff0c;然后又无情的抛弃&#xff0c;让它们无家可归&#xff0c;变成流浪狗&#xff0c;它们做错了什么&#xff1f;流浪动物被主人遗弃之后居无定所&#xff0c;时刻…
阅读更多...
C语言【趣编程】我们怎样便捷输出空心的金字塔

C语言【趣编程】我们怎样便捷输出空心的金字塔

目录 1问题&#xff1a; 2解题思路&#xff1a; 3代码如下&#xff1a; 4代码运行结果如下图所示&#xff1a; 5总结&#xff1a; r如若后续有不会的问题&#xff0c;可以和我私聊&#xff1b; 1问题&#xff1a; 2解题思路&#xff1a; 方法&#xff1a;找规律&#xff0…
阅读更多...
AI系统ChatGPT源码+详细搭建部署教程+AI绘画系统+支持GPT4.0+Midjourney绘画+已支持OpenAI GPT全模型+国内AI全模型

AI系统ChatGPT源码+详细搭建部署教程+AI绘画系统+支持GPT4.0+Midjourney绘画+已支持OpenAI GPT全模型+国内AI全模型

一、AI创作系统 SparkAi创作系统是基于OpenAI很火的ChatGPT进行开发的Ai智能问答系统和Midjourney绘画系统&#xff0c;支持OpenAI-GPT全模型国内AI全模型。本期针对源码系统整体测试下来非常完美&#xff0c;可以说SparkAi是目前国内一款的ChatGPT对接OpenAI软件系统。那么如…
阅读更多...
2023/11/11

2023/11/11

1. 实现字体渐变色 html <div class"gradient">实现字体渐变色 </div>css .gradient {display: inline-block;font-weight: 800;font-size: 40px;color: #fff;background: linear-gradient(90deg, #f00 0%, #000 50%, #00f 100%);background-clip: bo…
阅读更多...
积分上限函数

积分上限函数

定积分的形式 a&#xff1a;积分下限 b&#xff1a;积分上限 定积分的值与积分变量无关 积分上限函数的形式 x&#xff1a;自变量 t&#xff1a;积分变量 积分上限是变量&#xff0c;积分下限是常数 定积分的几何意义 x轴所围成面积 x轴以上面积为正 x轴以下面积为负 积分…
阅读更多...
Linux学习教程（第一章 简介）4

Linux学习教程（第一章 简介）4

第一章 简介 十一、Linux的主要应用领域有哪些？ Linux 似乎在我们平时的生活中很少看到，那么它应用在哪些领域呢？其实，在生活中随时随地都有 Linux 为我们服务着。 1、网站服务器 用事实说话！访问国际知名的 Netcraft 网站 http:// www.netcraft.com，在 "Whats …
阅读更多...
数据分析实战 | 泊松回归——航班数据分析

数据分析实战 | 泊松回归——航班数据分析

目录 一、数据及分析对象 二、目的及分析任务 三、方法及工具 四、数据读入 五、数据理解 六、数据准备 七、模型训练 八、模型评价 一、数据及分析对象 CSV文件&#xff1a;o-ring-erosion-only.csv 数据集链接&#xff1a;https://download.csdn.net/download/m0_7…
阅读更多...
网络安全基础之php开发文件上传的实现

网络安全基础之php开发文件上传的实现

前言 php是网络安全学习里必不可少的一环&#xff0c;简单理解php的开发环节能更好的帮助我们去学习php以及其他语言的web漏洞原理 正文 在正常的开发中&#xff0c;文件的功能是必不可少&#xff0c;比如我们在论坛的头像想更改时就涉及到文件的上传等等文件功能。但也会出…
阅读更多...
Leetcode刷题详解——单词搜索

Leetcode刷题详解——单词搜索

1. 题目链接&#xff1a;79. 单词搜索 2. 题目描述&#xff1a; 给定一个 m x n 二维字符网格 board 和一个字符串单词 word 。如果 word 存在于网格中&#xff0c;返回 true &#xff1b;否则&#xff0c;返回 false 。 单词必须按照字母顺序&#xff0c;通过相邻的单元格内的…
阅读更多...
基于ubuntu 22, jdk 8x64搭建图数据库环境 hugegraph--google镜像chatgpt

基于ubuntu 22, jdk 8x64搭建图数据库环境 hugegraph--google镜像chatgpt

基于ubuntu 22, jdk 8x64搭建图数据库环境 hugegraph download 环境 uname -a #Linux whiltez 5.15.0-46-generic #49-Ubuntu SMP Thu Aug 4 18:03:25 UTC 2022 x86_64 x86_64 x86_64 GNU/Linuxwhich javac #/adoptopen-jdk8u332-b09/bin/javac which java #/adoptopen-jdk8u33…
阅读更多...
最新文章

Copyright @ 2022~2023

友情链接： 我的编程经验分享 一条长河网 尧图网站开发 尧图建网站