darkstorm2111203的专栏

CUDA纹理存储器的特性及其使用

张舒 — Wed, 24 Jun 2009 20:23:00 +0800

2.3.5 纹理存储器
纹理存储器（texture memory）是一种只读存储器，由GPU用于纹理渲染的的图形专用单元发展而来，因此也提供了一些特殊功能。纹理存储器中的数据位于显存，但可以通过纹理缓存加速读取。在纹理存储器中可以绑定的数据比在常量存储器可以声明的64K大很多，并且支持一维、二维或者三维纹理。在通用计算中，纹理存储器十分适合用于实现图像处理或查找表，并且对数据量较大时的随机数据访问或者非对齐访问也有良好的加速效果。
纹理存储器在硬件中并不对应一块专门的存储器，而实际上是牵涉到显存、两级纹理缓存、纹理抓取单元的纹理流水线。纹理存储器提供了地址映射、数据滤波、缓存等功能，这些功能都是围绕着纹理渲染的需求设计的。关于GPU纹理流水线的介绍可以参考本书3.3.3节。在CUDA编程模型中，纹理缓存是透明的，编程人员不用去了解它的实现机制。
从CUDA的内核函数访问纹理存储器的操作被称为纹理拾取(texture fetching)。纹理拾取使用的坐标与数据在显存中的地址可以不同，两者通过纹理参照系（texture reference）约定从数据的地址到

通用计算程序在Tesla架构上的执行

张舒 — Sun, 26 Apr 2009 03:06:00 +0800

节选自正在写的书稿，还没有配图。这一部分是第三章硬件介绍的一部分，在之前的小节里已经介绍了显卡的组成和一般知识，以及GPU的架构简介。这一节专门介绍CUDA程序如何映射到硬件上，希望对大家有所帮助。

由nvcc生成的通用计算程序分为主机端程序和设备端程序两部分。那么，一个完整的CUDA程序是如何在CPU和GPU上执行的呢？在这一节，我们不仅将介绍CUDA的编程模型如何映射到硬件上，还会介绍GPU的硬件设计如何对CUDA程序效率产生影响。
通常，在计算开始前，需要将要计算的数据通过API从内存拷贝到显存中，再在计算结束后将数据从显存拷贝回内存。通过CUDA API的存储器管理功能进行数据传输，不需要SPA中的运算参与。CPU先通过存储器管理API在显存上开辟空间，将内存中的数据由北桥经过PCI-E总线传到显存中。在bandwidthTest例子中我们介绍过，主机端内存可以分为两种：pinned(page-locked)或者pageable memory。一般的操作系统使用了虚拟内存和内存分页管理，这种设计在带来种种好处的同时，也使得新开辟的空

CUDA SDK 例子分析（3）：reduction

张舒 — Sat, 28 Feb 2009 07:57:00 +0800

/* Parallel reduction kernels*/#ifndef _REDUCE_KERNEL_H_#define _REDUCE_KERNEL_H_#include #include "sharedmem.cuh"#ifdef __DEVICE_EMULATION__#define EMUSYNC __syncthreads()#else#define EMUSYNC#endif// Macros to append an SM version identifier to a function name// This allows us to compile a file multiple times for different architecture// versions// The second macro is necessary to e

Nvidia GT200:详解并行处理器（三）CUDA存储器模型

樟树 — Sun, 26 Oct 2008 07:28:00 +0800

CUDA存储器模型
除了执行模型以外，CUDA也规定了存储器模型（如图2所示）和一系列用于主控CPU与GPU间通信的不同地址空间。图中红色的区域表示GPU片内的高速存储器，橙色区域表示DRAM中的的地址空间。

图2 CUDA存储器模型
首先，是最底层的寄存器(register,REG)。对每个线程来说，寄存器都是私有的--这与CPU中一样。如果寄存器被消耗完，数据将被存储在本地存储器(local memory)。本地存储器对每个线程也是私有的，但是数据时被保存在帧缓冲区DRAM中，而不是片内的寄存器或者缓存中。线程的输入和中间输出变量将被保存在寄存器或者本地存储器中。
然后是用于线程间通信的共享存储器。共享存储器是一块可以被同一block中的所有thread（上节提到过，一个block最多可以有512个thread）访问的可读写存储器。访问共享存储器几乎和访问寄存器一样快，是实现线程间通信的延迟最小的方法。共享存储器可以实现许多不同的功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公用结果（例

Nvidia GT200：详解并行处理器（二）CUDA执行模型

樟树 — Sun, 26 Oct 2008 07:14:00 +0800

CUDA 执行模型Nvidia的并行编程模型被命名为CUDA（Computing Unified Device Architecture，统一计算架构模型）。CUDA的基本思想是尽量得开发线程级并行(Thread Level Parallel)，这些线程能够在硬件中被动态的调度和执行。CUDA编程模型的重点是将CPU做为终端(Host)，而GPU做为服务器(Server)或协处理器(Coprocessor)，或者设备（Device），从而让GPU来运行一些能够被高度线程化的程序。所以，GPU只有在计算高度数据并行任务时才能发挥作用。在这类任务中，需要处理大量的数据，数据的储存形式类似于规则的网格，而对这写数据的进行的处理则基本相同。这类数据并行问题的经典例子有：图像处理，物理模型模拟（如计算流体力学），工程和金融模拟与分析，搜索，排序。而需要复杂数据结构的计算如树，相关矩阵，链表，空间细分结构等，则不适用于使用GPU进行计算。找到程序中的计算并行度后，就能将一部分程序移植到GPU上。运行在GPU上的程序被称为Kernel(核)。核并不是完整的程序，而是整个程序

Nvidia GT200:详解并行处理器（一）简介

樟树 — Sun, 26 Oct 2008 07:08:00 +0800

注：本文是Nvidia's GT200:into the parallel processor的中文译文，原文作者为David Kanter博士翻译本文的目的旨在说明CUDA是如何在硬件上实现的，帮助编程人员编写和优化CUDA程序简介过去十年中，计算领域出现了一个新的发展趋势：GPU通用计算。由Intel,IBM,SUN,AMD和富士通生产的通用CPU虽然有了很大发展，但性能提高速度却已经不能与与上世纪八十年代末九十年代初相比。单线程处理的性能在很大程度上受到了限制。这些限制一方面来自于通用计算程序中过低的指令级并行;另一方面来自于“功率墙(Power Wall)”--集成电路的功率消耗的物理限制。摩尔定律能够为处理器提供数以亿计的晶体管，但在为了运行单线程程序而设计的处理器中，这些晶体管绝大多数都被用于制造高速缓存(Cache)。这样做虽然能把处理器功耗控制在合理的范围内，却阻碍了性能的进一步提高。与此同时，GPU（图形处理单元）却有效的利用了数量巨大的晶体管资源。由于图形渲染过程的高度并行性，GP

CUDA SDK例子分析（2）乾坤挪移：transpose

张舒 — Fri, 22 Aug 2008 20:50:00 +0800

Transpose是一个矩阵转置的例子，通过两个功能相同的核函数：transpose()

CUDA SDK 例子分析（1）万剑归宗：template

张舒 — Fri, 22 Aug 2008 20:49:00 +0800

Cuda sdk 例子分析但凡武林高手，都讲究“内修一口气，外练筋骨皮”。对于程序员来说，掌握思想和算法是必要的“内功”；

cuda心得体会

darkstorm2111203 — Wed, 16 Jul 2008 20:49:00 +0800

访问global虽然仍然需要最多数百个时钟周期，但是总的来说比cpu的内存还是快多了...如果使用得当的话。CUDA是符合SIMD-PRAM模型的，也就是n台功能相同的处理机（block看成处理机,每个thread看成处理机的一部分更好些），一个容量无限大的共享处理器M。其中M就是global memory，虽然在CUDA中也是有限的，但是比起可怜的shared还是多很多了。
global作为共享存储器，最重要的功能就是用来进行block之间的通信。这里能不能叫做通信其实有问题，因为block之间实际上没有通信...CUDA上的通信实际上就是刷新，运行完一个kernel，把所有的需要与其他block交换的数据都写进global，然后再launch一个kernel读入上次生成的数据...如此往复循环。
如果kernel的输入与输出是同一个矩阵，就是一个原地（in-place）操作，下个kernel还是读取这个矩阵。
如果输入矩阵与输出矩阵不同，而且需要反复操作这两个矩阵，就构成了一个乒乓操作，例如第一次kernel使用ping为输入，pong为输出，下次