基于TILE的渲染

缩写

缩写 说明
GPU Graphics Processing Unit，图形处理单元。
SC Shader Core，渲染核。GPU中的微处理器，负责执行顶点着色程序、片段着色程序和通用计算程序。
TBR Tile-based Rendering，基于TILE的渲染。
IMR Immediate Mode Rendering，立即渲染模式。

TILE

TILE的必应词典翻译是：地砖、瓦片。在这里，TILE是指将整个屏幕分成若干个小格。

TBR是将渲染一屏图像变换为渲染屏幕中的所有TILE。GPU中的一个SC渲染一个TILE，所以GPU中的SC越多，并行渲染TILE的能力就越强。

我们知道，屏幕显示的内容是帧缓冲区中的所有像素颜色，实际上，除了颜色之外，还有像素的深度和模板也会存储在显存中。

将整个屏幕划分成若干个TILE意味着，将颜色、深度、模板缓冲区划分成若干个小的存储。

TILE的大小

一般来说，TILE的大小为32x32或者16x16。有的情况下，TILE还可以是矩形。

总之，TILE的大小一般都不大。

TILE足够小则意味着这个TILE对应的颜色、深度、模板缓冲区的存储在片上也可以放一份。这样做有什么好处呢？

TILE的好处

假设GPU当前渲染的帧中包含100个三角形，且开启了深度模板测试和混合，那么，传统的渲染模式（立即渲染模式，IMR）每渲染一个三角形就需要从显存中读取这个三角形所覆盖的所有像素的颜色、深度和模板，渲染之后还需要立马写回显存。

耗时耗力！

而TBR会逐个遍历TILE，首先看都有哪些三角形覆盖了当前遍历到的TILE，最坏的情况下，某个TILE被100个三角形所覆盖，那么遍历这100个三角形，每渲染一个三角形(TILE覆盖部分而非三角形全部)则将结果放到这个TILE对应的片上高速缓存中，直到遍历完所有的三角形，最后将结果写回显存中的帧缓冲区。

可以看到，TBR将读写显存的频率大大降低。这首先带来的好处是低功耗，其次是高效率。

除此之外，TILE还可以提高多采样的效率，方便应用一些优化策略。

TBR

TBR的处理逻辑是这样的：

1
2
3
4
5
6
7
8
9
10
11
for (auto tile : tiles) {
    auto primitives = tile.getPrimitiveList();

    for (auto primitive : primitives) {
        auto fragments = primitive.getFragmentsInTile(tile);

        for (auto fragment : fragments) {
            fragment.render();
        }
    }
}

而IMR的处理逻辑是这样的：

1
2
3
4
5
6
7
8
9
10
11
for (auto drawcall : drawcalls) {
    auto primitives = drawcall.getPrimitives();

    for (auto primitive : primitives) {
        auto fragments = primitive.getFragments();

        for (auto fragment : fragments) {
            fragment.render();
        }
    }
}

可以看到TBR处理逻辑的第二行有一个操作是：tile.getPrimitiveList，这个操作是为了获取当前帧中覆盖了这个TILE的所有图元。

假设当前帧中有100个三角形，当第一个三角形所覆盖的TILE全部统计好了，也不能立即渲染，而需要等这100个三角形的覆盖情况都统计完成才能开始遍历TILE。这是因为如果直接开始遍历就和IMR的逻辑是一样的了，失去了TBR的优势。遍历TILE的前提便是每个TILE要能够完整表示屏幕中的一小部分，如果当前TILE中的内容还没准备好就直接开始渲染，其渲染结果虽然也可以放在片上，但是你不知道何时将片上存储的内容写回显存。

使用TBR渲染模式的GPU，通常会将染色好的顶点写回显存，然后再读出来将其构建成图元，根据图元的覆盖情况构建每个TILE的图元列表，图元列表也会被写回显存。一帧的所有TILE的图元列表构建完成后，再让SC读取出为其分配的TILE的图元列表。而采用IMR渲染模式的GPU是没有读写图元列表操作的。

虽然TBR多了读写图元列表的操作，但是相比于节省下来的读写缓冲区的操作，还是很划算的。

TBDR

由于要绘制的场景中存在大量的遮挡情况，而被遮挡的片段也是花费了很多时间和资源才计算出它们的位置和颜色，但最终我们看不到（最终的图像不会包含这些像素点的信息），所以，这部分工作就属于无用功。

TBDR（Tile Based Deferred Rendering）的核心思想是在确定好每个像素是否会最终显示到屏幕上之后，才开始渲染那些会显示到屏幕上的像素。

参考

GLSL 编译器处理逻辑

GLSL介绍

GLSL是多种紧密相关的语言，这些语言用于为OpenGL API处理管道中包含的每个可编程处理器创建着色器。包含如下的着色器：

  • Vertex
  • Tessellation Control
  • Tessellation Evalution
  • Geometry
  • Fragment
  • Compute

编译器处理阶段

Shader处理器的编译单元在编译的最后阶段可选地链接在一起之前被分别处理。其处理逻辑如下：

  1. 源字符串是以字节序列输入，’\0’被解释为终止符。
  2. 所有源字符串被串联以形成单个输入，’\0’字节被丢弃，所有其他值均保留。(每个shader可能有多个输入字符串)
  3. 每个字符串均根据UTF-8标准进行解释，不同之处在于，所有无效字节序列均以其原始形式保留以用于后续处理。
  4. 每个{回车，换行}和{换行，回车}序列都由单个换行符代替，所有剩余的回车符和换行符都将用单个换行符替代。
  5. 每个字符的行号等于前一个换行符的数量加一。请注意，此操作只能随后通过#line指令进行更改，并且不受编译阶段6中删除换行符的影响。
  6. 在换行符之前出现反斜杠（’'）的都将被删除（转义）。请注意，不会替换任何空格，从而允许单个预处理令牌跨越换行符。此操作不是递归的；不会删除任何新生成的{反斜杠 换行符}序列。
  7. 所有注释均替换为一个空格。注释中允许所有（非零）字符和无效的UTF-8字节序列。’//‘样式注释包括初始的’//‘标记，并一直延续到但不包括终止换行符。’/…/‘注释同时包含开始和结束标记。
  8. 源字符串将转换为一系列预处理Token。这些Token包括预处理编号，标识符和预处理操作。每一个Token的行号是Token开始的第一个字节的行号。
  9. 预处理器运行。执行指令并执行宏扩展。
  10. 空格和换行符将被丢弃。
  11. 预处理Token将转换为Token
  12. 语法根据GLSL ES语法进行分析。
  13. 根据语言的语义规则检查结果。
  14. 将着色器链接在一起以形成一个或多个程序或分离的程序。当将一对连续阶段的着色器链接到同一程序时，两个着色器中未使用的任何输出和相应的输入都可能会被丢弃。
  15. 生成二进制文件。

From：https://www.khronos.org/registry/OpenGL/specs/es/3.2/GLSL_ES_Specification_3.20.html#logical-phases-of-compilation

Fermi架构介绍

介绍

Fermi是NVIDIA公司的GPU架构的名称。取这个名字是为了纪念美籍著名意大利物理学家Erico Fermi。

Erico Fermi造出了人类第一台可控核反应堆—芝加哥一号堆，被称为“原子能之父”，在1938年获得了诺贝尔物理学奖。

NVIDIA声称Fermi架构是世界上第一个完整地支持通用计算的GPU架构。

架构模型

编程模型

Application

Application可以包含一个或多个Kernel，通常用来做顶点/像素染色或者通用计算。与常见的CPU程序相比，Application的分支跳转语句相对较少；程序长度相对较短；算数运算类语句相对较多。

Application支持OpenGL、CUDA、OpenCL和Direct Compute API。

Kernel

一个Grid执行一个Kernel。这是CUDA和OpenCL中的概念，意指计算算法单元。需要注意的是它和操作系统中的内核不是同一个概念。

Kernel在C语言的基础上扩展了并行运算的语法，以取代串行的循环等语句。除此之外，它还支持C++、Fortran、Java、MATLAB和Python。

Grid

Grid可以包含一个或多个ThreadBlock。与Kernel是一一对应的关系。

ThreadBlock

ThreadBlock可以包含最多48个Warp，也就是1536个线程。一个ThreadBlock里面的所有线程都运行在同一个SM上，他们之间可以协作和共享存储。

Warp

一个Warp包含32个Thread，不同ThreadBlock中的Warp可以并行执行。Warp是SM中的基本调度单元。

Thread

多个Thread可以并行执行，Fermi采用SIMT（Single Instruction Multiple Threads）架构。Warp可以快速切换，这是因为每个线程都有自己独立的寄存器和私有存储，Warp切换时不需要数据搬运。

GPGPU-Sim 安装过程

简介

GPU

GPU英文名称为Graphic Processing Unit，中文名称为图形处理器，主要用于计算机系统中的显示及图形处理，又称为Video Card(显卡)。

GPU具有一下特性：

  • 针对高度并行的工作负载进行优化
  • 高度可编程性
  • 桌面级的超级计算机

GPGPU

异构计算（Heterogeneous Computing）是指在异构计算系统上进行的并行计算。
GPGPU（General-purpose computing on graphics processing units）是一种利用处理图形任务的图形处理器来计算原本由中央处理器处理的通用计算任务。

GPGPU-Sim

GPGPU-Sim 是一个时钟级别的GPU仿真模型，可以运行使用cuda或者OpenCL编写的GPU计算程序。GPGPU的github如下：
https://github.com/gpgpu-sim/gpgpu-sim_distribution

安装 GPGPU Sim

本文介绍在虚拟机Centos7上配置安装GPGPU sim环境。
需要准备环境：

安装Cuda

GPGPU-Sim支持的Cuda版本有：4.2, 5.0, 5.5, 6.0, 7.5, 8.0, 9.0, 9.1。

1
2
yum localinstall cuda-repo-rhel7-7.5-18.x86_64.rpm
yum install cuda-tooklit-7-5

编译安装GPGPU-Sim

安装依赖库

安装依赖:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
# GPGPU-Sim dependencies:
yum install gcc
yum install gcc-c++
yum install make
yum install makedepend
yum install xorg-x11-utils
yum install bison
yum install flex
yum install zlib

# GPGPU-Sim documentation dependencies:
yum install doxygen
yum install graphviz

# AerialVision dependencies:
yum install python-pmw
yum install python-ply
yum install python2-numpy
yum install libpng12-devel
yum install python-matplotlib

编译

编译前需要设置环境变量

1
2
3
export CUDA_INSTALL_PATH=/usr/local/cuda
export PATH=${CUDA_INSTALL_PATH}/bin:${PATH}
source setup_environment

之后可以直接使用make进行编译：

1
make -j8

运行demo

cuda的helloworld程序如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
/* file: hello.cu */
#include <stdio.h>
#include <cuda_runtime.h>

__global__ void add(int a, int b, int *c)
{
    *c = a + b;
}

int main()
{
    int c;
    int *dev_c;
    cudaMalloc((void **)&dev_c, sizeof(int));
    add<<<1, 1>>>(2, 7, dev_c);
    cudaMemcpy(&c, &dev_c, sizeof(int), cudaMemcpyDeviceToHost);
    printf("2 + 7 = %d\n", c);

    return 0;
}

编译程序时指定--cudart shared确保可执行程序动态链接到CUDA runtime库。

1
nvcc --cudart shared -o hello hello.cu

上文2.2.2中通过source setup_environment设置可执行程序连接到GPGPU-Sim编译的libcudart.so中。
运行ldd确保链接正确的libcudart.so:

1
ldd hello

运行程序之前需要拷贝GPGPU-Sim路径下的configs/tested-cfgs/SM2_GTX480/的配置文件到当前运行demo的路径下。
最后运行编译的可执行程序即可正常仿真:

1
./hello