执行

在C++程序的执行阶段，程序将按照main()函数开始执行，然后根据程序中的逻辑顺序执行相应的语句和函数调用。

内存管理

容易发生内存问题的包括栈溢出和内存泄漏，分别对应内存模型中的栈和堆，但需要注意以下问题

• 不要返回局部变量的引用或指针
• 不要对已经释放的内存进行操作

异常处理

线程安全

资源管理

代码优化

边界条件

错误处理

测试与调试

动态联编和静态联编

• 指计算机程序不同部分彼此关联，分为静态和动态
• 静态联编是指联编⼯作在编译阶段完成的，这种联编过程是在程序运⾏之前完成的，⼜称为早期联编。要实现静态联编，在编译阶段就必须确定程序中的操作调⽤（如函数调⽤）与执⾏该操作代码间的关系，确定这种关系称为束定，在编译时的束定称为静态束定。静态联编对函数的选择是基于指向对象的指针或者引⽤的类型。其优点是效率⾼，但灵活性差
• 动态联编指联编在程序运⾏时动态地进⾏，根据当时的情况来确定调⽤哪个同名函数，实际上是在运⾏时虚函数的实现。这种联编⼜称为晚期联编，或动态束定。动态联编对成员函数的选择是基于对象的类型，针对不同的对象类型将做出不同的编译结果
• C++中⼀般情况下的联编是静态联编，但是当涉及到多态性和虚函数时应该使⽤动态联编。动态联编的优点是灵活性强，但效率低。动态联编规定，只能通过指向基类的指针或基类对象的引⽤来调⽤虚函数，其格式为：指向基类的指针变量名->虚函数名（实参表）或基类对象的引⽤名.虚函数名（实参表）

实现的条件

• 必须把动态联编的⾏为定义为类的虚函数；
• 类之间应满⾜⼦类型关系，通常表现为⼀个类从另⼀个类公有派⽣⽽来；
• 必须先使⽤基类指针指向⼦类型的对象，然后直接或间接使⽤基类指针调⽤虚函数；

何时需要初始化

• 引用成员变量
• const
• 基类构造有一组参数
• 成员构造有一组参数
• 初始化顺序：按照成员声明顺序决定

C++初始化

类型

• 编译初始化
• 动态初始化

编译初始化

• 静态初始化在程序加载的过程中完成
• 包括全局变量初始化和constexpr类型的初始化
  • zero initialization 的变量会被保存在 bss 段
  • constexpr initialization 的变量则放在 data 段内
  • 其次全局类对象也是在编译器初始化。

动态初始化

出现时机：出现在编译期和运行期的局部位置初始化

• 动态初始化也叫运行时初始化
• 需要经过函数调用才能完成的初始化、类初始化
  • 局部静态类对象的初始化
  • 局部静态变量的初始化
• 动态初始化一般出现在

动态初始化中静态局部变量2个问题

线程安全问题

实现方法

• 一个线程在初始化 m 的时候，其他线程执行到 m 的初始化这一行的时候，就会挂起而不是跳过
  • 局部静态变量在编译时，编译器的实现是和全局变量类似的，均存储在bss段中。
  • 然后编译器会生成一个保证线程安全和一次性初始化的整型变量，是编译器生成的，存储在 bss 段。
    • 它的最低的一个字节被用作相应静态变量是否已被初始化的标志
      • 若为 0 表示还未被初始化，否则表示已被初始化(if ((guard_for_bar & 0xff) == 0)判断)。
    • __cxa_guard_acquire 实际上是一个加锁的过程，
      • 相应的 __cxa_guard_abort 和__cxa_guard_release 释放锁。

内存泄漏问题

原因

• 在局部作用域消失时，data区仍然保存其内存空间
• 执行路径不明
  • 对于局部静态变量，构造和析构都取决于程序的执行顺序。程序的实际执行路径不可预知的
• 关系不明
  • 局部静态变量分布在程序代码各处，彼此直接没有明显的关联，很容易让开发者忽略它们之间的这种关系

建议

• 减少使用局部静态变量

实例化

过程

• 分配内存空间
• 执行构造

不可实例化的类

• 带有一个或以上virtual的函数的类
• 工具类（直接通过static调用函数）

如何阻止实例化

• 包含纯虚函数
• 构造函数私有

实例化中的变量初始化时机与顺序

时机

• 类中const初始化必须在构造函数初始化列表中初始化
• 类中static初始化必须在类外初始化
• 成员变量初始化顺序按照类中声明顺序，而构造函数初始化顺序按照成员变量在构造函数中位置决定

顺序

• 初始化base类中的static部分（按程序出现顺序初始化）
• 初始化派生类中的static部分（按程序出现顺序初始化）
• 初始化base类的普通成员变量和代码块，再执行父类的构造方法；
• 初始化派生的普通成员变量和代码块，在执行子类的构造方法；

问题

把异常完全封装在析构函数内部，决不让异常抛出函数之外

栈上的指针什么时候析构

.h .cpp .hpp 关系

函数指针

• 指向函数的指针变量。函数指针本身是一个指针变量，变量指向一个函数。
  • 有了这个指针，可以用这个指针变量调用函数
  • 除了调用外还可以做函数的参数

char * fun(char * p) {…}  // 指向char的指针
char * (*pf)(char * p);   // pf函数指针
pf = fun;                 // 函数指针指向函数
pf(p);                    // 调用

bool、int、float、指针类型变量a与0的比较语句

if(!a) or if(a)

if (a ==0)

if(a <= 0.000001 && a >=-0.000001)

if(a != NULL ) or if(a == NULL)

helloworld程序开始到打印到屏幕的过程

• ⽤户告诉操作系统执⾏ HelloWorld 程序（通过键盘输⼊等）；
• 操作系统：找到 HelloWorld 程序的相关信息，检查其类型是否是可执⾏⽂件；并通过程序⾸部信息，确定代码和数据在可执⾏⽂件中的位置并计算出对应的磁盘块地址；
• 操作系统：创建⼀个新进程，将 HelloWorld 可执⾏⽂件映射到该进程结构，表示由该进程执⾏ HelloWorld程序；
• 操作系统：为 HelloWorld 程序设置 cpu 上下⽂环境，并跳到程序开始处；
• 执⾏ HelloWorld 程序的第⼀条指令，发⽣缺⻚异常；
• 操作系统：分配⼀⻚物理内存，并将代码从磁盘读⼊内存，然后继续执⾏ HelloWorld 程序;
• HelloWorld 程序执⾏ puts 函数（系统调⽤），在显示器上写⼀字符串;
• 操作系统：找到要将字符串送往的显示设备，通常设备是由⼀个进程控制的，所以，操作
• 系统将要写的字符串送给该进程;
• 操作系统：控制设备的进程告诉设备的窗⼝系统，它要显示该字符串，窗⼝系统确定这是⼀个合法的操作，然后将字符串转换成像素，将像素写⼊设备的存储映像区;
• 视频硬件将像素转换成显示器可接收和⼀组控制数据信号;
• 显示器解释信号，激发液晶屏;

printf实现原理

函数参数通过压入堆栈的方式来传递参数 而栈是从内存高地址向低地址生长，因此最后压栈的在堆栈指针的上方，printf第一个被找到的参数就是字符指针。函数通过判断字符串控制参数的个数来判断参数个数与数据类型，进而算出需要的堆栈指针偏移量

C++代码的编译流程

• 步骤：预处理器→编译器→汇编器→链接器
• 预处理阶段：在这个阶段，编译器会处理预处理指令（以 ‘#’ 开头），如 #include、#define 等。预处理器会将这些指令替换为相应的内容，生成一个被称为 “translation unit” 的中间文件。
• 编译阶段：在这个阶段，预处理后的代码会被编译成汇编语言。编译器会将源代码转换成汇编语言的表示形式，这个过程包括了语法和语义分析，以及生成中间表示（如抽象语法树）编译器可能会进行一系列的优化操作，以提高程序的性能和效率。这些优化包括但不限于循环展开、内联函数、常量折叠等。该阶段会构建继承关系。编译工具包括GCC、Clang、MSVC等
• 汇编阶段：汇编器将汇编代码转换为机器代码，也称为目标代码或对象代码。 目标代码是特定于体系结构的，这意味着为x86架构生成的代码不能直接在ARM架构上运行。
• 链接阶段：如果程序由多个源文件组成，那么编译后会生成多个目标文件（Object Files）。链接器将这些目标文件以及所需的库文件链接在一起，生成最终的可执行文件。这个过程包括解析符号引用、地址重定向和符号重定向等步骤。链接器可以生成两种类型的输出：可执行文件（可以直接运行）和库文件（包含可以在其他程序中使用的代码和数据）。

静态链接和动态链接

静态链接

.a 定义

• 静态链接是将程序的所有模块在编译时链接到一个单独的可执行文件中的过程。
• 在静态链接中，目标文件中的所有模块和库都会被复制到最终的可执行文件中，使得可执行文件独立于系统环境。

过程

• 在编译和链接程序时，链接器将目标文件中的所有符号（如函数和变量）解析并合并到一个单独的可执行文件中。
• 静态链接器将所有依赖项（如库文件）的代码和数据都复制到可执行文件中，因此生成的可执行文件比较大。

优点

• 独立性：生成的可执行文件可以在没有外部依赖的情况下在任何环境中运行。
• 性能：静态链接可以避免运行时的库加载和解析过程，因此可以提高程序的启动速度。

缺点

• 可执行文件较大：每个可执行文件都包含了所有依赖项的代码和数据，因此可能会占用较多的磁盘空间。
• 更新困难：如果库更新了，所有依赖该库的程序都需要重新编译和链接以使用新版本。

动态链接

.so，动态链接库的地址：LD_LIBRARY_PATH，GCC默认为动态链接 定义

• 动态链接是在程序运行时将程序的模块和所需的库链接到一起的过程。
• 在动态链接中，程序运行时所需的库被加载到内存中，而不是在编译时被合并到可执行文件中。

过程

• 在编译和链接程序时，只生成包含程序自身代码和对外部库的引用的可执行文件。
• 当程序运行时，操作系统的动态链接器（如Linux中的ld.so）将程序所需的库加载到内存中，并将程序中的引用解析为动态库中的符号。

优点

• 节省空间：由于可执行文件只包含程序自身的代码和数据，因此通常较小。
• 管理便捷：多个程序可以共享同一个动态库，节省了系统资源并方便了库的更新和维护。

缺点

• 运行时开销：程序运行时需要加载和解析动态库，可能会稍微降低启动速度。
• 系统依赖：程序运行需要确保所依赖的动态库在系统中可用，否则可能会导致运行时错误。

加载

• 在加载阶段，操作系统负责将可执行文件的内容加载到内存中，并为程序分配必要的资源。这个过程包括以下几个关键步骤：
  • 内存分配：操作系统会为程序分配内存空间，这个空间通常包括代码段、数据段、堆和栈等区域。代码段用于存储程序的指令，数据段用于存储静态变量和全局变量，堆用于动态分配内存，栈用于存储函数调用和局部变量。
  • 装载程序：操作系统会将可执行文件的内容从存储设备（如硬盘）中读取到内存中。这包括将程序的指令、全局变量的初始值等加载到相应的内存区域。
  • 地址重定位：在加载阶段，操作系统还会执行地址重定位的操作。由于程序可能会被加载到内存的不同位置，其中的地址引用需要被调整，确保程序能够正确地访问内存中的数据和指令。
  • 动态链接库加载：如果程序依赖于动态链接库（DLL），那么这些库也会在加载阶段被加载到内存中，并与程序进行链接。
• 一旦加载完成，操作系统会将程序的控制权交给程序的入口点，即 main() 函数

加载&运行时的对象构建与析构顺序

顺序原因

• 在编译阶段会产生继承树

构建顺序

• 基类对象构建
• 基类对象成员按照声明顺序构建，不按照初始化列表
• 对象构建

析构顺序

• 析构顺序与构建顺序相反

#include <iostream>

class Member1 {
public:
    Member1() {
        std::cout << "Member1 constructor\n";
    }
    ~Member1() {
        std::cout << "Member1 destructor\n";
    }
};

class Member2 {
public:
    Member2() {
        std::cout << "Member2 constructor\n";
    }
    ~Member2() {
        std::cout << "Member2 destructor\n";
    }
};

class Base1 {
public:
    Base1() : member111(), member222() {
        std::cout << "base1 constructor\n";
    }
    ~Base1() {
        std::cout << "base1 destructor\n";
    }
private:
    Member1 member111;
    Member2 member222;
};

class Base: public Base1 {
public:
    Base() {
        std::cout << "Base constructor\n";
    }
    ~Base() {
        std::cout << "Base destructor\n";
    }
};

class Derived : public Base {
public:
    Derived() : member1(), member2() {
        std::cout << "Derived constructor\n";
    }
    ~Derived() {
        std::cout << "Derived destructor\n";
    }
private:
    Member2 member2;
    Member1 member1;

};
Derived d;
int main() {
    std::cout << "main\n";
    return 0;
}

// Member1 constructor
// Member2 constructor
// base1 constructor
// Base constructor
// Member2 constructor
// Member1 constructor
// Derived constructor
// main
// Derived destructor
// Member1 destructor
// Member2 destructor
// Base destructor
// base1 destructor
// Member2 destructor
// Member1 destructor

C++加载的内存分配模型

代码段

• 可执行文件中的代码被复制到代码段中
  • 只读，防止程序意外修改自身的代码
  • 共享，多个进程同时执行相同的程序，可以共享同一个代码段，节省内存

数据段

• 包含程序中的全局变量和静态变量，以及一些常量数据。
• 在加载过程中，数据段被初始化并分配内存空间，以存储这些变量的值。
• 数据段通常分为两部分：初始化的数据段（Initialized Data Segment）和未初始化的数据段（Uninitialized Data Segment，也称为BSS段）。
  • 初始化的数据段存储已经初始化的全局变量和静态变量的值
  • 未初始化的数据段则存储全局变量和静态变量的声明，但尚未被初始化的值。程序开始执行时将其初始化为零或空值。

堆

• 动态内存分配的区域，用于存放程序运行时动态分配的内存。
• 在加载过程中，操作系统会为堆分配一块初始大小的内存空间，通常称为堆的起始地址。
• 程序可以通过调用new和delete等动态内存管理函数来在堆上分配和释放内存，堆的大小可以根据程序的需要动态增长或缩小。

栈

• 栈是用于函数调用和局部变量存储的内存区域。
• 在加载过程中，操作系统会为每个线程分配一块栈空间，用于存储函数调用时的参数、局部变量和函数调用的返回地址等信息。栈是一种先进后出（LIFO）的数据结构，函数调用时会将调用函数的参数和局部变量压入栈中，函数返回时会从栈中弹出这些数据。

加载中的程序装载

• 在C++中，加载时初始化主要涉及全局变量（数组或对象）和静态变量（数组或对象）。这些变量在程序加载到内存时就会被初始化。
  • 全局变量（数组或对象）在程序的任何部分都可见，它们的初始化在程序启动时完成，在main函数执行之前进行。

// 变量
int globalVar = 42; // 加载时初始化

// 数组
int arr[4] = {1,2,3,4}; // 加载时初始化

// 对象
class MyClass {
public:
    MyClass() {
        std::cout << "MyClass Constructor" << std::endl;
    }
};

MyClass globalObject; // 加载时初始化，构造函数在main之前调用

• 静态变量（数组或对象）
  • 全局静态变量、命名空间作用域中的静态变量和类的静态成员变量：通常在加载时初始化。这些变量的作用域仅限于声明它们的文件。
  • 静态局部变量：
    • 如果初始化表达式是常量表达式（带const关键字），则在加载时初始化。
    • 否则，在第一次使用时初始化。
  • 静态变量的初始化顺序不确定，特别是针对不同的单元时，会导致使用未定义的静态变量。解决办法是使用函数的本地静态变量，在第一次使用时才进行初始化，同时线程安全

// 全局静态变量
static int staticGlobalVar = 10; // 加载时初始化

// 命名空间下的全局静态变量
namespace MyNamespace {
    static int staticVar = 10; // 命名空间作用域中的静态变量，在加载时初始化
}

// 类的静态成员变量
class MyClass {
public:
    static int staticMemberVar;
};

int MyClass::staticMemberVar = 300; // 加载时初始化

// 局部静态变量
void exampleFunction() {
    static const int staticLocalVar = 20; // 如果是常量表达式，加载时初始化
    static int staticconst = 10; // 加载时初始化
}

void func() {
    static int localStaticVar = someFunction(); // 第一次调用func时初始化
}

// 静态全局对象
class MyClass {
public:
    MyClass() {
        std::cout << "MyClass Constructor" << std::endl;
    }
};

static MyClass staticGlobalObject; // 加载时初始化，构造函数在main之前调用

• 初始化列表或联合体：使用初始化列表初始化的全局或静态变量在加载时初始化，联合体同理

// struct
struct Data {
    int x;
    int y;
};

Data data = {1, 2}; // 加载时初始化

// union
union MyUnion {
    int a;
    float b;
};

MyUnion myUnion = { 1 }; // 加载时初始化

• 外部库与自定义初始化
  • 有些外部库会在加载时初始化全局状态。这些初始化通常通过库的构造函数或初始化函数进行。
  • 使用自定义初始化函数：通过自定义初始化函数，程序员可以确保某些初始化在加载时完成。比如使用GCC的构造函数属性

// 外部库初始化
extern "C" void __attribute__((constructor)) myInitFunction() {
    // 加载时初始化代码
}

// 自定义初始化
void __attribute__((constructor)) initFunction() {
    // 加载时初始化代码
}

• 常量表达式的初始化：
  • 对于局部静态常量来说，在加载时初始化
  • 对于编译时能确定的初始值的常量，会在编译期间进行初始化，是直接插入exe文件中
  • 对于某些需要运算得到的值，初始化在运行时进行

执行

在C++程序的执行阶段，程序将按照main()函数开始执行，然后根据程序中的逻辑顺序执行相应的语句和函数调用。

global memory 内存优化

global memory 内存合并思想

DRAM局部性思想与 global memory 的局部性原理

• GPU对于内存数据的请求是以wrap为单位，而不是以thread为单位。如果数据请求的内存空间连续，请求的内存地址会合并为一个warp memory request，然后这个request由一个或多个memory transaction组成。具体使用几个transaction 取决于request 的个数和transaction 的大小。
  • 该方法从数据存取->数据利用中利用局部性思想，与DRAM硬件实现中在数据存储->数据读取中利用局部性思想同理

GPU与CPU在局部性原理的不同

• cache大小：CPU的片上有很大的cache，相对于CPU，GPU相对较少。
• thread对于cache的处理：
  • 组织形式：CPU中的不同thread访问连续的内存会被缓存（局部性），GPU的thread通过warp进行封装，warp访问内存通过L1/shared memory等进行
  • 数据影响性：CPU的不同thread之间数据相互不影响，GPU的thread之间的数据会存在互相影响的问题
• 充分利用内存方法：
  • CPU：由于thread相对稀缺，为充分利用core性能，每个core负责一段连续的内存。（e.g. thread 1 : array 0-99; thread 2 : array 100-199; thread 3 : array 200-299.）
  • GPU：由于cache相对稀缺，为充分发挥带宽，应当在warp的每个iteration中保证花费全部cache line。

global memory带宽检测

带宽包括理论带宽和实际带宽，实际带宽检测包括CPU方法和GPU方法

实际带宽计算

CPU计算带宽方法

相比起GPU来说，比较粗糙

#include <cuda_runtime.h>
#include <chrono>
#include <iostream>

int main() {
    // 假设我们有一个足够大的数组
    size_t bytes = 1024 * 1024 * 1024; // 1GB
    float* host_ptr = new float[bytes / sizeof(float)];
    float* device_ptr;
    cudaMalloc((void**)&device_ptr, bytes);

    // 填充host_ptr...（如果需要）

    // 开始计时
    auto start = std::chrono::high_resolution_clock::now();
  
    // 同步设备（可选，但在这里确保没有之前的操作干扰）
    cudaDeviceSynchronize();

    // 执行内存传输
    cudaMemcpy(device_ptr, host_ptr, bytes, cudaMemcpyHostToDevice);

    // 同步设备以确保传输完成
    cudaDeviceSynchronize();

    // 结束计时
    auto end = std::chrono::high_resolution_clock::now();

    // 计算时间差
    auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count();

    // 计算带宽
    double bandwidth = (bytes / 1024.0 / 1024.0) / (duration / 1000.0); // MB/s

    std::cout << "Estimated bandwidth: " << bandwidth << " MB/s" << std::endl;

    // 清理内存
    cudaFree(device_ptr);
    delete[] host_ptr;

    return 0;
}

GPU计算带宽方法

GPU时间与操作系统无关，主要通过GPU的event计时完成

#include <cuda_runtime.h>  
#include <stdio.h>  
  
// 假设的kernel函数声明（这里只是示意）  
__global__ void kernel(float* d_odata, float* d_idata, int size_x, int size_y, int NUM_REPS) {  
    // ... kernel的实现 ...  
}  
  
int main() {  
    cudaEvent_t start, stop;  
    float time;  
    cudaError_t err;  
  
    // 创建事件  
    err = cudaEventCreate(&start);  
    checkCudaErrors(err);  
    err = cudaEventCreate(&stop);  
    checkCudaErrors(err);  
  
    // 记录开始事件  
    err = cudaEventRecord(start, 0);  
    checkCudaErrors(err);  
  
    // 假设的GPU内存分配和数据传输（这里只是示意）  
    // float* d_odata, d_idata; // 需要提前分配  
    // cudaMemcpy(...); // 需要进行数据传输  
  
    // 执行kernel  
    int grid = ...; // 定义grid大小  
    int threads = ...; // 定义block大小  
    kernel<<<grid, threads>>>(d_odata, d_idata, size_x, size_y, NUM_REPS);  
  
    // 记录结束事件  
    err = cudaEventRecord(stop, 0);  
    checkCudaErrors(err);  
  
    // 等待结束事件  
    err = cudaEventSynchronize(stop);  
    checkCudaErrors(err);  
  
    // 计算时间（毫秒）  
    float milliseconds;  
    err = cudaEventElapsedTime(&milliseconds, start, stop);  
    checkCudaErrors(err);  
  
    // 假设我们知道数据传输的总量（以字节为单位）  
    size_t totalBytesTransferred = size_x * size_y * sizeof(float) * 2; // 假设输入和输出都是float数组，并且大小相同  
  
    // 计算带宽（字节/秒）  
    double bandwidth = (double)totalBytesTransferred / (milliseconds * 1e-3);  
  
    // 输出结果  
    printf("Execution time: %f ms\n", milliseconds);  
    printf("Estimated bandwidth: %f GB/s\n", bandwidth / (1024 * 1024 * 1024)); // 转换为GB/s  
  
    // 销毁事件  
    err = cudaEventDestroy(start);  
    checkCudaErrors(err);  
    err = cudaEventDestroy(stop);  
    checkCudaErrors(err);  
  
    // ... 其他清理代码 ...  
  
    return 0;  
}  
  
// 辅助函数来检查CUDA错误  
void checkCudaErrors(cudaError_t err) {  
    if (err != cudaSuccess) {  
        fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));  
        exit(EXIT_FAILURE);  
    }  
}

理论带宽计算

以V100为例的理论带宽计算

• V100使用HBM2作为带宽（时钟频率877MHz、为double data rate RAM 、内存带宽为4096 bit。得到的带宽为

$ 0.877 * 10^9 * 4096 / 8 * 2 / 10^9 = 898GB/s $

• 代码得到GPU带宽

cudaDeviceProp dev_prop;
CUDA_CHECK( cudaGetDeviceProperties( &dev_prop, dev_id ) );
printf("global memory bandwidth %f GB/s\n", 2.0 * dev_prop.memoryClockRate * ( dev_prop.memoryBusWidth / 8 ) / 1e6 );

理论带宽的影响因素：GDDR中的ECC导致带宽下降

• ECC是一种用于检测和纠正数据传输或存储中错误的编码技术。通过在数据中添加冗余位（即ECC位），ECC技术能够检测和纠正某些类型的错误，从而提高数据的完整性和准确性。然而，这些额外的ECC位会占用原本用于数据传输的带宽，导致所谓的ECC overhead
• 当在GDDR中加入ECC时，为了提供错误检测和纠正的能力，需要在每个数据传输单元（如一个字节或一组字节）中添加额外的ECC位。这些ECC位并不直接用于传输有效数据，因此会占用原本可以用于数据传输的带宽。这就是ECC overhead的来源，它会导致GDDR的理论带宽下降。
• HBM2是一种高带宽内存技术，其设计初衷就是为了提供极高的数据传输速率。为了支持ECC而不牺牲带宽，HBM2采用了专门的硬件设计，为ECC位分配了独立的存储空间。这意味着ECC位不会占用原本用于数据传输的带宽，因此不会导致理论带宽的下降。这种设计使得HBM2能够在保持高带宽的同时，提供强大的数据错误检测和纠正能力。
• 总结来说，ECC在GDDR中会导致ECC overhead，进而降低理论带宽，是因为ECC位占用了原本用于数据传输的带宽。而HBM2通过专门的硬件设计，为ECC位分配了独立的存储空间，从而避免了ECC对带宽的影响。这种设计使得HBM2在保持高带宽的同时，能够提供强大的数据错误检测和纠正能力。

Visual profiler 内核性能分析工具

requested throughput 和 global throughput

• 系统的global throughput 相当于物理理论带宽（考虑到cache line 要一起传送的带宽）
• 系统的requested memory 相当于实际带宽（并未考虑 cache line）
• 应当让requested throughput尽可能接近global throughput。

GPU中的DRAM硬件技术

本部分主要讲述global memory的硬件实现DRAM

Bit Line & Select Line

原理

• 一个电容决定一个bit， 通过select线选择读取哪个bit line，然后通过bit line选择对应的bit。
• 需要不停检查和电容充放电，因此叫做DRAM

特点

• bit line的容量大，存在电容耦合与信号衰减，同时更容易遇到数据冲突，同时bit line的带宽有限。
• 每个bit电容，需要信号放大器放大，进一步限制传统DRAM设计

Core Array & Burst的DRAM数据传输

改进

• 在bit line -> bus的传输中，添加column buffer将bit line数据进行备份，提高bit line其他数据传输到bus的效率（数据局部性）

传输过程

• 数据传输分为两个部分
  • core array -> column latches / buffer
  • column latches / buffer-> mux pin interface

传输耗时

• core array -> column latches / buffer 耗时久
• buffer -> mux pin interface 的耗时相对较小（原先耗时长）

burst与burst size/ line size

• burst：当访问一个内存位置的时候，select line选择的bit line的数据会全部从core array传输到column latches/ buffer中。使用数据根据mux来确定传输给bus哪些数据，这样可以加速
• burst size/ line size：读取一次memory address，会有多少个数据从core array被放到buffer中

DRAM 设计思想（局部性原理）

• 在DRAM硬件实现中，为了提高DRAM的速度，充分利用局部性思想，通过添加buffer缓存，将core array -> column latches / buffer的时间转换为buffer -> mux pin interface时间，极大地提高速度
  • 该方法从数据存储->数据读取中利用局部性思想

Multiple Banks技术

引入原因

• 从core array到buffer的时间长，单个传输导致实际使用的bus interface数据bandwidth未被完全利用($T_c:T_b = 20:1$)。如果只使用一个bank导致带宽大部分时间为空闲。
• 所以需要在一个bus 上使用多个banks，来充分利用bus bandwidth。如果使用多个bank，大家交替使用interface bus，保证bus不会空闲，保证每个时间都有数据从bus传送过来。

一个bus需要多少个bank？

• 如果访问core array与使用bus传输数据的时间比例是20:1，那么一个bus至少需要21个bank才能充分使用bus bandwidth。
• 一般bus有更多的bank，不仅仅是ratio+1
  • 并发访问：如果同时对多个bank进行操作，这些操作可以并发进行，因为它们访问的是不同的物理存储体。假设一个DRAM模块有两个bank，并且内存控制器同时接收到两个读请求。如果这两个请求都针对同一个bank，那么它们将顺序执行，导致等待时间和带宽的减少。但是，如果这两个请求分别针对两个bank，那么它们可以并发执行，从而提高带宽。
  • 避免bank冲突：当多个请求尝试同时访问同一个bank时，会发生冲突。为尽可能减少冲突可能性，使用更多的banks。

banks与burst的速度对比

• burst的速度快于banks的速度

Multiple Channels技术

• 一般的GPU processor要求带宽达到128GB/s，HBM2要求带宽为898GB/s，在使用multiple banks后仍未满足要求（使用DDR DRAM clock为1GHz，每秒能传送8 bytes/ 1 words，得到的传输速率为16GB/s）因此需要多个channels

数据Interleaved（交织）分布技术

数据交织原因

• 充分利用channel的带宽，实现max bandwidth。
• 允许在core array -> buffer 的时间进行转化成多个channels获取数据（并行），减少总的时间。

如何实现数据交织

将数据平均分布在channels对应的banks中（上图中）

内存结构

内存延迟

内存分类

• linear

  memory

• CUDA arrays

  -> 常用于texture

各种内存的物理实现

• global memory -> DRAM
• shared memory -> SRAM
• const memory -> 专门存储器
• texture memory -> DRAM，专门用于高速图像缓存，具备二维结构
• L1/L2 cache -> 利用SRAM和逻辑电路实现

NULL和nullptr

NULL 在c++与c中，使用NULL代表0或void*，但可能会产生类型不匹配的警告或错误

  #ifdef __cplusplus
  #define NULL 0
  #else
  #define NULL ((void*)0)
  #endif

nullptr

• 为解决NULL的二义性，使用nullptr替代空指针，对应参数为 void*