NVIDIA Fermi GF100 及 GF1XX 架构讨论
Fermi 架构的初步概况已经发布，大家可以就相关的话题在本主题中展开讨论（原来的讨论串作为猜测性质，现在已经关闭，你可以在这个连接里回顾：http://we.pcinlife.com/thread-969042-1-1.html）。http://www.pcinlife.com/article_photo/fermi_preview/fermi_architecture-s.png
Fermi 的体系结构亮点如下：
[*]第三代 Streaming Multiprocessor（流式多处理器）[*]每个 SM 有 32 个内核，4 倍于 GT200 的 SM。[*]双精度浮点性能达到 GT200 的 8 倍。[*]配备双 Warp Scheduler（Warp 调度器），能每个周期对两个宽度为 32 线程的 Warp 进行排程和分发。 [*]每个 SM 有 64kB 可配置为 shared memory 和 cache 的随机存取内存[*]第二代并行线程执行（PTX）ISA[*]统一寻址空间，提供完全 C++ 支持[*]针对 OpenCL 和 DirectCompute 作最佳化[*]完全的 IEEE 754-2008 32-bit/64-bit 精度支持[*]透过 Predication（论断）提升性能[*]改进的内存子系统[*]拥有可配置 L1 cache 和统一化 L2 cache 的 NVIDIA 并行数据高速缓存（PDC）[*]第一枚提供 ECC 内存支持的 GPU[*]显著提升原子内存操作性能[*]NVIDIA GigaThread 3.0 引擎[*]应用程序上下文切换性能达到 10 倍于上代产品[*]同时执行多个核心程序（kernel，指的是程序中在 GPU 上执行的功能）[*]线程块（CTA）乱序执行[*]双重叠式内存传输引擎NVIDIA 官方白皮书：
http://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIAFermiArchitectureWhitepaper.pdf
本站文章：
http://www.pcinlife.com/article/graphics/2009-09-29/1254197427d834.html
推荐阅读 Realworldtech 的报道：
Inside Fermi: Nvidia's HPC Push Predication（论断）是什么？分支预测？ http://www.nvidia.com/content/PDF/fermi_white_papers/P.Glaskowsky_NVIDIAFermi-TheFirstCompleteGPUComputingArchitecture.pdf
按照这篇报道，Fermi 的 DRAM ECC 实现机制和传统 CPU 每 8-bit 增加一个位元的方式有些不一样，是一种专利的方式 ：
Fermi is the first GPU to provide ECC (error correcting code) protection for
DRAM; the chip’s register files, shared memories, L1 and L2 caches are also ECC
protected. The level of protection is known as SECDED: single (bit) error correction
double error detection. SECDED is the usual level of protection in most ECC‐r
equipped systems.
Fermi’s ECC protection for DRAM is unique among GPUs; so is its
implementation. Instead of each 64゜it memory channel carrying eight extra bits fo
ECC information, NVIDIA has a proprietary (and undisclosed) solution for packing
the ECC bits into reserved lines of memory. Predication（论断）是什么？分支预测？
voodoo12345 发表于 2009-10-2 16:30 http://we.pcinlife.com/images/common/back.gif
Prediction（预测） 和 Predication（论断） 是两回事情，虽然有点相似。
http://en.wikipedia.org/wiki/Branch_predication
http://en.wikipedia.org/wiki/Branch_prediction
branch predication 的话，你是 branch 的两边都跑，当确定一个方向后（用 predication 算出一个"筛子"），就只拿这个方向的结果，另一方向的就抛弃掉。
branch prediction 就是猜测哪个方向会发生状况，然后就卯足劲跑这个方向，猜中了就中奖，猜不中就重新跑。 从这个Fermi100的推出可以看出NV今后的方向，他们走的这一步算是很激进的。如果保守的话他们是可以通过重点发展消费图形处理慢慢蚕食的。不过基于目前的状况来看这一步激进路线前途很坎坷。 这里是架构讨论区，着重讨论体系架构以及相关的技术讨论，所以如果是发表市场策略看法，请到下面的显卡区发表。 问个关于C++的问题，这个GPU对C++的完全支持，是应用在开发者利用GPU编写C++代码提高开发效率呢？还是应用在用户执行C++程序时GPU提高程序运行效率呢？ 能否达到和晶体管数量相称的游戏性能增长？或者说，游戏性能增长相对晶体管性能增长，打了几折？ 问个关于C++的问题，这个GPU对C++的完全支持，是应用在开发者利用GPU编写C++代码提高开发效率呢？还是应用在用户执行C++程序时GPU提高程序运行效率呢？
voodoo12345 发表于 2009-10-2 19:23 http://we.pcinlife.com/images/common/back.gif
应该是需要 CUDA 的编译器编译为 PTX 2.0 代码，然后驱动编译为本机代码来执行。
例如，以前的 GPU 是没有 function 的，函式调用都是 inline 的，不是真正的函式。
Fermi 实现了这些特性的支持。 有中文资料没？ 能否达到和晶体管数量相称的游戏性能增长？或者说，游戏性能增长相对晶体管性能增长，打了几折？
westlee 发表于 2009-10-2 20:45 http://we.pcinlife.com/images/common/back.gif
目前还没有 gaming 方面更详细的资料看到，不过我想得看具体的游戏吧，当前的许多游戏都不能从 cache 获益，但是并不意味着那些采用了 gpu computing 的游戏不能因此而获益。 本帖最后由 Prescott 于 2009-10-3 00:20 编辑
真正的亮点就是
1. 大幅增加DP单元
2. 统一地址空间
3. ECC
可以说完全是针对GPGPU的。
对于3D图形渲染性能的改进就是double SM的SIMD宽度到16，这个带来的渲染性能提高还有待检验。肯定是不到一倍的，我估计乐观一点在60-80%之间。 指令指针在以前的CUDA肯定也能实现(其实就是一个jmp指令), 但是统一地址空间是新的.这样做有好有坏.
1: 指针长度翻番, 32bit增加到64bit, 意味着寄存器/代码长度消耗增大
2: 即便地址空间统一了, 也只是逻辑上放在一起. 不能写的东西, 我估计还是不能写. 比如self-modify code.
3: 对于性能而言, 这个特性并没有任何帮助.
至于SM长度x4, P大可能理解错了. 因为WARP的长度还是和原来一样, 所以只是硬件上出现了变化, 逻辑上完全没有变. 并且, 虽然SM长度加大了, SM个数却减少了. 因此我个人认为这个特性不会对性能有什么影响. 决定的东西还是SP的数量.
对于性能影响对大的东西, 我认为是Atomic的性能提升. 这个可能很多人都忽略了. 其他什么C++虚函数之类的特性, 简单地说, 只会让性能变得更差. 而且, 大部分的改进完全是软件编译器前端的.
这个尝试确实很勇敢，还是要鼓掌的。NV看来要豪赌一把了。 另外， 统一地址空间很容易诱惑程序员写出巨慢的代码 。C++也是。
这样我想起了PS3的统一地址空间，程序员不小心用CPU读写到显存出现性能巨幅下降，搞不清楚状况骂娘的情形。 RacingPHT，我同意你关于性能的意见。这些不会有性能方面的提高，但是对于提高通用编程能力是至关重要的。
关于统一寻址：在实现统一寻址之前，很多C/C++语法根本无法编译成PTX，比如对一个float*的defererence，在编译的时候，编译器很难知道这个指针到底指向什么东西，它这次有可能指向share memory，下次就有可能指向global memory，这个时候，指令生成就成了一个大问题，因为这是完全不同的指针。除非对指针加上修饰，比如 __share__ float*，然后限制不同指针互相赋值也不能相互转换，这显然限制更大，func(__share__ float*)和func(__global__ float*)变成了不同的函数。[wacko> 所以，统一寻址对于提高可编程能力至关重要。至于原先不能写的东西，这个和C/C++的语义是符合的，const float *本来就是不能写，你非要强制转成float*然后写自然是后果自负，Fermi当然是理直气壮的死给你看。可以说CUDA从此之后可以有一个很大的进步，虽然现在还没有发布新的CUDA版本，但是可以值得期待。
关于SM的宽度：对于我来说，SM（Stream Multiprocessor）中SP的个数就是SIMD宽度。所以我的理解是，Fermi现在有32个SM，每个SM的宽度是16(SP)，前面我好像弄错了。我还没来的及弄明白Dual Wrap Scheduler是怎么回事，看起来像是把16个SP又分成了两组。Wrap只不过是一个CUDA上的概念，应该和硬件实现无关。
我对Atomic操作在图形学中的重要性没有概念，但是对于Fermi针对的高性能计算市场，应该不是特别重要。 另外， 统一地址空间很容易诱惑程序员写出巨慢的代码 。C++也是。
这样我想起了PS3的统一地址空间，程序员不小心用CPU读写到显存出现性能巨幅下降，搞不清楚状况骂娘的情形。
RacingPHT 发表于 2009-10-2 23:00 http://we.pcinlife.com/images/common/back.gif
这个，所有硬件都有这样的问题，或多或少，CPU一个Core比GPU Core大那么多，不都是为了避免这些情况嘛。
这方面Nehalme >> LRB > Fermi P大的说法我赞同 , 关于指针语义上确实是统一了. 但是即便最终实现了C++的语法, 也大概还是486的级别的性能. 最终, 程序员还是要非常清楚里面究竟是什么. 因此, C++这种东西, 可能是减少了一开始的门槛, 但是最终要获得高性能, 还是一步也少不了(没有人愿意写出来的CUDA程序比CPU C++还慢, 虽然这是极其常见的情况). 从编程语言上掩盖这种东西, 利弊参半, 因为我不相信GPU的直接C++移植能获得可用的性能. 这个和SPU的C++状况不同. CELL的SPU上的C++实现不仅完整，而且不需优化就和PPU性能相似。
至于Atomic操作， 这个是CUDA的至关重要的东西。GPU线程之间的全局同步，唯一同时是性能最高的办法就是原子操作。线程同步只能在block内进行。目前还似乎没有一个衡量该操作性能的办法? GPU在乎的是并行计算, 关键在于把问题并行化. 只要这点做好了, 程序自然会快. 就单个进程而言, SM的IPC性能估计还不如Pentium Pro. But, who cares? 如果要追求特别高效, 那就不用C++特性嘛. 这个在CPU上也是一样的道理. 所以那些做embedded real-time systems 的还是大多用C而不是C++. 如果要在开发的进度和代码效率之间找到一个平衡点的话, C++是个很好的选择. 当然也有很多人码得巨烂, 这个是普通现象, 不是GPU特有的.
总之您说这些, IMHO放到CPU上面也适用.
P大的说法我赞同 , 关于指针语义上确实是统一了. 但是即便最终实现了C++的语法, 也大概还是486的级别的性能. 最终, 程序员还是要非常清楚里面究竟是什么. 因此, C++这种东西, 可能是减少了一开始的门槛, 但是最终要获 ...
RacingPHT 发表于 2009-10-3 00:41 http://we.pcinlife.com/images/common/back.gif "IMHO放到CPU上面也适用."
如果您说的CPU是x86的话, 我认为是不适用的. x86做的这么多事情, 就是为了用来吃烂代码的. 同样烂的代码放到Power, CELL之类的架构上, 一般都会变得非常糟糕, 到CUDA上则是只能说惨不忍睹. 为了一个庞大的市场, 成就了x86这个最智能的CPU体系.
现在CUDA追求C++, 我只能说NV的野心远远不在高性能运算领域, 而是寄希望家用PC的庞大市场. 可惜这个连stack都几乎没办法支持的硬件结构目前是不适合运行这种通用语言的. 最终的演化如果可以看到CUDA"指令集"的多核"CPU"我也不必惊奇. 当然, 线程比现在要少得多, 也要高效得多. 有可能是类似UltraSparc的东西.
很高兴看到NV终于发现线程数量和效率之间的严重不均衡了. 这次CUDA减少线程同时提高效率个人认为是一个好的开始. 片面追求几万个线程, 也许fans会觉得很高兴, 但是却没有任何实际意义. 目前还不是很清楚 Fermi 的双精度实现成本，但是从 Cell 这边看，Cell eDP 和 Cell 相比，eDP 的成本主要在 SPE 上增加了 10%，而 Cell 的 DP 实现是类似于 GT200 也就是有还专门的 DP 运算单元实现。Fermi 的 eDP 是单精度/双精度 单元一起跑实现的，所以在 eDP 的实现上，我想 Fermi 应该不会高到哪里去(+10% per 8 SP？）吧。
当然，这样的 10% 有 n 个存在的话，那就是很可观的成本增加了。 7# voodoo12345
natively support c++ = cuda compiler can compile c++ program. In gt200, it can only compile c program. That is the only difference. Optimization is not there initially. But as an open source community, anyone can contribute to an optimized extension. So one day, there might be a huge open library that provides optimized api. Currently, video transcoding, adobe ps4, flash, matlab, etc, are such examples. The selling point of gpu in the future would like this, if one bot a SOC (junior cpu + senior gpu), one could speed up one of his existing/developing C++ programs by calling a function in an c++ cuda extension. If that gives you even 2x speedup, that sounds valuable. Let alone, future operation systems running faster in SOC than intel only. GPU在乎的是并行计算, 关键在于把问题并行化. 只要这点做好了, 程序自然会快. 就单个进程而言, SM的IPC性能估计还不如Pentium Pro. But, who cares? 如果要追求特别高效, 那就不用C++特性嘛. 这个在CPU上也是一样的道 ...
lik 发表于 2009-10-3 12:43 http://we.pcinlife.com/images/common/back.gif
IPC性能估计还不如Pentium Pro. But, who cares?
其实这句话是很有问题的.
即便有无穷多个运算单元可以将并行的时间减少为0, 一个程序总有一些部分是无法并行的. 这部分决定了一个程序的最高速度.
在CUDA上, 就是和CPU通信的时间, 串行部分运行的时间, 以及初始化的时间等等. 串行部分运行的时间便取决于CUDA的单线程性能.
所以NVIDIA最喜欢演示的demo, 是那些上述三部分时间几乎可以忽略不计的程序. 很遗憾在现实世界中这部分时间是不可忽略的.
Atomic运算子的性能改进, 便可以减少多个SP争用一个地址的情况下的串行执行成本.
Atomic运算子的性能改进, 便可以减少多个SP争用一个地址的情况下的串行执行成本.RacingPHT 发表于 2009-10-4 21:51 http://we.pcinlife.com/images/common/back.gif
Can you elaborate it? Here is my understanding about nv atom, say, threads A, B, C are accessing address A0, the order of excution for the 3 threads is random. If B is accessing it first, then A, C are halted until it is done. So the best total execution time for the 3 threads are A exec time + B exec time + C exec time. Is this close to what you meant here? hd4770:
我不知道NV的具体实现, 有没有一些其他的优化, 例如atomic操作会不会导致线程切换, 因此这个时间可以掩盖掉.
不过在大量访问的情形下, 是这个意思. hd4770:
我不知道NV的具体实现, 有没有一些其他的优化, 例如atomic操作会不会导致线程切换, 因此这个时间可以掩盖掉.
不过在大量访问的情形下, 是这个意思.
RacingPHT 发表于 2009-10-5 09:42 http://we.pcinlife.com/images/common/back.gif
Thanks for clarifying it. From the published document, we can see threads are dymanically loading into SPs. So Yes, if tons of threads are active there, the halt of a few threads due to atomic operation confliction can be easily hidden by other threads. I think the improtance of atimics is to theoretically allow a big task to be broken down into small pieces. Although it sounds odd for multiple pieces working on the same address, but it cannot be ruled out. Today's operation systems tend to break tasks into thousands of small pieces, which could be speeded up by GPUs. Anyhow, your observation is very interesting.