【CUDA学习】全局存储器

全局存储器,即普通的显存,整个网格中的任意线程都能读写全局存储器的任意位置。

存取延时为400-600 clock cycles  非常容易成为性能瓶颈。

访问显存时,读取和存储必须对齐,宽度为4Byte。如果没有正确的对齐,读写将被编译器拆分为多次操作,降低访存性能。

多个half-warp的读写操作如果能够满足合并访问,则多次访存操作会被合并成一次完成。

合并访问的条件,GT200放宽了合并访问的条件。

支持对8 bit、16 bit、32 bit、64 bit数据字的合并访问 相应传输32Byte 64Byte 128Byte,大于128Byte,分两次传输。

在一次合并传输的数据中,不要求线程编号和访问的数据字编号相同。

当访问128Byte数据时,如果地址没有对齐到128Byte时,在GT200会产生两次合并访存。根据每个区域的大小,分为两次合并访存,如图所示32Byte和96Byte。

PS:图片来源于网上

关于访存合并以及访存冲突,关键就是要理解,GPU是以half-warp进行访存时,即16个线程一起访问存储器,到这16个线程的访问的地址在同一块区域(指硬件上可以一起传送宽

度)时,并且没有冲突产生时,则这块区域的数据可以被线程同时,提升了访存的效率.

时间: 2024-10-11 00:05:38

【CUDA学习】全局存储器的相关文章

CUDA学习之二:shared_memory使用,矩阵相乘

CUDA中使用shared_memory可以加速运算,在矩阵乘法中是一个体现. 矩阵C = A * B,正常运算时我们运用 C[i,j] = A[i,:] * B[:,j] 可以计算出结果.但是在CPU上完成这个运算我们需要大量的时间,设A[m,n],B[n,k],那么C矩阵为m*k,总体,我们需要做m*n*k次乘法运算,m*(b-1)*k次加法运算,并且是串行执行,总体的复杂度为O(m*n*k) . 矩阵类: 1 class Matrix 2 { 3 public: 4 int cols; /

CUDA学习之一:二维矩阵加法

今天忙活了3个小时,竟然被一个苦恼的CUDA小例程给困住了,本来是参照Rachal zhang大神的CUDA学习笔记来一个模仿,结果却自己给自己糊里糊涂,最后还是弄明白了一些. RZ大神对CUDA关于kernel,memory的介绍还是蛮清楚,看完决定写一个二维数组的加法.如果是C++里的加法,那就简单了,用C[i][j] = A[i][j] +B[i][j]就可以. 1 void CppMatAdd(int A[M][N],int B[M][N],int C[M][N]){ 2 for(int

CUDA学习5 常量内存与事件

当线程束中的所有线程都访问相同的只读数据时,使用常量内存将获得额外的性能提升. 常量内存大小限制为64k. 以下摘自hackairM的博文CUDA学习--内存处理之常量内存(4). 常量内存其实只是全局内存的一种虚拟地址形式,并没有特殊保留的常量内存块.常量内存有两个特性,一个是高速缓存,另一个是它支持将单个值广播到线程束中的每个线程.但要注意的是,对于那些数据不太集中或者数据重用率不高的内存访问,尽量不要使用常量内存. 当常量内存将数据分配或广播到线程束中的每个线程时(注意,实际上硬件会将单次

CUDA学习,第一个kernel函数及代码讲解

前一篇CUDA学习,我们已经完成了编程环境的配置,现在我们继续深入去了解CUDA编程.本博文分为三个部分,第一部分给出一个代码示例,第二部分对代码进行讲解,第三部分根据这个例子介绍如何部署和发起一个kernel函数. 一.代码示例 二.代码解说 申明一个函数,用于检测CUDA运行中是否出错. kernel函数,blockIdx.x表示block在x方向的索引号,blockDim.x表示block在x方向的维度,threadIdx.x表示thread在x方向的索引号. 这里也许你会问,为什么在x方

cuda学习3-共享内存和同步

为什么要使用共享内存呢,因为共享内存的访问速度快.这是首先要明确的,下面详细研究. cuda程序中的内存使用分为主机内存(host memory) 和 设备内存(device memory),我们在这里关注的是设备内存.设备内存都位于gpu之上,前面我们看到在计算开始之前,每次我们都要在device上申请内存空间,然后把host上的数据传入device内存.cudaMalloc()申请的内存,还有在核函数中用正常方法申请的变量的内存.这些内存叫做全局内存,那么还有没有别的内存种类呢?常用的还有共

CUDA学习和总结1

一. 基本概念 1. CUDA 2007年,NVIDIA推出CUDA(Compute Unified Device Architecture,统一计算设备架构)这个编程模型,目的是为了在应用程序中充分利用CPU和GPU各自的优点,实现CPU/GPU联合执行.这种联合执行的需要已经在最新的集中编程模型(OpenCL,OpenACC,C++ AMP)中体现出来了. 2. 并行编程语言和模型 使用比较广泛的是为可扩展的集群计算设计的消息传递接口(Message Passing Interface,MP

CUDA学习ing..

0.引言 本文记载了CUDA的学习过程~刚开始接触GPU相关的东西,包括图形.计算.并行处理模式等,先从概念性的东西入手,然后结合实践开始学习.CUDA感觉没有一种权威性的书籍,开发工具变动也比较快,所以总感觉心里不是很踏实.所以本文就是从初学者的角度,从无知开始探索的过程.当然在学习过程中避免不了出现概念性的理解错误,出现描述模糊不确切的地方还望指出,共勉共勉~ 1.CUDA的概念 2.CUDA的模型 CUDA的运行模型,让host中的每个kernel按照线程网格的方式(Grid)在显卡硬件(

【CUDA学习】共享存储器

下面简单介绍一些cuda中的共享存储器和全局存储器 共享存储器,shared memory,可以被同一块中的所有线程访问的可读写存储器,生存期是块的生命期. Tesla的每个SM拥有16KB共享存储器. 在编程过程中,有静态的shared memory 动态的shared memory 静态的shared memory 在程序中定义 __shared__ type shared[SIZE]; 动态的shared memory 通过内核函数的每三个参数设置大小 extern __shared__

CUDA学习之从CPU架构说起

最近要学习GPU编程,就去英伟达官网下载CUDA, 遇到的第一个问题就是架构的选择 所以我学习的CUDA的第一步是从学习认识CPU架构开始的,x86-64简称x64,是64位版的x86指令集,向前兼容与16位版和32位版的x86架构.x64最初是由AMD于1999年设计完成,AMD首次公开64位集以扩充给x86,称为“AMD64”.后来也被Intel所采用,又被intel 叫做“Intel 64”. 那么ppc64le又指什么呢,下面引用wiki上的两段话“ ppc64 是Linux和GCC开源