[Lecture Notes] CMU 15-618 Parallel Computer Architecture & Programming

Lec 2. A modern multi-core processor

4 key concepts: 两个与parallel execution有关, 两个与challenges of accessing memory 有关

Parallel execution

  • Superscalar processor: Instruction level parallelism (ILP)
    • ILP读未来的指令(每个周期读两条指令),有两个fetch/decode单元和两个exec单元,能够同时执行两条指令
  • Multi-core: 多个processing cores
    • 多核之前,处理器提升重点在更大缓存,更多分支预测predictor等;同时更多晶体管(才能放得下更多缓存和更多predictor和乱序执行逻辑)促生更小的晶体管,促进更高的计算机主频
    • 2004年多核出现之后,人们在一个chip上放多个processor,用更多晶体管放更多核心。
  • SIMD processing (aka Vector processing): 多个ALU(同一个core内)
    • 仍然只需要一个fetch/decode单元,多个ALU。
    • conditional execution: 如果想simd的程序块有if else,要通过mask处理
    • 手写avx代码(cpu指令)是explicit SIMD; 而GPU是implicit SIMD,因为compiler生成的并不是并行指令(是普通的scalar instructions),只有GPU硬件运行才是SIMD的

Accessing memory

  • cache: reduce latency

  • prefetching reduces stalls: hides latency

  • Multi-threading, interleave processing of multiple threads

    • 跟prefetching一样,也是hide latency,不能reduce latency
    • 指的是:开多个线程,在一个线程卡住的时候执行别的线程
    • 在下图中,创建thread1的时候不仅仅创建thread1,还会告诉电脑创建了thread 2 3 4,硬件检测线程是否发生了stall(被等待内存操作卡住),如果发生了stall会很快切换到别的线程,想juggling一样。硬件决定如何juggle这些线程。
    • 这样memory latency仍然存在,但是被hide了。memory latency在后台发生,前台CPU一直在执行有用的工作。
    • 这种操作会导致单个线程的执行时间变长(因为thread1从runnable到重新开始执行有一段空挡(这段空隙在执行thread 2 3 4)。
    • 需要更多硬件资源,存储线程的register等状态信息,这样切换线程才会快。且需要较大的memory bandwidth。

GPU设计成处理大量数据(远大于核内缓存的数据量)。

与CPU内存相比,GPU显存带宽更高,但延迟更高。

Lec 3. Programming Models

Abstraction vs Implementation

abstraction和implementation不是一个东西!

ISPC (Intel SPMD Program Compiler)

SPMD: single program multiple data

一个ISPC计算sin(x)的例子

  • Interleaved

    • export void sinx(
        uniform int N, uniform int terms,
        uniform float* x, uniform float* result) {
        // assume N % programCount = 0
        for (uniform int i=0; i<N; i+=programCount) {
          int idx = i + programIndex;
          // 不重要
        }
      }
      <!--code0-->
      
      

在这个示例中blocked比interleaved更好。

因为每个iteration工作量完全相同,SIMD指令load连续内存(直接_mm_load_ps1)比load不连续内存(这种操作叫gather,只在AVX2及以后才支持)更快。

我们可以使用foreach来代替。

foreach表示循环中的每一次iteration都是独立的,由ISPC决定如何分配

1
2
3
foreach(i = 0 ... N) {
// index ...
}
  • ISPC: abstraction VS. Implementation
    • Programming model: SPMD
      • 程序员想的是,我们有programCount个逻辑指令流,写代码也按照这样的abstration去写
    • Implementation: SIMD
      • ISPC输出SSE4或AVX这种指令来实现逻辑操作

Four Programming Models

  • Shared Address Space model

    • 共享内存,不同线程通过读写同一块内存来通信

    • 很多人认为这是最方便易用的model,因为这和sequential programming最相近

    • 每个处理器都能访问任意内存地址

    • Uniform memory access time; Symmetric shared-memory multi-processor (SMP), : 指每个处理器都能访问内存且访问内存所需时间一致。Scalability不好。

    • Non-uniform memory access (NUMA): 每个处理器都能访问完整内存,但所需时间不一致。这样Scalability比较好。但是performance tuning可能需要更多精力。

  • Message Passing model

    • 线程之间不共享内存(address space不共享),只能通过发送和接收messages通信
    • 相比shared memory的优点:不需要别的硬件,可以纯软件做,实现起来简单。常用的库是MPI (message passing interface)
    • 现代很常用的操作是,在一个节点(节点内是多核的)内用shared address,在不同节点间用message passing

这里很搞:abstraction can target different types of machines.

分清abstractionimplementation的区别!

比如说:

message passing 这个 abstraction 可以使用硬件上的 shared address space 来 implement

  • 发送/接收消息就是读写message library的buffer。

shared address space 这个 abstraction 在不支持硬件 shared address space 的机器上也可以用软件来 implement (但是低效)

  • 所有涉及共享变量的page都标记成invalid,然后使用page-fault handler来处理网络请求
  • The data-parallel model

    • 上面的shared address space和message passing更general
    • data-parallel更specialized, rigid。
    • 在不同的数据(数据块)上执行相同的操作。
    • 通常用SPMD的形式:map (function, collection),其中对所有的数据都做function的操作,function可以是很长的一段逻辑,比如一个loop body。collection是一组数据,可以包含多个数据。
    • gather/scatter: gather是把本来不连续的数据按照index放到一起,scatter是把本来连续的数据分散开。
  • The systolic arrays model

    • 读内存太慢了,memory bandwidth会成为瓶颈

    • 所以要避免不必要的内存读取,尽量只读一次内存,做完所有需要用到这块内存的操作,再写回去。

    • 示例:矩阵乘法 https://www.youtube.com/watch?v=2VrnkXd9QR8

Lec 5. Parallel Programming Basics

  • Decomposition

    把问题分解成 tasks

    Main idea: 创造至少能把机器占满的tasks数量,通常对于t个processor会给多于t个task,并且要让这些task尽可能 independent.

    Amdahl’s Law: 程序中有S部分只能顺序运行(无法用并行加速)则整个程序的speedup \leq 1/S.

    通常是programmer负责decomposition。

  • Assignment

    Goal: balance workload, reduce communication costs

    can be performed statically or dynamically

    • statically: e.g. ISPC foreach
    • dynamically: e.g. ISPC launch tasks 运行的时候会维护线程池,线程池中的线程从任务队列中读。这样做的优点是runtime workload balance.
  • Orchestration

    Goal: reduce communication/sync cost, preserve locality of data reference, reduce overhead

    需要考虑机器的特性(上面的decomposition和assignment不用太考虑)。

    包括

    • structuring communication: 信息传递模型 e.g. 传一个chunk数据而不是只传一个byte,节约overhead
    • adding synchronization to preserve dependencies
    • organizing data structures in memory
    • scheduling tasks
  • Mapping to hardware

    对程序员来说是optional的。programmer可以显式制定哪个thread跑在哪个processor上。

    • mapping by OS: e.g. pthread
    • mapping by compiler: e.g. ISPC maps ISPC program instances to vector instruction lanes
    • mapping by hardware e.g. GPU map CUDA threads to GPU cores

    Mapping 还能有不同的decisions,比如说

    • Place related threads on the same processor: 最大化locality,共享数据,减少通讯成本
    • Place unrelated threads on the same processor: 可能一个thread受制于内存带宽,另一个受制于计算时间,这两个thread放在一起可以让处理器利用率更高

A parallel programming example: 2D-grid based solver

TODO here.

Lec 6. Work Distribution & Scheduling

key goals:

  • balance workload
  • reduce communication
  • reduce extra work (overhead)

workload balance

  • Static assignment

    任务分配在运行之前就已经 pre-determined

    例如之前讲的blocked assignment和interleaved assignment.

    Zero runtime overhead

    当任务量可预测时可以使用。不一定要任务量相同,可预测不会变就行。

  • Semi-static assignment

    可预测未来短期内的任务量

    一边运行一边profile并调整任务分配(periodically profiles itself and re-adjusts assignment)

  • Dynamic assignment

    任务量unpredictable.

    • while (1) {
        int i;
        lock(counter_lock);
        i = counter++;
        unlock(counter_lock);  // 或使用 atomic_incr(counter); 代替
        if (i >= N) break;
        // do with index i
      }
      <!--code2-->
      
      

Scheuling fork-join programs

  • Bad idea: cilk_spawn --> pthread_create, cilk_sync --> pthread_join

    因为创建kernel thread开销很大。

    应该用线程池。

  • 让idle thread 从别家thread的queue里steal work to do.

    continuation first:

    • record child for later execution
    • child is made available for stealing by other threads (child stealing)
    • 在遇到spawn的时候,自己执行spawn后面的任务,并把spawn出来的放在自己的work queue里,等待别的线程(如果别的线程有空闲)steal自己的任务。
    • 如果没有stealing,那么(相比于去除所有spawn语句)执行顺序全都是反的

    child first:

    • record continuation for later execution
    • continuation is made available for stealing by other threads (continuation stealing)
    • 遇到spawn的时候,只创建一个可被steal的项目。
  • work queue可以用dequeue (double-ended queue)实现

    每一个线程有自己的work queue,针对自己的work queue,在尾部添加,从尾部取出

    如果要steal别的线程的work queue,从头部取出

Lec 7. Locality, Communication, and Contention

Lec6讲如何平均分配任务,Lec7讲如何降低communication开销.

  • synchronous (blocking) send and receive

  • non-blocking asynchronous send and receive

send()和recv()函数会立即返回 在后台做事

Pipeline

使用Pipeline: Latency 不变, Throughput 增加

例子:

  • Communication = Overhead(橙色) + Occupancy (蓝色) + Network delay (灰色)

  • 最长的部分是瓶颈,决定了throughput上限

  • Overlap: communication和其它工作同时运行的时间。

    我们希望能尽可能增加overlap这样communication cost才会降低。

    降低overlap的方法

    • Example 1: Asynchronous message send/recv 异步消息
    • Example 2: Pipelining 发送多条消息时让这个发送过程overlap

Communication

Communication包含inherent和artifactual

  • Inherent communication: 程序算法写好的,必须发生的通信

    • Communication-to-computation ratio: 通信量/计算量 的比值。越低越好。
    • arithmetic intensity: 1/communication-to-computation ratio. 越高越好。
  • Artifactual communication: 所有别的通信,因为memory hierarchy导致额外的通信,例如L1/L2/L3/内存/网络之间的通信。包括:

    ① 系统有minimum granularity of transfer: 即使只需要读取4byte数据,也需要复制64-byte整条cache line

    ② 系统有rules of operation: 例如,写入内存需要先把内存读到cache line中(write-allocate)之后踢出cache line再写入内存,导致一次写入操作需要访问两次内存

    Poor placement of data in distributed memories: 被某个processor访问最多的数据并没有放在这个processor附近

    ④ Finite replication capacity: 因为cache太小放不下,会被踢掉,所以有一些数据频繁被踢出/放入cache

    提高locality对降低artifactual communication很重要

提高temporal locality的例子

  • by changing grid traversal order

  • by fusing loops

  • by sharing data

提高spatial locality

  • false sharing 不好

  • 4D array layout (blocked data layout): Embedding a 2D array within another 2D array allows page granularities to remain within a tile, making it practical to map data to local portions of physical memory (thereby reducing cache miss latencies to main memory).

Contention

Contention: 在短时间内很多人请求同一个resource

Example: distributed work queues (让每个线程有自己的work queue)可以降低contention

Summary

  • 降低communication costs
    • Reduce overhead: 发更少的消息数量,更长的消息内容(合并短消息)
    • Reduce delay:提高locality
    • Reduce contention: 把contended resource分开,例如local copies, fine-grained locks
    • Increase overlap: 用异步消息、pipeline等 提高communication和computation的overlap

Lecture 9. Workload-Driven Perf Evaluation

  • Super-linear Speedup:
    • processor足够多的时候,每个processor分到的数据fits in cache
  • Decreasing Speedup:
    • 随着processor增多,communication占比太大了
  • Low speedup:
    • Increasing contexts are hyperthreaded contexts (?)

Resource-oriented scaling properties

  • Problem constrained scaling (PC)
    • 更快速解决同一个问题
  • Memory constrained scaling (MC)
    • 不爆内存的情况下运行最大能完成的任务
  • Time constrained scaling (TC)
    • 同样的时间内完成更多任务

Simulation

Execution-driven simulator

  • 模拟内存,模拟内存访问
  • 模拟器的performance通常与模拟的细节数量成反比

Trace-driven simulator

  • 在real machine上运行real code得到内存访问的trace,或者用execution-driven simulator生成trace
  • 然后在模拟器上运行trace

Lec 10. Interconnects

Interconnect terminology

Terminology

  • Network node: 网络终端,会产生或消耗traffic,例如processor cache

  • Network interface: 把nodes和network相连

  • Switch/Router: 将固定数量的input links与固定数量的output links相连

  • Link: 传输信号的线缆

设计interconnection需要考虑的因素

  • topology: 怎么相连

    • topology的属性:

      • routing distance: nodes之间的长度,nodes相连需要多少个links (hops)

      • diameter: 最大routing distance

      • average distance: 平均routing distance

      • direct / indirect networks

      • bisection bandwidth

      • blocking vs non-blocking: 如果任何两个pairs of nodes可以同时传输,不相干扰,则为non-blocking。大部分network都是blocking的

  • routing: 消息沿什么路线传输到达目的地?可以static可以adaptive

  • buffering and flow control

Buffering and Flow control

和14740的第一节课讲的很像

Lec 11. Perf Tools

性能测试工具

GProf

  • compiler flag -pg
  • places a call into every function --> call graph (total time in each function)
  • 先跑程序,然后单独使用 gprof 命令(不传参数)

Perf

  • 有硬件指令测量性能计数器:cache misses, branch mispredicts, IPC, …
  • perf stat (同时只能开启4个counter)

VTune

  • similar to perf: analysis across counters
  • 有图形界面和解析

Debug工具

Valgrind

  • heavy-weight, 需要 shadowing
  • 有大量的overhead,不要用它测试performance
  • valgrind --tool=memcheck

Address Sanitizer

  • GCC and LLVM support, 有编译器支持
  • overhead比valgrind小一些
  • -fsanitize=address

Advanced analysis

Pin (Pintool)

  • acts as a virtual machine: reassembles instructions
  • can record every single instruction/block(无跳转)/trace(可能跨函数)

Contech

  • compiler-based (uses clang+LLVM)
  • record control flow, mem access, concurrency
  • traces analyzed AFTER collection

Summary questions

  • Reproducible?
    • Do you have a workload? Is the system stable? (Stable是说每次运行性能差距不能太大)
  • Workload at full CPU?
    • Other users using CPU? Does workload rely heavily on IO?
    • 使用time / top 看cpu占用时长
  • Is CPU time confined to a small number of functions?
    • 占用时长最长的函数?算法复杂度?
    • gprof / perf
  • Is there a small quantity of hot functions?
    • perf / VTune

Lec 12. Snooping-based Cache Coherence

Recap:

  • write-allocate: 如果写入的内存不在cache中,则需要先把memory读到cache中再写入cache
    • write-allocate与write-through和write-back都可以配合使用,但通常write-allocate与write-back搭配
  • no-write-allocate: 直接写内存
  • write-through: 同时写入cache和memory
  • write-back: 只写入cache,之后再flush to memory

Memory coherence

A memory system is coherent if:

  • the results of a parallel program’s execution are such that for each memory location, there is a hypothetical serial order of all program operations (executed by all processors) to the location that is consistent with the results of execution.

    与某一个serial的内存访问顺序的结果一致

Said differently

Definition: A memory system is coherent if

  • obeys program order: 一个processor先write再read一定读到新值
  • write-propagation: P1先write,一段时间后(suffciently separated in time) P2再read,则P2一定读到新值。注意此处需要相隔多久并没有定义。
  • write serialization: 两个processor写入同一个位置,则大家都必须agree on one order, 大家都同意这个顺序

可以用软件或硬件的方法解决

软件解法:OS采用page fault来propagate writes

硬件解法:Snooping based (本节课), Directory-based (下节课)

Snooping

Cache controllers monitor (snoop) memory operations.

任意一个processor修改cache都会broadcast通知其它所有人

现在cache controller不仅需要响应处理器,还需要响应其它cache的broadcast

Assume write-through:

Write-through is inefficient: 因为每次write操作都需要写入内存,也是因为此原因write-through不常见

MSI write-back invalidation protocol

Cache line加上dirty bit,如果dirty bit=1,代表处理器拥有这条cache line的exclusive ownership (Dirty = Exclusive)

如果其它processor想读同一条cache line,能从具有exclusive ownership的processor的cache中读(owner is responsible for supplying the line to other processors)

Processor也只能写入M-state的cache line; Processor能随时修改M-state的cache line不用通知他人

Cache controller监听是否有别人想要exclusive access,如果有,则自己必须要invalidate

MESI invalidation protocol

在MSI中read会有两种情况

  • BusRd: 从I转成S,即使并没有真正shared
  • BusRdX: 从S转成M

添加一个新的E-state (exclusive clean) 代表exclusive但尚未修改

读取时询问别的cache有没有这条cache line,如果别人也有则从I进S,否则从I进E

从E升级到M不需要通知他人

Other (MESIF, MOESI)

  • MESIF: F = Forward
    • 类似与MESI,但是在多个cache都shared的时候,有一个cache不在S而在F
    • F holded my most recent requester
    • F负责service miss(给别人提供数据),作为对比,MESI中所有S都会做出响应
    • I不能直接进入S,你要么进E(如果没人有), 要么进F(别人也有cache,但你是most recent requester所以你要负责给别的cache line提供data)。E和F随后有可能进入S。
    • Intel处理器用的是MESIF
  • MOESI: O = Owned but Not Exclusive
    • 在MESI中,从M转成S需要flush to memory,MOESI添加O-state,从M转成O但不flush内存
    • 在O时,这条cache line负责给别人提供data,但不flush进内存

Inclusive property

L1 L2 L3多级缓存

如果让所有L1缓存间和L2缓存间都interconnecting,那么效率低,所以让L2之间interconnect,并让L2 inclusive,即L2缓存中包含所有L1缓存,由L2控制L1的invalidate等

Inclusive property of caches: all lines in closer (to processor) cache are also in farther (from processor) cache. e.g. contents of L1 are a subset of L2

如果单纯让L2比L1大,不能自动保证inclusion

需要让L1和L2相互交流,L2中维护一个“是否在L1中”的bit

GPU don’t have cache coherence

每个cache都必须监听并对所有broadcast做出反应,这样interconnect开销会随着processor数量增长而增长

所以Nvidia GPU没有cache coherence,而是用atomic memory operation绕过L1 cache访问global memory

Lec 13. Directory-based Cache coherence

上一节课讲Snooping

Snooping-based cache coherence需要依赖broadcast工作,每次cache miss时都要与其它所有cache通信

存在scalability问题

One possible solution: hierarcical snooping。缺点是root会成为瓶颈;延迟;不能用于不同的拓扑结构

Scalable cache coherence using Directories

在一个地方存directory,每条cache line的directory信息存储着所有cache中这条cache line的状态

用point-to-point messages代替broadcast来传数据

Directory中包含 dirty bitpresence bit, 第k个presence bit代表第k个processor是否有这条cache line

Distributed directory: directory与memory大小同步增长

Example 1: Read miss to clean line

Processor 0把位于1的内存数据读到了自己的local cache 里,对应的directory记录P0有值

Example 2: Read miss to dirty line

本来P2的local cache中有dirty的数据

P0想读,发送read miss消息,P1告诉P0目前P2有dirty数据,P0收到后去向P2请求数据,P2将数据发给P0并将状态设置为shared,位置1的directory presence bit记录目前P0和P2有数据

Example 3: Write miss

P0有一条cache line,将要写入这条cache line,因此先请求找出有哪些Processor目前有这条cache line(找出sharer ids)然后向它们(P1和P2)发送invalidate请求,收到P1和P2的ack之后代表它们两个已经invalidate,此时再进行写入内存操作。

Advantage of directories:

  • 在read时,directory能直接告诉节点应该去问谁要数据,仅需要点对点通信:如果line is clean, 从home node要;如果line is dirty, 从owner node要。
  • 在write时,directory告诉sharer id,工作量取决于有多少节点在共享数据。极端情况,如果所有cache都在共享数据,则需要与所有节点通信,像broadcast一样。

Limited pointer schemes

presense bit需要占用存储空间,会导致storage overhead

Reducing storage overhead

  • increase cache line size: 让占比减小(M减小)
  • group multiple processors into a single directory node (让P减小)
  • 除此之外还能使用 limited pointer scheme (降低P) 和 sparse directories

Limited pointer schemes: 只存指针(指针=processor的id)

如果指针溢出,有几种不同的实际方法

  • 指针溢出时改为broadcast(添加一个additional bit代表指针不够用)
  • 设置最大共享者数量,不允许超出,如果超出,老的sharer被移除
  • 指针溢出时改为bit vector representation

Sparse directories

Key observation: majority of memory is NOT resident in cache.

Sparse directories只存一个指针,而在processor的cache line上存prev和next指针

优化:Intervention forwarding, Request forwarding


一些基础知识

ISPC

ISPC代码调用时会生成多个program instances, 可以利用 programCountprogramIndex 来获取instance总数和当前instance编号。

uniform 表示在一个SIMD程序块中,变量对所有SIMD通道都是相同的值。仅仅是一种优化,不影响正确性(因为uniform变量只需要加载一次或执行一次,编译器可以做出优化,不加uniform可能造成不必要的重复计算)。

非uniform (varying) 表示变量在不同SIMD通道可能有不同的值。

所以说 programCount 是 uniform, programIndex 是 varying.


ISPC可以通过tasks来实现多核加速,利用多线程。

Contrary to threads, tasks do not have execution context and they are only pieces of work. ISPC编译器接受tasks并自行决定启动多少个threads。

通常我们应该启动比cpu逻辑线程数更多的tasks数量,但也不要太多,否则会有scheduling的overhead。

task自带 taskIndex

CUDA

host是CPU, device是GPU

__device__: 在device上执行,只能在device中调用

__global__: 在device上执行,只能在host中调用。叫做kernel,返回值必须是void

__host__: 在host上执行且只能在host上调用

cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost)


Threads, Blocks, Grids

threads grouped into blocks

需要指明blocks的数量,和每个block中threads的数量。

假设n是总的threads数量, t是每个block中threads的数量。

KernelFunction<<<ceil(n/t), t>>>(args)

每一个thread都会运行同样的kernel,每一个thread由blockID和这个block中的threadID来标识。


Example:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
__global__ void vecAddKernel(float* A, float* B, float* C, int n) {
int i = threadId.x + blockDim.x * blockId.x;
if (i<n) C[i] = A[i] + B[i];
}
void vecAdd(float* A, float* B, float* C, int n) {
int size = n * sizeof(float);
float *d_A, *d_B, *d_C;

cudaMalloc((void **) &d_A, size);
cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
cudaMalloc((void **) &d_B, size);
cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);
cudaMalloc((void **) &d_C, size);

vecAddKernel<<<ceil(n/256), 256>>>(d_A, d_B, d_C, n);

cudaMemcpy(C, d_C, size, cudaMemcpyDeviceToHost);
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
}

注: 为什么cudaMalloc第一个参数是二级指针,而不直接使用返回值来赋值给指针?

因为 cudaMalloc 的返回值已经用来返回 cudaError_t


grid和blocks可以是1D, 2D, 3D的。上面这个例子是1D,所以是".x"

2D的例子:假设要把一个WIDTH x WIDTH的矩阵P分成几块。

WIDTH=8, TILE_WIDTH为2的话,就是把8x8的矩阵分成16个小块(grid),每一个小块大小是2x2(4个thread)。

1
2
3
dim3 dimGrid(WIDTH / TILE_WIDTH, WIDTH / TILE_WIDTH, 1);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
MatrixMulKernel<<<dimGrid, dimBlock>>>(args);

每一个thread可以用以下方式来标识

1
2
row    = blockId.y * blockDim.y + threadId.y;
column = blockId.x * blockDim.x + threadId.x;

为什么用两层threads?因为组成多个grid的thread blocks比一个很大的单个thread block更好管理。

GPU有很多很多核心,核心group成SM(streaming multiprocessors),每一组SM有自己的内存和调度。

GPU不同时启动所有100万个threads,而是把大约1000个thread装进一个block里,并分发给SM。

assign给SM的thread block会使用SM的资源(寄存器和共享内存)。这些资源已经pre-allocated,且由于寄存器数量很多,在切换threads时不需要register flush。


不同的block可以用任何顺序运行,因此不能assume block2在block1之后运行。如果真的要这么做,需要放在不同的kernel里(启动kernel比较耗资源)

同一个block中的thread可以使用 __syncthreads() 来做barrier synchronization。

但是通常不建议使用 __syncthreads()


如何选择合适的block size?

  • Consideration 1: hardware constraints
    • 例如:每一个SM分配小于1536个thread,小于8个block;每一个block小于512个thread
  • Consideration 2: complexity of each thread
  • Consideration 3: thread work imbalance.

GPU memory

Global memory很慢,所以同时运行大量线程,线程因为内存IO卡住的时候切换其它线程,这是massive multi-threading (MMT).

这样总的throughput很高,即使每个thread的延迟也很高。

每个SM有自己的scheduler,每个SM存储了所有thread的context(PC, reg等),所以SM内能做到零开销线程切换。同时,SM scheduler有一个scoreboard追踪哪些thread是blocked/unblocked,所以SM有大约30个核但可以运行大约1000个线程。


Tiled MM是一种进行矩阵乘法 内存友好的方法。

CUDA类型关键词

  • __device__ __shared__ memory: shared; scope: block; lifetime: block
  • __device__ memory: global; scope: grid; lifetime: application
  • __device__ __constant__ memory: constant; scope: grid; lifetime: application

Race conditions:

CUDA中难以实现mutex,而且包含critical sections的代码在GPU上本来就运行得不好。

CUDA中有一些原子操作,可以在global或shared memory变量上操作

  • int atomicInc(int *addr): 加一,返回旧值
  • int atomicAdd(int *addr, int val): 加val, 返回旧值
  • int atomicMax(int *addr, int val): 让*addr=max(*addr, val) 并返回旧值
  • int atomicExch(int *addr1, int val): set
  • int atomicCAS(int *addr, old, new): Compare and swap.
    • if (*addr == old) *addr = new;

[Lecture Notes] CMU 15-618 Parallel Computer Architecture & Programming
https://www.billhu.us/2025/054_cmu_15618/
Author
Bill Hu
Posted on
January 7, 2025
Licensed under