NVIDIA Fermi GF100 及 GF1XX 架构讨论

Edison

目前对 Fermi 图形架构测试最深入的报道，强烈推荐：

http://www.beyond3d.com/content/reviews/55/1

Edison

Fermi 架构的初步概况已经发布，大家可以就相关的话题在本主题中展开讨论（原来的讨论串作为猜测性质，现在已经关闭，你可以在这个连接里回顾：http://we.pcinlife.com/thread-969042-1-1.html）。



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

voodoo12345

Predication（论断）是什么？分支预测？

Edison

http://www.nvidia.com/content/PD ... ingArchitecture.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.


Edison

Predication（论断）是什么？分支预测？
voodoo12345 发表于 2009-10-2 16:30

Prediction（预测） 和 Predication（论断） 是两回事情，虽然有点相似。

http://en.wikipedia.org/wiki/Branch_predication

http://en.wikipedia.org/wiki/Branch_prediction

branch predication 的话，你是 branch 的两边都跑，当确定一个方向后（用 predication  算出一个"筛子"），就只拿这个方向的结果，另一方向的就抛弃掉。

branch prediction 就是猜测哪个方向会发生状况，然后就卯足劲跑这个方向，猜中了就中奖，猜不中就重新跑。

iphone

从这个Fermi100的推出可以看出NV今后的方向，他们走的这一步算是很激进的。如果保守的话他们是可以通过重点发展消费图形处理慢慢蚕食的。不过基于目前的状况来看这一步激进路线前途很坎坷。

Edison

这里是架构讨论区，着重讨论体系架构以及相关的技术讨论，所以如果是发表市场策略看法，请到下面的显卡区发表。

voodoo12345

问个关于C++的问题，这个GPU对C++的完全支持，是应用在开发者利用GPU编写C++代码提高开发效率呢？还是应用在用户执行C++程序时GPU提高程序运行效率呢？

Edison

问个关于C++的问题，这个GPU对C++的完全支持，是应用在开发者利用GPU编写C++代码提高开发效率呢？还是应用在用户执行C++程序时GPU提高程序运行效率呢？
voodoo12345 发表于 2009-10-2 19:23

应该是需要 CUDA 的编译器编译为 PTX 2.0 代码，然后驱动编译为本机代码来执行。

例如，以前的 GPU 是没有 function 的，函式调用都是 inline 的，不是真正的函式。

Fermi 实现了这些特性的支持。

太虚公

有中文资料没？

Edison

能否达到和晶体管数量相称的游戏性能增长？或者说，游戏性能增长相对晶体管性能增长，打了几折？
westlee 发表于 2009-10-2 20:45

目前还没有 gaming 方面更详细的资料看到，不过我想得看具体的游戏吧，当前的许多游戏都不能从 cache 获益，但是并不意味着那些采用了 gpu computing 的游戏不能因此而获益。

Prescott

真正的亮点就是
1. 大幅增加DP单元
2. 统一地址空间
3. ECC
可以说完全是针对GPGPU的。

对于3D图形渲染性能的改进就是double SM的SIMD宽度到16，这个带来的渲染性能提高还有待检验。肯定是不到一倍的，我估计乐观一点在60-80%之间。

Prescott

RacingPHT，我同意你关于性能的意见。这些不会有性能方面的提高，但是对于提高通用编程能力是至关重要的。
关于统一寻址：在实现统一寻址之前，很多C/C++语法根本无法编译成PTX，比如对一个float*的defererence，在编译的时候，编译器很难知道这个指针到底指向什么东西，它这次有可能指向share memory，下次就有可能指向global memory，这个时候，指令生成就成了一个大问题，因为这是完全不同的指针。除非对指针加上修饰，比如 __share__ float*，然后限制不同指针互相赋值也不能相互转换，这显然限制更大，func(__share__ float*)和func(__global__ float*)变成了不同的函数。 所以，统一寻址对于提高可编程能力至关重要。至于原先不能写的东西，这个和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针对的高性能计算市场，应该不是特别重要。

Prescott

另外， 统一地址空间很容易诱惑程序员写出巨慢的代码 。C++也是。

这样我想起了PS3的统一地址空间，程序员不小心用CPU读写到显存出现性能巨幅下降，搞不清楚状况骂娘的情形。
RacingPHT 发表于 2009-10-2 23:00

这个，所有硬件都有这样的问题，或多或少，CPU一个Core比GPU Core大那么多，不都是为了避免这些情况嘛。
这方面Nehalme >> LRB > Fermi

lik

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