cuda并行归约算法（不断优化过程）

归约

方法

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[101], int b[101]) {__shared__ int sdata[101];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x;sdata[idx] = a[x];__syncthreads();for (int s = 1; s < blockDim.x; s *= 2) {if (idx % (s * 2) == 0) {sdata[idx] += sdata[idx + s];}__syncthreads();}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[101];int *da,*db;int hb[101];for(int i = 0; i <= 100; i++)ha[i] = i;cudaMalloc((void**)&da, sizeof(int) * 101);cudaMalloc((void**)&db, sizeof(int) * 101);cudaMemcpy(da, ha, sizeof(int) * 101, cudaMemcpyHostToDevice);dim3 Block(101);jia << <1, Block >> > (da,db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

优化一

在上一方法中对于线程的使用由于间隔使用会造成很大的资源浪费

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[101], int b[101]) {__shared__ int sdata[101];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x;sdata[idx] = a[x];__syncthreads();for (int s = 1; s < blockDim.x; s *= 2) {int indx = 2 * s * idx;if (indx<blockDim.x) {sdata[indx] += sdata[indx + s];}__syncthreads();}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[101];int *da,*db;int hb[101];for(int i = 0; i <= 100; i++)ha[i] = i;cudaMalloc((void**)&da, sizeof(int) * 101);cudaMalloc((void**)&db, sizeof(int) * 101);cudaMemcpy(da, ha, sizeof(int) * 101, cudaMemcpyHostToDevice);dim3 Block(101);jia << <1, Block >> > (da,db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

优化二

提高内存访问效率

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[128], int b[128]) {__shared__ int sdata[100];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x;sdata[idx] = a[x];__syncthreads();for (int s = blockDim.x / 2; s > 0; s /= 2) {if (idx < s) {sdata[idx] += sdata[idx + s];}__syncthreads();}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[128] = { 0 };int* da, * db;int hb[128];for (int i = 0; i < 100; i++)ha[i] = i + 1;cudaMalloc((void**)&da, sizeof(int) * 128);cudaMalloc((void**)&db, sizeof(int) * 128);cudaMemcpy(da, ha, sizeof(int) * 128, cudaMemcpyHostToDevice);dim3 Block(128);jia << <1, Block >> > (da, db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

优化三

在加载共享内存时可以在一个共享内存中同时加载两个数值进去，先进行一次归约，提高效率

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<math.h>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[128], int b[128]) {__shared__ int sdata[100];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x * 2;sdata[idx] = a[x] + a[x + 1];__syncthreads();for (int s = blockDim.x / 2; s > 0; s /= 2) {if (idx < s) {sdata[idx] += sdata[idx + s];}__syncthreads();}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[128] = { 0 };int* da, * db;int hb[128];for (int i = 0; i < 100; i++)ha[i] = i + 1;cudaMalloc((void**)&da, sizeof(int) * 128);cudaMalloc((void**)&db, sizeof(int) * 128);cudaMemcpy(da, ha, sizeof(int) * 128, cudaMemcpyHostToDevice);dim3 Block(128);jia << <1, Block >> > (da, db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

优化四

当我们执行到只有32个线程在计算值的时候，此时只有一个warp在使用而其他的warp没有使用此时会造成资源的浪费，我们可以单独将最后一个warp展开运行

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<math.h>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[128], int b[128]) {__shared__ int sdata[100];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x * 2;sdata[idx] = a[x] + a[x + 1];__syncthreads();for (int s = blockDim.x / 2; s > 32; s /= 2) {if (idx < s) {sdata[idx] += sdata[idx + s];}__syncthreads();}if (idx < 32) {sdata[idx] += sdata[idx + 32];sdata[idx] += sdata[idx + 16];sdata[idx] += sdata[idx + 8];sdata[idx] += sdata[idx + 4];sdata[idx] += sdata[idx + 2];sdata[idx] += sdata[idx + 1];}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[128] = { 0 };int* da, * db;int hb[128];for (int i = 0; i < 100; i++)ha[i] = i + 1;cudaMalloc((void**)&da, sizeof(int) * 128);cudaMalloc((void**)&db, sizeof(int) * 128);cudaMemcpy(da, ha, sizeof(int) * 128, cudaMemcpyHostToDevice);dim3 Block(128);jia << <1, Block >> > (da, db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

优化五

根据线程数量必须是2的n次幂，我们可以在同一个block里面将for循环完全拆除，从而加快运行速率

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<iostream>
#include<math.h>
#include<device_functions.h>
using namespace std;
__global__ void jia(int a[128], int b[128]) {__shared__ int sdata[100];int idx = threadIdx.x;int x = blockIdx.x * blockDim.x + threadIdx.x * 2;sdata[idx] = a[x] + a[x + 1];__syncthreads();if (idx < 32) {sdata[idx] += sdata[idx + 32];sdata[idx] += sdata[idx + 16];sdata[idx] += sdata[idx + 8];sdata[idx] += sdata[idx + 4];sdata[idx] += sdata[idx + 2];sdata[idx] += sdata[idx + 1];}if (idx == 0)b[0] = sdata[0];
}
int main() {int ha[128] = { 0 };int* da, * db;int hb[128];for (int i = 0; i < 100; i++)ha[i] = i + 1;cudaMalloc((void**)&da, sizeof(int) * 128);cudaMalloc((void**)&db, sizeof(int) * 128);cudaMemcpy(da, ha, sizeof(int) * 128, cudaMemcpyHostToDevice);dim3 Block(128);jia << <1, Block >> > (da, db);cudaMemcpy(hb, db, sizeof(int), cudaMemcpyDeviceToHost);cout << hb[0] << endl;cudaFree(da);cudaFree(db);
}

cuda并行归约算法（不断优化过程）相关推荐

cuda编程_CUDA编程入门（四）并行归约算法
这一篇我们一起学习一下如何使用CUDA实现并行归约算法. 首先我们要知道什么是并行归约.并行归约(Reduction)是一种很基础的并行算法,简单来说,我们有N个输入数据,使用一个符合结合律的二元操作 ...
GEMM算法及优化流程详解
目录前言 im2col+GEMM算法简介 GEMM算法优化 optimize1 optimize2 optimize3 前言神经网络前向耗时主要由卷积的耗时决定,参考賈杨青毕业论文,那么如何对卷积 ...
CUDA程序性能优化并行归约
归约算法基本思想是,对于一个输入数组执行某种计算,然后产生一个更小的结果数组.当大量的数进行加和运算时,可以利用归约算法,多线程进行求和运算例如串行实现需要7步,性能比较差成对方式 ...
基于CUDA的N-Body问题并行程序设计及性能优化
目录 N-body问题原理串行代码 CUDA并行程序设计并行的基本思路并行的详细设计 Step1:申请CPU和GPU内存空间并对数据进行初始化和拷贝操作. Step2:设计bodyForce函 ...
短视频推荐算法过程分享，论如何针对推荐算法来优化短视频内容
短视频推荐算法过程分享,论如何针对推荐算法来优化短视频内容相信做短视频的小伙伴一定知道"短视频推荐算法",简单理解就是短视频平台都自有一套推荐机制,决定我们发布的短视频是否可以获 ...
基于串行并行ADMM算法的主从配电网分布式优化控制研究
基于串行并行ADMM算法的主从配电网分布式优化控制研究关键词:ADMM 串行并行算法主动配电网无功优化分布式优化参考文档:非复现,仅参考部分模型: 1)<主动配电网分布式无功优化控制方 ...
水平集图像分割并行加速算法设计与实现(串行、OpenMP、CUDA)——串行实现篇
本次水平集图像分割并行加速算法设计与实现包含:原理篇.串行实现篇.OpenMP并行实现篇与CUDA GPU并行实现篇四个部分.具体各篇章链接如下: 水平集图像分割并行加速算法设计与实现--原理篇水平 ...
【OpenCL】求矩阵所有元素的和——归约算法
任务分析矩阵内求和任务在GPU上实现并不直观.假设每个线程load一个数据并累加到最后的结果标量中,这样的处理过程会使得整个任务退化到串行处理.原因是同一时刻只有一个线程可以操作目的标量,其他线程只 ...
CUDA学习笔记之程序优化
CUDA学习笔记之程序优化标签: cuda优化conflict存储算法数学计算 2010-01-05 17:18 5035人阅读评论(4) 收藏举报分类: CUDA(6) 版权声明:本文为博主 ...

cuda并行归约算法（不断优化过程）

归约

方法

优化一

优化二

优化三

优化四

优化五

cuda并行归约算法（不断优化过程）相关推荐

最新文章

热门文章