CUDA循环展开

串行循环展开

loop unrolling是一种牺牲程序的尺寸来加快执行速度的优化方法。拿数组来说，数组的数据在内存中是连续存储的，每次取数据的时候可以一次抓取相邻的多个数据，从而减少从内存中读取数据的时间，优化程序。例子：

假设n可以被4整除

未展开程序for(int i = 0; i < n; i++){c[i] = a[i] + b[i];}4层展开程序for(int i = 0; i < n/4; i++){c[i] = a[i] + b[i];c[i+1] = a[i+1] + b[i+1];c[i+2] = a[i+2] + b[i+2];c[i+3] = a[i+3] + b[i+3];}

GPU并行循环展开

在GPU上，同样的也可以进行循环展开优化, 更重要的是展开可以减少warp总得分支，下面是一个简单的程序，将value写入out的N个元素中。

global_write函数未展开版:template <class T>__global__ void Global_write(T*out,T value,size_t N){for ( size_t i = blockIdx.x*blockDim.x+threadIdx.x;i < N;i += blockDim.x*gridDim.x ) {out[i] = value;}global_write函数展开,展开n层  template<class T, const int n>__global__ void Global_write(T* out, T value, size_t N){size_t i;for(i = n*blockDim.x*blockIdx.x + threadIdx.x;i < N - n*blockDim.x*blockIdx.x;i += n*gridDim.x*blockDim.x;)for(int j = 0; j < n; i++){size_t index = i + j * blockDim.x;outp[index] = value;}// 为了不在循环里加入控制语句，将最后一次循环单独的写for ( int j = 0; j < n; j++ ) {size_t index = i+j*blockDim.x;if ( index<N ) out[index] = value;}}

向量运算例子
计算 out = alpha * x + beta * y

一般形式:__global__ voidsaxpyGPU(float *out,const float *x,const float *y,size_t N,float alpha，float beta ){for ( size_t i = blockIdx.x*blockDim.x +threadIdx.x;i < N;i += blockDim.x*gridDim.x ) {out[i] = alpha * x[i]+beta * y[i];}}unroll:template<const int n>__device__ void saxy_unrolled(float* out,const float* px,const float* py,size_t N,float alpha, float beta){float x[n],y[n];size_t i;for(i = n*blockDim.x*blockIdx.x + threadIdx.x;i < N - n*blockDim.x*blockIdx.x; i += gridDim.x*blockDim.x){for(int j = 0; j < n; j++){size_t index = i + j * blockDim.x;x[j] = px[index];y[j] = py[index];}for( int j = 0; j < n; j++){size_t index = i + j*blockDim.x;out[index] = alpha * x[j] + beta * y[j];}  }for ( int j = 0; j < n; j++ ) {for ( int j = 0; j < n; j++ ) {size_t index = i+j*blockDim.x;if ( index<N ) {x[j] = px[index];y[j] = py[index];} }for ( int j = 0; j < n; j++ ) {size_t index = i+j*blockDim.x;if ( index<N ) out[index] = alpha*x[j]+y[j];}}__global__ voidsaxpyGPU( float *out, const float *px, const float *py,size_t N,float alpha ){saxpy_unrolled<4>( out, px, py, N, alpha );}

在cuda-handbook中说这样的优化能够提高 10% 的效率。

[berkeley并行循环展开课件](http://www.cs.berkeley.edu/~volkov/volkov11-unrolling.pdf)

CUDA笔记2-循环展开相关推荐

cuda笔记-一个Block多线程求卷积
最近在学cuda,发现自己数学方面的知识不太够,C语言的知识也有待加强. 这里记录个笔记对矩阵求卷积. 逻辑是这样的: 1. 先CUDA生成一个16*16的矩阵: 2. 将这16*16的矩阵,外面包一 ...
[转]Cuda笔记【1】GPU计算DEMO
声明本文是看小破站某cuda入门教程留下来的笔记,多上PPT上内容,夹杂一点自己的理解,和代码注释教程地址:https://www.bilibili.com/video/av74148375 gi ...
cuda笔记-GPU多线程的奇偶排序
首先说明奇偶排序: 算法的思路是先排奇数序号的相邻2个,或者偶数序号的相邻两个,然后一直到序列有序为止,如下代码: #include<stdio.h> #include<stdlib ...
cuda笔记-初始化矩阵及thread,block,grid概念
thread:一个CUDA的并行程序会被许多threads来执行: block:多个threads组成一个block,同一个block中threads可以使用_syncthreads()同步,也可以通 ...
cuda笔记-第一个cuda程序
这里先说明下一些基本概念: 释放GPU中的内存cudaFree() CUDA函数的定义: __global__:定义在GPU上,可以在CPU上调用的函数: __device__:定义在GPU上,由GP ...
CUDA笔记-卷积计算
1. 一维卷积一维卷积的计算是源数据根据卷积核将每个元素周围的数值与卷积核加权乘积的总和:卷积核一般是奇数个: 下图是还有16个元素的数组作为源数据,5个元素的卷积核: 计算卷积的过程如下图,每个节 ...
CUDA笔记--GPU的结构与SM（流处理器）结构
GPU的结构简单的可以分为: 一个连接GPU和PCIe总线的主机接口: 0~2个复制引擎: 一个连接GPU与GPU内存设备的DRAM接口: 一定数目的TPC或者GPC(纹理处理集群或图形处理集群),每 ...
cuda笔记-流的使用（定义、创建、消耗、同步）
定义流: cudaStream_t stream; 创建流: cudaStreamCreate(cudaStream_t *s) 销毁流: cudaStreamDestroy(cudaStream_t ...
CUDA入门3.2——使用CUDA实现鱼眼转全景图（CUDA环节）1227更
算法算法借鉴了Converting a fisheye image into a panoramic, spherical or perspective projection,核心内容如下: Sof ...

CUDA笔记2-循环展开

CUDA循环展开

串行循环展开

GPU并行循环展开

向量运算例子
计算 out = alpha * x + beta * y

CUDA笔记2-循环展开相关推荐

最新文章

热门文章

CUDA笔记2-循环展开

CUDA循环展开

串行循环展开

GPU并行循环展开

向量运算例子 计算 out = alpha * x + beta * y

CUDA笔记2-循环展开相关推荐

最新文章

热门文章

向量运算例子
计算 out = alpha * x + beta * y