【Cuda并行编程之二】Cuda Memory Hierarchy_Cuda内存层次结构

要想编写高效的程序,那么一定要对内存结构有比较深刻的认识,就像C/C++里面的堆内存,栈内存,全局存储区,静态存储区,常量区等。Cuda是并行计算框架,而GPU的内存有限,那么如果想编写高效的Cuda程序,首先要对其内存结构有一个简单的认识。

首先我们先上一张图,然后通过解释一些名词和代码来进行解释。

各种存储器比较:

存储器  位置 拥有缓存 访问权限 变量生存周期
register GPU片内 N/A device可读/写 与thread相同
local memory 板载显存 device可读/写 与thread相同
shared memory  GPU片内 N/A device可读/写 与block相同
constant memory 板载显存 device可读,host可读写 可在程序中保持
texture memory  板载显存 device可读,host可读写 可在程序中保持
global memory  板载显存 device可读写,host可读写 可在程序中保持
host memory  主机内存 host可读写  可在程序中保持
pinned memory 主机内存 host可读写 可在程序中保持

registers:寄存器。它是GPU片上告诉缓存器,执行单元可以以极低的延迟访问寄存器。寄存器的基本单元是寄存器文件(register file),每个寄存器文件大小为32bit。寄存器文件数量虽然客观,但是平均分给并行执行的线程,每个线程拥有的数量就非常有限了。编程时,不要为每个线程分配过多的私有变量。下面程序中,aBegin,aEnd,aStep,a等变量都是寄存器变量,每个线程都会维护这些变量。

__global__ void registerDemo(float *B,float *A ,int wA)
{
    int aBegin = wA*BLOCK_SIZE * blockIdx.y;
	int aEnd = aBegin + wA - 1;
	int aStep = BLOCK_SIZE;

	for(int a=aBegin;a<=aEnd;a+=aStep)
	{
	    //...
	}
}

local memory:局部存储器。对于每个线程,局部存储器也是私有的。如果寄存器被消耗完,数据将被存储在局部存储器中。如果每个线程用了过多的寄存器,或声明了大型结构体或数组,或者编译期无法确定数组的大小,线程的私有数据就有可能被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据将被保存在显存中,而不是片上的寄存器或者缓存中,因此对local
memory的访问速度比较慢。

shared memory:共享存储器。共享存储器也是GPU片内的告诉存储器。它是一个块可以被同一block中的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快。是实现线程间通信的延迟最小的方法。共享存储器可用于实现多种功能,如用于保存共用的计数器(例如计算循环迭代次数)或者block的公共结果(例如规约的结果)。

static variable使用shared memory:

#include<iostream>
#include<stdio.h>

#if 1
__global__ void example(float *u)
{
	int i = threadIdx.x;
	<strong>__shared__ int tmp[4];</strong>
	tmp[i] = u[i];
	u[i] = tmp[i] * tmp[i] + tmp[3-i] ;
}
#endif

#if 1

int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u ;
	size_t size = 4*sizeof(float);

	cudaMalloc(&dev_u , size);
	cudaMemcpy(dev_u,host_u,size,cudaMemcpyHostToDevice);

	example<<<1,4>>> (dev_u);

	cudaMemcpy(host_u , dev_u , size , cudaMemcpyDeviceToHost);

	cudaFree(dev_u);

	for(int i=0;i<4;i++)
		printf("%f\n",host_u[i]);
	return 0;
}

#endif

dynamic variable使用shared memory:

#include<iostream>
#include<stdio.h>

<strong>extern __shared__ int tmp[];</strong>

__global__ void example(float *u)
{
	int i = threadIdx.x;
	tmp[i] = u[i];
	u[i] = tmp[i] * tmp[i] + tmp[3-i];
}

int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u;
	<strong>size_t size = 4*sizeof(float);</strong>

	cudaMalloc(&dev_u,size);
	cudaMemcpy(dev_u , host_u ,size , cudaMemcpyHostToDevice);
	example<<<1,4,<strong>size</strong>>>>(dev_u);

	cudaMemcpy(host_u, dev_u,size,cudaMemcpyDeviceToHost);
	cudaFree(dev_u);
	for(int i=0;i<4;i++)
		printf("%f ",host_u[i]);
	return 0;
}

global memory:全局存储器位于显存(占据了显存的绝大部分),CPU、GPU都可以进行读写访问。整个网格中的任意线程都能读写全局存储器的任意位置由于全局存储器是可写的。全局存储器能够提供很高的带宽,但同时也具有较高的访存延迟。显存中的全局存储器也称为线性内存。线性内存通常使用cudaMalloc()函数分配,cudaFree()函数释放,并由cudaMemcpy()进行主机端与设备端的数据传输。

此外,也可以使用__device__关键字定义的变量分配全局存储器,这个变量应该在所有函数外定义,必须对使用这个变量的host端和device端函数都可见才能成功编译。在定义__device__变量的同时可以对其赋值。

static variable使用global memory:

global_mem_static.cu:

#include<stdio.h>
#include<iostream>

__device__ float devU[4];
__device__ float devV[4];

//__global__ function
__global__ void addUV()
{
	int i = threadIdx.x;
	devU[i] += devV[i];
}

int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};

	int size = 4* sizeof(float);

	//cudaMemcpyToSymbol:将数据复制到__constant__或者__device__变量中
	//cudaMemcpyFromSymbol:同上相反
	//cudaMalloc:在设备端分配内存
	//cudaMemcpy:数据拷贝
	//cudaFree():内存释放
	//cudaMemset():内存初始化
	cudaMemcpyToSymbol(devU,hostU,size,0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(devV,hostV,size,0,cudaMemcpyHostToDevice);

	addUV<<<1,4>>>();

	cudaMemcpyFromSymbol( hostU,devU,size,0,cudaMemcpyDeviceToHost );

	for(int i=0;i<4;i++)
		printf("hostU[%d] = %f\n",i,hostU[i]);
	return 0;
}

结果:

dynamic variable使用global memory:

global_mem_dynamic.cu:

#include<iostream>
#include<stdio.h>

__global__ void add4f(float *u , float *v)
{
	int i = threadIdx.x;
	u[i] += v[i];
}

void print(float * U ,int size)
{
	for(int i=0;i<4;i++)
	{
		printf("U[%d] = %f\n",i,U[i]);
	}
}

int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};

	float * devU ;
	float * devV ;
	int size = sizeof(float) * 4;

	//在设备内存上分配空间
	cudaMalloc( &devU,size );
	cudaMalloc( &devV,size );
	//数据拷贝
	cudaMemcpy( devU ,hostU ,size ,cudaMemcpyHostToDevice );
	cudaMemcpy( devV ,hostV ,size ,cudaMemcpyHostToDevice );

	add4f<<<1,4>>> (devU,devV);
	//数据返回
	cudaMemcpy(hostU,devU,size,cudaMemcpyDeviceToHost);

	print(hostU,size);
	//释放空间
	cudaFree(devV);
	cudaFree(devU);
	return 0;
}

结果:

host memory : 主机端内存,即CPU对应的我们普通意义上的内存。主机端内存分为两种:可分页内存(pageable memory)和页锁定(page-locked 或pinned)内存。可分页内存即为通过操作系统API(malloc(),new())分配的存储器空间:而页锁定内存始终不会被分配到低俗的虚拟内存中,能够保证存在于屋里内存中,并且能够通过DMA加速与设备端的通信。

constant memory:常数存储器。它是只读的地址空间。常熟存储器中的数据位于显存,但拥有缓存加速。常数存储器的空间较小,在Cuda程序中用于存储需要频繁访问的只读参数。当来自同一half-warp的线程访问常数存储器中的同一数据时,如果发生缓存命中,那么只需要一个周期就可以获得数据。常数存储器有缓存机制,用以节约带宽,加快访问速度。每个SM拥有8KB的常数存储器缓存。常数存储器是只读的,因此不存在缓存一致性问题。

constant memory的使用:

#include<iostream>

using namespace std;

__constant__ int devVar = 100;

__global__ void xminus(int *a)
{
	int i = threadIdx.x;
	a[i] = devVar+i;
}

int main()
{
	int *h_a = (int*)malloc(4*10) ;
	int *d_a ;
	cudaMalloc(&d_a, 4*10) ;
	cudaMemset(d_a, 0, 40) ;

	xminus<<<1,4>>>(d_a);

	cudaMemcpy(h_a, d_a, 4*10, cudaMemcpyDeviceToHost) ;

	for(int i = 0; i < 4 ; i++)
		cout << h_a[i] << " " ;
	cout << endl ;
}

结果:

texture memory:纹理存储器。在此不做过多介绍。

再结合一个图解释一下GPU里面的一些概念作为结尾:

Thread:线程。即一个GPU核心处理的单个单元。

Block:块。一个块里面有多个线程组成。block是软件级别的概念。

Grid:格。即二维的格,由多个block组成。

SM:Streaming multiprocessor。一个GPU里面有多个SM,一个SM里面有多个SP(streaming processor),是硬件级别的概念。

Warp:Warp是SM调度和执行的基本单位。warp是32个并列的线程,软件级别的概念。CPU上执行一条指令时候都是一个线程的,但是GPU则是以warp为单位。SM执行一条指令,那么这条指令使得32个线程同时执行,而每个线程都会操作自己的内存处理自己相应的数据,因此就达到了执行一条指令操作多个数据,也就是SIMD(single instruction
and multiple data)

注明出处:http://blog.csdn.net/lavorange/article/details/20465869

时间: 12-09

【Cuda并行编程之二】Cuda Memory Hierarchy_Cuda内存层次结构的相关文章

【CUDA并行编程之八】Cuda实现Kmeans算法

本文主要介绍如何使用CUDA并行计算框架编程实现机器学习中的Kmeans算法,Kmeans算法的详细介绍在这里,本文重点在并行实现的过程. 当然还是简单的回顾一下kmeans算法的串行过程: 伪代码: 创建k个点作为起始质心(经常是随机选择) 当任意一个点的簇分配结果发生改变时 对数据集中的每个数据点 对每个质心 计算质心与数据点之间的距离 将数据点分配到距其最近的簇 对每一个簇,计算簇中所有点的均值并将均值作为质心 我们可以观察到有两个部分可以并行优化: ①line03-04:将每个数据点到多

【CUDA并行编程之七】数组元素之和

现在需要求得一个数组的所有元素之和,之前感觉似乎不太可能,因为每个线程只处理一个元素,无法将所有元素联系起来,但是最近学习了一段代码可以实现,同时也对shared memory有了进一步的理解. 一.C++串行实现 串行实现的方法非常之简单,只要将所有元素依次相加就能够得到相应的结果,实际上我们注重的不是结果,而是运行的效率.那么代码如下: array_sum.cc: #include<iostream> #include<stdio.h> #include "kmean

【CUDA并行编程之四】矩阵相乘

前面介绍了基本的Cuda编程的相关知识,那么这一篇在此基础之上来看看GPU在处理数据计算上的高效能,我们拿矩阵相乘来作为例子. 1.CPU上执行矩阵相乘以及性能. 在CPU上进行矩阵相乘运算的代码: mat_mul.cc: <span style="font-family:Microsoft YaHei;font-size:18px;">//a[i]*b[i] + c[i] = d[i] #include<iostream> #include<vector

cuda并行编程之求解ConjugateGradient(共轭梯度迭代)丢失dll解决方案

在进行图像处理过程中,我们经常会用到梯度迭代求解大型现在方程组:今天在对奇异矩阵进行求解的时候,出现了缺少dll的情况: 报错如下图: 缺少cusparse32_60.dll 缺失cublas32_60.dll 解决方案: (1)将cusparse32_60.dll和cublas32_60.dll直接拷贝到C:\Windows目录,但这样在一直的时候,还会出现同样错误,为了避免麻烦,最好采用方法(2) (2)将cusparse32_60.dll和cublas32_60.dll拷贝到你所在项目的文

CUDA C编程入门-介绍

CUDA C编程入门-介绍 1.1.从图形处理到通用并行计算 在实时.高清3D图形的巨大市场需求的驱动下,可编程的图形处理单元或者GPU发展成拥有巨大计算能力的和非常高的内存带宽的高度并行的.多线程的.多核处理器.如图1和图2所示. 图 1 CPU和GPU每秒的浮点计算次数 图 2 CPU和GPU的内存带宽 在CPU和GPU之间在浮点计算能力上的差异的原因是GPU专做密集型计算和高度并行计算-恰好是图形渲染做的-因此设计成这样,更多的晶体管用于数据处理而不是数据缓存和流控制,如图3所示. 图 3

CUDA编程(二) CUDA初始化与核函数

CUDA编程(二) CUDA初始化与核函数 CUDA初始化 在上一次中已经说过了,CUDA安装成功之后,新建一个工程还是十分简单的,直接在新建项目的时候选择NVIDIA CUDA项目就可以了,我们先新建一个MyCudaTest 工程,删掉自带的示例kernel.cu,然后新建项,新建一个CUDA C/C++ File ,我们首先看一下如何初始化CUDA,因此我命名为InitCuda.cu 首先我们要使用CUDA的RunTime API 所以 我们需要include cuda_runtime.h

CUDA C编程入门-编程模型

这章节介绍CUDA编程模型的主要的概念. 2.1.kernels(核函数) CUDA C扩展了C语言,允许程序员定义C函数,称为kernels(核函数).并行地在N个CUDA线程中执行N次. 使用__global__说明符声明一个核函数,调用使用<<<...>>>,并且指定执行的CUDA线程数目.执行的每个线程都有一个独一的ID,在核函数中可以通过变量threadIdx获取. 例子,两个向量的加,A加B,并把结果存入C,A.B和C的长度为N. __global__ vo

GPGPU OpenCL/CUDA 高性能编程的10大注意事项

转载自:http://hc.csdn.net/contents/content_details?type=1&id=341 1.展开循环 如果提前知道了循环的次数,可以进行循环展开,这样省去了循环条件的比较次数.但是同时也不能使得kernel代码太大. 1 #include 2 using namespace std; 3 4 int main(){ 5 int sum=0; 6 for(int i=1;i<=100;i++){ 7 sum+=i; 8 } 9 10 sum=0; 11 fo

CUDA 标准编程模式

前言 本文将介绍 CUDA 编程的基本模式,所有 CUDA 程序都基于此模式编写,即使是调用库,库的底层也是这个模式实现的. 模式描述 1. 定义需要在 device 端执行的函数.( 函数声明前加 _golbal_ 关键字 ) 2. 在显存中为待运算的数据以及需要存放结果的变量开辟显存空间.( cudaMalloc 函数实现 ) 3. 将待运算的数据传输进显存.( cudaMemcpy,cublasSetVector 等函数实现 ) 4. 调用 device 端函数,同时要将需要为 devic