Why Parallelism? Why Efficiency? CMU 15-418

• what is a parallel computer?
  • a parallel computer is a collection of processing elements that cooperate to solve problems quickly
• baseline number computation
  • add up 16 numbers
  • first we gonna use single digit numbers
    • with two processors
  • start adding it up
    • why it isn’t twice as fast
  • more realistic than the others
• parallel software
  • if work isn’t divided into equal stuff, the efficiency of the machine will be slower.
  • the load imbalance is one of the major challenge in the parallel software
• quick review
  • a big light bulk be put into a big box
  • anyway that sort of
  • communication latency
    • two second pieces
  • over and over again
  • sort of computer architecuture
  • extract the parallelism
  • a lot of raw computing hardware
  • adding more and more hardware to the system
• spread into two processor
  • 25% slower than the former one
  • splitting this into k different threads
  • lower level of hardware
  • has this general assignment into it
  • communication cost
• commercial designs have given their own concept
  • 32 bytes = 256 bits
  • __mm256
  • load ts–ps
• how to extract parallelism
• each SIMD is talk about something
• relatively low cost
• GCC automatically vectorize this code
  • this have been a lot of code in this
• the interesting feature about this

sticky point

it doesn’t matter what this will do

set up a mask

flatten out some of the data we will

general idea of handle this

• SIMD operation works
• main parts of GPU
  • coherent coherence in this sense means
  • as i mentioned before, like the Gates

GPU

• SPMD single program multiple data
• n value each and n value result
• multi-core and simd

getting data in and out of this system

• link the memory to the processor
• one gigahertz

multiple levels of caches

• private to the L3 shared the cache of the course
• go through the L3 cache the DRAM
• caches actually help the
• 代码只是注意到了一部分地址，并且开始猜测部分需求，顺便泄漏各种信息
• 基于所击中的内容

ISPC

ISPC is kind of language download for

each of logical thread is doing the same code

completely different procedures

but in our class, we’re doing the same code

• Recall: example program from last class

  • lot’s of values in this

• ISPC: the way its works

  • if we animate what’s this look like
  • in the main file here, you are executing sequentially
  • sinx is the entrance of the parallelism
  • there are a couple of valuable program count and program index
  • how many things are we running in parallel
  • the runtime system decide what the time should be
  • So you can see what the program count here
  • your instance number from zero to instance one

• instead of going round robin

• 

• contiguous

  • load 0 4 8 12
  • it turns out there is an instruction that allow you to do this
  • into one vector

• for each

  • a for loop
  • tell the system you can execute the loop in parallel
  • the source of paralleism
  • in the case of ISPC, the code we do is the first code we see

• how was that implemented

• kernel support for that

• ISPC

  • the operating system did not involve
  • the compiler did a lot of things

Three models of communication

1. Shared address space
2. Message passing
3. Data parallel

Shared address space

各个线程能够互相作用是因为机器有公共的地址空间，它们之间可以拥有私有变量。如何我能够知道我现在得到的是不是原来的版本，如果我们在增加一个变量，尝试增加相同的变量的话。

我们有一个共同的硬件空间

• 从物理层面分享硬件的空间
  • 有一条总线，多个处理器插入在总线当中，并且所有的都可以去访问总线
  • 内存不会离任何处理器更近
    • 那么如何连接处理器和内存之间呢？
    • 如果我们增加更多处理器的数量呢？
    • 总线上永远只能够传一个东西
      • 那么你可以做很多东西
  • 真实硬件是什么样的

• Intel的Core会有一个Shared L3 Cache

• 没有做共享的内存

  • 而是让每个人都拥有了一块自己的共享内存

• 我们可以去获取自己的本地内存

• 

• 两个好处

  • 更快，因为离自己的处理器更近
  • 从带宽角度更快，流量不会更大而减慢其他人的速度

cache

当我们尝试cache的时候，可能会有一些协同上的问题

部分原因是实现缓存一致性而需要的成本

所以我们接下来要谈论的是信息的传输的成本

send it over to the other thread

你想要做很多一样的事情，可以充分利用高端机器的支持，因为它们仅仅在物理上共享内存，会在共享内存空间。

不要混淆抽象和实现，我们确实谈论了关于共享地址空间的硬件和消息传输的硬件，但事实证明可以在任何类型的硬件上实施这两个模型，这两种不是被严格绑定到一种硬件类型之上。

对于消息模型实际上可以在共享内存空间的模型上实现，只需要传递一个指向内存的指针就可以。

同时也可以在消息模型中实现内存的相互传递和转换，但是速度会比较差。

The data-parallel model

GPU是围绕着数据构建的并行化而产生的

ISPC，有一些关于数据并行性的

• 比如foreach

  • 是控制的并行性质

• 会造成完全不同的

• Streams

  • 元素的集合，元素会被独立处理
  • 一种实现数据并行性并避免一些数据竞争或问题的方法

• Kernels

  • 内核，无副作用的函数

  

• gather

  • 将input转为indices版本的tmp_input

• scatter

  • 将tmp_output转为indices版本的output

• 我们需要生成结果

• 如何将事物连续地粘合在一起

Recall

Imagine the ocean is a three dimension thing

a very thin dimension object

他们将其表示为一组二维平面

如何足够远离某颗恒星，那么可以忽略其问题，并且可以对一些聚簇进行同化，将其视作一颗较大的恒星-n-body simulaiton

Barnes-Hut tree

二维是四叉树，三维是八叉树

这是一个identify things that we can do parallel

最后我们需要拥有这些不同的线程

• 海洋任务
  • 元素的数量比处理器多
• n-body simulation
• 将工作分为各个小块
• 根据具体情况编排工作
  • 你需要比线程数量更多的任务，数量足够多，但是这样管理就需要一些成本或者资源
• 艾姆代尔法则

• 我们正在顺序运行的一部分程序，最终会对最大值产生一个上限，我们可以通过并行性得到更好的性能
• a simple example
  • 对于每个像素的计算都是可以并行的
• 一种平衡负载的方法是随机将一些分配分配给任何其他的处理器

分配

• 对分配阶段而言，可能是静态的，也可能是动态的
• 静态分配是在程序运行之前就已经进行了任务的分配
• 动态分配是在程序运行之后在沿途进行改变
• a simple example for static
  • 假设只有两个处理器，我们正在做一个数组，我们计算输入值，将数组一分为二，pthread_create创建了一个子线程，子线程处理数组左边的，父线程处理数组右边的
    • 代码是硬编码的，如果我们在一个four cores的cpu上进行操作的话可能出现问题
    • 缺点与风险：工作量可能会有很多的变化，左手边的变化可能比右手边的变化有问题，我们可能没有很好地平衡工作量的问题
  • 这只是一套工作，就像在工作时从队列当中
• 如果我们尝试使用一个队列来安排任务的话
  • 那么需要考虑多个处理器同时取队列中任务的情况，可能需要在队列中加一把锁，这样就增加了使用动态调度和动态分配时候的运行时开销。

编排

in GPU this mapping is done by the hardware

花哨的偏微分方程，所以我们会极度简化计算，只是将元素加在一起

const int n;
float* A;
void solve(float* a){
	float diff, prev;
	bool done = false;
	while(!done){
		diff = 0.f;
		for(int i = 1; i<n;i++){
			prev = A[i,j];
			A[i,j] = 0.2f * (A[i,j] + A[i,j-1] + A[i-1,j] + A[i,j+1] + A[i+1,j]);
            diff += fabs(A[i,j] - prev);
		}
	}
	if(diff/(n*n) < TOLERANCE) done = true;
}

他会迭代直到完成，将所有的差异相加，最终将这个值与某个阈值相比较，在内部循环中，有这个diff值。

如果我们完全并行处理它的话，那么会出现racing conditions?

Solutions:

• 将二进制速率分割为各个部分，然后在20 rounds之后再进行通信

• 我们可以重写循环结构，然后利用对角线的循环结构来处理程序，这样就没有racing condition的问题

  • 但是对于某些对角线来说，工作量仍然是不一样的，所以会出现一些问题
  • 但这样就可以引入cache performace

• 创建两个矩阵，创建另一个矩阵的副本

  • 对于科学程序员来说，这样复制数据是不好的

• red-black ordering 红黑排序

  • 我们更新所有棋盘当中的红色的元素，黑色元素暂时不变

  • 

  • 有确定性的优势

  • question

    • 当我们完成对红色数据的更新时候，黑色数据实际上是被更新过后的红色数据所包围的
    • 所以黑色的会得到更新过后的红色数据
    • 需要保证红色计算完成和对黑色的更新是完全分离的，所以我们会有一个障碍了（barrier）

• 所以红色全部生成之后，需要一个同步步骤来保证所有的都已经生成，并且需要开始交流沟通信息（最小化交流）

  • 
  • 所以对于这个而言，左边的可以在自己的cache或者内存当中直接得到这些数值
  • 最多只需要进行两个thread之间的通信

• const int n;
  float* A = allocate(n+2, n+2)；
  void solve(float* A){
      bool done = false;
      float diff = 0.f;
      while(!done){
          for_all(red cells (i,j)){
              float prev = A[i,j];
              A[i,j] = 0.2f * (A[i,j] + A[i,j-1] + A[i-1,j] + A[i,j+1] + A[i+1,j]);
              reduceAdd(diff, abs(A[i,j] - prev))；
              if(diff/(n*n) < TOLERANCE) done = true;
          }
      }
  }
  - 我们只是在外部循环中进行并行化
  
  - for_all原语告知可以进行并行化，并将其交给系统处理
  
      - 当我们在处理绝对值差异的时候，是一种特殊的减去操作。本地计算和并且加起来。
      - 基本上，代码的改变非常小
  
  - Shared address space
  
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240120180801.png)
      - 需要等到所有人吧赶上来才能继续前进
      - 进程之间的通信则仅仅需要一些load和store操作
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240120180923.png)
      - 通过线程id得到对应的最小值和最大值，得到需要进行操作的行的最小值和最大值
      - lock操作则用来保证diff的更新不会产生data race的问题
      - 瓶颈的问题在于
          - 我们在循环内存在一组lock&unlock，这导致了性能的瓶颈
      - stick with regular locks
      - 所以每个线程可以存储一个local diff，等到所有进行完毕之后再进行更新
  
  - 为什么有三个barrier
  
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240120181515.png)
      - 如果没有最后的障碍barrier的话，那么如果有线程跳出来就会返回到上方的diff处，从而reset了diff的值，从而可能部分线程会对终止条件产生误解和误差
      - 而第一个barrier是为了保证mydiff在进入下述循环的时候已经被重置
  
  - 其他的一种tricks
  
      - 永远不要使用普通的内存来作为同步工具，那样将无法工作
  
  - Message-passing expression of solver
  
      - 当使用消息传递的时候，进通过消息进行沟通，唯一能够访问的内存是私有内存。
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240120182501.png)
      - 我们会建立所谓的ghost row，来保证代码逻辑的完备性
      - 我们会上下传递ghost row
      - 对于最后的相加环节，会先确定一个特殊的处理器，在处理逻辑时，如果不是这个处理器，那么就将最终myDiff发给这个处理器，如果是这个处理器，那么就对这些结果进行处理
  
  - send和recv到底做了什么？
  
      - 当我们假设存在这种消息传送机制，我们用同步发送时候，知道receiver获得的时候，我们会block自己。
      - 解决方法
          - 让偶数的处理器先发送，让奇数的处理器先接受
          - 或者使用非阻塞式的发送和接受
  
  # Performance Optimization Part 1: Work Distribution and Scheduling
  
  ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121155644.png)
  
  当你尝试编写并行软件的时候，始终从实施开始
  
  永远尝试先去实施哪些更简单的解决方式
  
  我们如何应用静态分配？任务的运行时间需要可预测，简单的方法需要预测的话。
  
  我们展示了循环迭代的任务
  
  - 这段代码有什么问题吗
  - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121162247.png)
  - 问题：彼此或花费大量时间尝试夺取锁定
  - 解决方法：如果我们只是做一次迭代，那么可能会很慢，那么我们可以去设定多个工作数量，这是我们会得到多个迭代，这时候我们才会去检查。
      - 缺点：如果你将任务安排的太大，就可能会产生负载均衡的问题，所以我们需要一个较好的临界点。
  - 如果我们出现，之前的task都非常小，最后的task非常大的时候，这时候该怎么办？
      - 解决方法：如果知道这是一个大任务，那么也行先去调度它比较好
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121163054.png)
      - 你在调度的最后能拥有小的任务是最好的结果
  - 随着你的任务更多，希望一切都会顺利地变平滑
  - 动态调度实际实现是在哪个阶段
  - 分布式队列实现
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121164104.png)
      - 你可以直接去排队，没有其他可以进行竞争
      - 解决方法：steal task from other queues
          - 可能会涉及到一些通信上的问题
  - 集中化的事物
      - 如果我们偷工减料的话，最好完成多个任务
  - 如果我们到了任务的结束的时候，可能只有一个线程，那么会出现只有一个任务的时候，那么这种时候不会出现线程的steal方式，这里的处理方式可能会查看对应的queue当中的任务数量，再决定是否进行任务的偷取。
  - 总结
      - 挑战：做好负载平衡
      - 静态分配和动态分配
      - 半静态，定期衡量事物并做一些静态的分配事务。
  
  ### Cilk plus
  
  - `cilk_spawn`是将并行性暴露给语言的原语，表示这个运行的事务和之后的代码都是同时进行的
  - `cilk_sync`当所有的calls spawned by current function才会返回，将所有的并发的线程汇聚到一起
  - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121170614.png)
  - Cilk的生成不会强制系统实际执行或创建任何并发，可以直接忽略任何并发的Cilk生成，即使忽略任何Cilk的宏也可以保证程序正确运行
  - 从底层来看Cilk的运行情况
      - spawn is not a pthread_create
      - 并不意味着我们正在创建一个新的线程，实际上我们在告诉系统这是潜在的并行性。
      - 在这张图中，可以通过分而治之得到更大的因子，在仅仅几层递归之后，我们可以获得很多的任务
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121182739.png)
  - a naive implementation of Cilk
  - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121183027.png)
  - thread是由操作系统所定义的，如果仅仅使用pthread_create来代替cilk_spawn的话，那么不认为应该做我们可能会获得很多线程，会维护寄存器、堆栈和指令指针，当我们软件生成的thread大于硬件生成的thread之后，那么会需要暂停一个线程而启动另一个线程，这就需要进行状态保存和跳转，而这是代价非常昂贵的。
  
  ### 例子
  
  - foo and bar
  - 如果我们有两个物理上的线程，我们应该如何做来让他能够完成调度工作
  
  ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121183643.png)
  
  ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121184421.png)
  
  ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121184449.png)
  
  - 问题：为什么先做foo而不是bar
      - foo是创建出来的子任务
      - bar是原来的任务
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121184644.png)
      - 那么这个执行顺序有关系吗？
          - 到底是先执行子任务让其它线程抢夺本任务继续执行的机会呢还是先执行本任务让其它线程去抢夺子线程？
          - 先做自己的版本
              - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121184808.png)
              - 更像是广度优先的版本
          - 先做子进程的版本
              - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121185004.png)
              - 更像是深度优先的
              - 如果自己的事情被其他人做了，那么你自己就有空闲了，需要等待知道其他线程有事情需要你自己去做，**因为自己的队列当中可能没有需要做的事情了**。
              - 那么它会做什么呢？是一个问题
          - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121185429.png)
          - 从执行层面来看，后一种更加lazy一点，是一种懒加载的形式进行的操作，前一种放了很多的队列空间进行存储。后一种会更快一些。
      - 例子
          - 200元素的快速排序
          - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121185929.png)
          - steal的时候会去偷大的任务还是小的任务
              - 大的任务，因为大的任务实际上会放在前面做会更好
              - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121190140.png)
          - 从双向队列来说，不会产生一些并行上的问题，因为插入队列的地方和取出队列的地方是不一样的。
  - cilk_sync
      - 有一个数据结构，在跟踪其他事情，并保证每一个thread都完成
      - 当我们完成的时候，我们会与哪一个硬件线程进行交互，在哪个硬件线程上跑？
          - 最开始的那个线程。
              - 简单的实现但是速度慢
          - 同步后继续进行的不一定是最开始的硬件线程，同步的数据结构可以进行改变位置，是贪婪思想
          - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240121190816.png)
  
  # /Proc/CPU
  
  当你想要从操作系统内核获取信息时，你实际上需要做的是进行系统调用，但这会带来一些开销。
  
  - /proc是虚拟的文件系统
  - `ls /proc/37684`这可以看到关于文件描述符的信息，可以看到一个代表我的标准输入和标准输出的文件
  - slash proc
  - 可能有是处理器但实际上不是真正的处理器的
      - 40个处理器和20个siblings
      - cpu Mhz - speed now(Speed step)
      - cache size - outermost(L3) 最外层的cache
      - siblings - number of hyperthreads
      - processor - id of a hyperthreads
      - cpu cores
      - core id
      - physical id - socket
      - flags - note avx, avx2, sse, etc.
  - grep
  - 优化代码的时候，我觉得我能够拿到某些能力
  - 执行一次指令需要多少时钟周期？
      - 我尽力进行了优化，进我所能，我有一个非常不错的答案
  - 阿姆达尔定律
      - 如果你的份额只占到份额的10%，那么这样的常见活动的小优化倾向于达到更大的优化活动。
      - 即使在内部工作循环当中也会有一些小的改进，所以你会看到这样的东西，你想要真正发现的
  - proxy code
      - S trace traces the system calls
      - benchmark很重要，进行瓶颈测试
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240123135422.png)
      - ![](https://boogiepop1230.oss-cn-shanghai.aliyuncs.com/pics20240123135455.png)
      - uniform variable
          - 在所有实例当中都存在的变量
          - pipeline 如果从串行做了
  
  # GPU Architecture & CUDA Programming
  
  ## 什么是cuda
  
  **统一计算设备架构**（Compute Unified Device Architecture, **CUDA**），是由**NVIDIA**推出的通用并行计算架构。解决的是用更加廉价的设备资源，实现更高效的并行计算。
  
  几乎所有的编程语言，不使用特定框架，都只能实现CPU编程——std::thread也是将线程开在CPU中。而不同于每一位程序员都接触过的CPU编程，GPU编程可以使用**更多的流处理器**、**更多的线程数**。
  
  想象矩阵运算？
  
  如果是多线程编程的话？
  
  多个线程同时写入同一个变量是会很危险的，但同时读取同一个变量不会出现问题，所以如果不想引入非常耗时的互斥锁，我们就必须根据可写入变量的数量来划分线程数。
  
  > 假设是2x2的方阵，我们可以分成四个线程来计算，每个线程只计算结果矩阵的一个元素
  
  ```c
  template <typename T>
  void thrd_MatrixProduct(matrix<T>& C /* Out */, matrix<T> const& A /* In */, matrix<T> const& B /* In */,
  			size_t row, size_t col) {
  //每一个线程因row和col的不同而被区分
  	T ans = 0;
  	for(size_t k = 0; k < A.Width(); ++k) {
  		ans += A[row][k] * B[k][col];
  	}
  	C[row][col] = ans;
  }
  
  template <typename T>
  void MatrixProduct(threadPool& tp, matrix<T>& C /* Out */, matrix<T> const& A /* In */, matrix<T> const& B /* In */) {
  	assert(A.Width() == B.Height());
  	C.Set_size(A.Height(), B.Width());
  //将每个计算单元的任务放进线程池中
  	for(size_t i = 0; i < A.Height(); ++i) {
  		for(size_t j = 0; j < B.Width(); ++j) {
  			tp.addTask(thrd_MatrixProduct, C, A, B, i, j);
  		}
  	}
  //线程池启动并阻塞等待
  	tp.start();
  	tp.join();
  }
  

事实上，这个操作在数据量较小的时候看不出优势；而数据量过大时，线程调度又是操作系统的难题。我们需要可以容纳更多线程的廉价计算资源。

cuda正是给显卡计算这一廉价而高效的并行计算方式提供了接口，同时也不需要线程池的维护。比如上述问题，用cuda实现，或许过程有点复杂，但核心还是相当容易的：

template <typename T>
__global__ void kern_MatrixProduct(T* C_ptr, const T* A_ptr, const T* B_ptr,
				   size_t A_Height, size_t A_Width, size_t B_Width) {
//获取当前线程调用者的Id
	size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
	if(Idx >= A_Height * B_Width) return;
	
//规定计算第i行j列的线程Id为(i * width + j)
	size_t row = Idx / B_Width;
	size_t col = Idx % B_Width;
//核心计算
	T ans = 0;
	for(size_t k = 0; k < A_Width; ++k) {
		ans += A_ptr[row * A_Width + k] * B_ptr[k * B_Width + col];
	}
	C_ptr[row * B_Width + col] = ans;
}

template <typename T>
void MatrixProduct(matrix<T>& C /* Out */, matrix<T> const& A /* In */, matrix<T> const& B /* In */) {
	assert(A.Width() == B.Height());
	C.Set_size(A.Height(), B.Width());
	
	T *A_ptr, *B_ptr, *C_ptr;
//在显卡中申请内存（显存）
	assert(cudaSuccess == cudaMalloc(&A_ptr, sizeof(T) * A.Height() * A.Width()));
	assert(cudaSuccess == cudaMalloc(&B_ptr, sizeof(T) * B.Height() * B.Width()));
	assert(cudaSuccess == cudaMalloc(&C_ptr, sizeof(T) * A.Height() * B.Width()));
//将数据从内存复制到显存
	assert(cudaSuccess == cudaMemcpy(A_ptr, A.pos_ptr(), sizeof(T) * A.Height() * A.Width(), cudaMemcpyHostToDevice));
	assert(cudaSuccess == cudaMemcpy(B_ptr, B.pos_ptr(), sizeof(T) * B.Height() * B.Width(), cudaMemcpyHostToDevice));
//调用核函数
	size_t thread_count = 1024;
	size_t block_count = (A.Height() * B.Width() - 1) / thread_count + 1;
	kern_MatrixProduct<<<block_count, thread_count>>> (C_ptr, A_ptr, B_ptr, A.Height(), A.Width(), B.Width());
	cudaDeviceSynchronize();
	assert(cudaSuccess == cudaGetLastError());
//将数据从显存复制到内存
	assert(cudaSuccess == cudaMemcpy(C.pos_ptr(), C_ptr, sizeof(T) * C.Height() * C.Width(), cudaMemcpyDeviceToHost));
//释放临时内存
	assert(cudaSuccess == cudaFree(A_ptr));
	assert(cudaSuccess == cudaFree(B_ptr));
	assert(cudaSuccess == cudaFree(C_ptr));
}

1. CUDA是一种用于GPU编程的语言，可以将计算问题转化为图形操作。

2. CUDA使用了层次化的块结构来组织计算任务。

3. CUDA的线程模型与传统的线程模型有所不同，它更加紧密地耦合在一起，并且有一些限制。

4. CUDA可以在CPU和GPU之间进行数据交换和任务调度，实现并行计算。

5. CUDA是一种专有的编程语言，但在性能和广泛使用方面优于开源版本。

6. 在CUDA中，数据被分割成小的块，每个块中有1024个线程，但是有时候代码的执行不会刚好是块大小的倍数，需要进行边界测试和禁用操作。

7. CUDA使用SPMD（Single Program Multiple Data）模型，即多个线程同时执行相同的内核代码。

8. 在CUDA中，有全局内存、共享内存和每个线程的小内存，需要显式地在代码中进行内存分配和数据传输。

9. CUDA的硬件会自动处理块的分配和调度，不需要程序员手动创建调度器。

10. CUDA中使用同步操作来保证线程的执行顺序和数据的一致性，还有一些原子操作可以保证操作的原子性。

11. GPU的并发性与pthread模型不同，可以支持大量的线程块。

12. GPU硬件上有一个硬件调度器，可以将线程块映射到实际处理器上，并支持多个线程块的计算。

13. 编译的CUDA代码会生成机器码，描述了内核的具体操作、每个线程块的线程数、每个线程需要的本地数据量以及共享内存的最大使用空间。

14. GPU的执行单元可以通过超线程技术进行时间分复用，可以并行执行多个线程块的计算。

15. GPU的编程模型中，线程块内部可以共享内存，但线程块之间不能共享内存。

如何表现球面的反光是一个问题？

OpenGL你可以给你的系统添加参数

纹理映射是什么，可以将一些统一的模式来进行操作，营造一些奇怪的图案，这些都表明是所谓的光线追踪，这是一个更加复杂的事情。

GPGPU

“general purpose” computation on GPUs

早期的GPU并不局限于这一操作，将其转化为其他的GPU进行操作

如何在硬件上跑代码

CUDA programming language

• is CUDA a data-parallel programming model?
• is CUDA an example of the shared address space model?
• Or the message passing model?
• Can you draw analogies to ISPC instances and tasks? What about pthreads?

后文所提到的线程都是CUDA线程

三维向量矩阵数组表示法的数据，所以已经内置了多维数量。

多维度的业务是由CUDA直接支持的

CUDA内存模型

CUDA malloc

代码会被分为代码块，每个代码块将负责适当的计算和更新。

这就是比较简单的版本，它能够正常运行，但是速度不太快。

同步线程

_syncthreads()保证所有同时执行的线程都已经完成了对全局内存的读取

原子操作

float atomicAdd(float* addr, float amount)

Atomic operations on both global memory and shared memory variables

他们称之为，看起来像是我们在执行一个程序，但是与pthread完全不同，你不能分配那么多

Wrap

warp是一个线程块的32个元素块，SIMD处理技术来处理warp，它会动态地安排这些，并且提取出来这些操作。

也存在一些硬件上的问题

总结：当有一个计算块来的时候，我将其分为32个元素，每次称为一个warp，每个warp都是一个SIMD执行，但我可以在一个块内运行它的所有线程，我会有同样的事情要做，如果有某种条件，就像标准SIMD中一样，它应用选择性执行selective execution来决定要执行哪些操作或不执行哪些操作，如果某个warp的价值和某个条件测试相同，那么就不需要去做。所以不是整个块上的SIMD模型，但是规则必须是我可以按任意顺序执行代码块

Performance Optimization Part II：Locality, Communication, and Contention

主要关注如何将一个问题分解为多个问题，并保证不会因为各个模块之间的通信成本所淹没，我们专注如何将所谓的并行计算模块转换为某些并行计算。

Message-passing solver review

(since it makes communication explicit)

除了顶部的线程以外，所有线程都会将其数据发送到上方，除了底部的线程以外，所有的线程都会将他们的数据发送到下方，所有的线程的空间都会增加两行。

Synchronous (blocking) send and receive

当你发送消息的时候，你的信息发送消息的时候，你会阻止你的代码块知道你的对象收到这条消息，收到之后才会将数据进行复制和搬运。就像这样，可以容纳一条消息的问题

non-blocking send

发送即是立即发送，这样代码就可以，如果事情发展太快的话

如果一个发送方的线程一直比另一个线程更快，那么会出现无限增长的消息队列（？），所以会申请接收方会返回一个句柄（token），让发件人可以查询和找到发送是否实际发生了

可以做一些其他类型的工作，这里有一些类型参数，远离电脑

一次只能由一辆车获得的资源，指定了一项规则，规定一次只可能有一辆车，他们会希望能够在

Pittsburgh story

想象旧金山到这里车程是4000km，车的速度是100km/h

通常我们更关注带宽而不是仅仅减少其延迟

流水线技术

处理器中的流水线技术

现代处理器中可能会有15-20个阶段，每个阶段处理自己的事情

如果能够使得我们的管道阶段细粒度化的话，那么我们可以将我们的存储量加倍。

一个非流水线交流的简单模型

我们的沟通和周长都是相同的，你的算数强度和

cache miss的三种方式

cold miss

• 第一次访问即将错过的一些数据，因为数据没有在缓存中，且这样的情况无法避免

capacity miss

• 当我们有一些较大的计算时，被称之为工作集，处理器在一段较长的时间内操作的数据集合的类型

conflict miss

• 由于布局的原因，缓存用的是集合关联寻址或索引，这只是因为基本上这样做

communication miss

• 在系统的更高层次上，我们需要付出额外通信的代价

减少交流的技术

想象缓存容量是24个元素，但是必须按照每个盒子四个元素进行对齐，当你绕到一行的末尾并进行下一行的时候，你就已经被删除了

这是什么miss呢？

• 这是一种capacity miss