优达学城-并行编程-Unit2 通信模块、同步机制、原子操作

(一). Parallel communication Patterns

在上一章CUDA系列学习（二）CUDA memory & variables中我们介绍了memory和variable的不同类型，本章中根据不同的memory映射方式，我们将task分为以下几种类型：Map, Gather, Scatter, Stencil, transpose.

1.1 Map, Gather, Scatter

  • Map: one input - one output
  • Gather: several input - one output 
    e.g image blur by average
  • Scatter: one input - several output 
    e.g add a value to its neighbors 
    (因为每个thread 将结果scatter到各个memory，所以叫scatter)

图为Map, Gather & Scatter示意图:

fififififififiifififififiififif  :AOS 结构数组

转置后

ffffffffffffiiiiiiiiiiiii  ：SOA 数组结构

对需要大量操作其中一个元素时贼有用

1.2 Stencil, Transpose

  • stencil: 对input中的每一个位置， 
    stencil input：该点的neighborhood 
    stencil output：该点value 
    e.g image blur by average 
    这样也可以看出，stencil和gather很像，其实stencil是gather的一种，只不过stencil要求input必须是neighborhood而且对input的每一个元素都要操作 
    图示：

    1. 2D stencil: (示例为两种形式) 
       
    2. 3D stencil: 
  • transpose 
    input：matrix M 
    output: M^T 
    图示：
    1. Matrix transpose 
    2. Transpose represents in vector 
    3. fififififififiifififififiififif  :AOS 结构数组

      转置后

      ffffffffffffiiiiiiiiiiiii  ：SOA 数组结构

      对需要大量操作其中一个元素时贼有用

几种机制：

映射/转置：一对一的每个输入映射到单个唯一的输出

收集操作：很多对一很多可能的输入被选来计算一个输出

分散：每个线程在很多可能的目的地输出

模板：特殊输出 多对一

归约：全对一

扫描/排序：全对全

Exercise 
Q: 
看这个quiz图，给每个蓝线画着的句子标注map/gather/scatter/stencil/transpose: 

A:四个位置分别选AECB。 
这里我最后一个选错成B&D, 为什么不选D呢？看stencil的定义：如果是average，也应该对每一个位置都要进行average，而题目中有if(i%2)这个condition。


那么对于不同的Parallel communication Patterns需要关注哪些点呢？ 
1. threads怎样高效访问memory？- 怎样重用数据？ 
2. threads怎样相互交互部分结果？（通过sharing memory）这样安全吗？

我们将在下一节中首先回顾讲过的memory model，然后结合具体问题分析阐述how to program。

(二). Programming model and Memroy model

第一讲和第三讲中我们讲过SM与grid, block, thead的关系：各个grid, block的thread组织（gridDim，blockDim，grid shape, block shape）可以不同，分别用于执行不同kernel。 

如我们第一章所讲，不同GPU有不同数量的硬件SM（streaming Multiprocessors），GPU负责将这些block分配到SMs，所有SM独立，并行地跑。

2.1 Memory model

第二讲中我们讲了memroy的几种形式，这里我们先来回顾一下memory model.

每个thread都可以访问： 
1. 该thread独占的local memory 
2. block内threads共享的shared memory 
3. GPU中所有threads（包括不同SM的所有threads）共享的global memory

下面复习一下，做两个quiz。


Quiz -1 : 

Ans：选择A,B,D 
解读：根据定义，一个block只能run在一个SM；SM中不同blocks的threads不能cooperate


Quiz - 2 : 

Ans: 都不选~~~ 
解读：block执行时间及顺序不可控；block分配到哪个SM是GPU做的事情，并非programmer能指定的；

2.2 Memory in Program

How to write Efficient Programs from high level
  1. maximize arithmetic intensity 
    arithmetic intensity = calculation/memory 
    即要maximize calculation per thread 并 minimize memory per thread(其实目的是minimize memory access的时间) 
    方法：经常访问的数据放在可快速访问的memory（GPU中不同memory在硬件层的介绍参考第二章），对于刚才讲的local, shared and global memory的访问速度, 有 
    local > shared >> global >> CPU memory 
    所以，比如我想经常访问一个global memory，那可以在kernel中先将该global memory variable赋值给一个shared memory variable, 然后频繁访问那个shared memory variable.
  2. minimize memory access stride 
    如coalesce memory access图所示: 

    如果GPU的threads访问相邻memory，我们称为coalesced，如果threads间访问memory有固定步长(蹦着走)，我们称stripped，完全没规律的memory访问称为random。访问速度，有 
    coalesced > strided > random

  3. avoid thread divergence 
    这个我们在前两讲中有过相应说明。

Exercise:

给下面这段代码中5,6,7,8行的几句话执行速度排序（1最快，4最慢）：

1   __global__ void f(float* x, float* y, float* z){
2       float s,t,u;
3       __shared__ float a,b,c;
4       ...
5       s = *x;
6       t = s;
7       a = b;
8       *y = *z;
9   }

速度还是本地的T、S最快啦！！！！！！

但是本地内存是share还是global呢？

*y和*z才是global

 

而且 a=b是share memory既然是比本地慢的为什么还要用share来传递数组让速度变快呢？？不懂啊

懂了 核函数里传入参数（带指针的部分）是global内存里的 自己设的shared 参数是share内存里的

然后核函数里不带开头符号的参数是local 内存里的 也就是每一条线程里的东西

如果我不用share初始化一个数组，直接初始化一个数组 那么就表示每个线程里都有一个数组了 麻辣鸡 而我不需要这么做啊~

Ans: 5,6,7,8行执行速度为：3,1,2,4。

下面一节我们来看具体programming问题中的流程控制与同步。

(三). Control flow and synchronisation

3.1 program 运行顺序

在讲流程控制之前我们首先看一个例子，用来测试不同block的运行顺序。

Demo code:

#include <stdio.h>
#define Num_block 16
#define Num_thread  1

__global__ void print(){
     printf(“Num: %d\n”,blockIdx.x);
}

int main(){
    //launch the kernel
    print<<<Num_block, Num_thread>>>();
    cudaDeviceSynchronize();// what is the function of this sentence? - force the printf()s to flush， 不然运行时显示不出来
    return 0;
}

编译命令： 
nvcc -arch=sm_21 -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc print.cu

运行两次结果：

可见程序执行每一次的结果都不同，也就是不同block之间的执行顺序是不可控的，正如刚才quiz的ans。那么如果我们希望同步各个threads呢？

3.2 同步机制

第二章中我们在一个例子中引入并使用了同步函数syncthreads(), 即设置一个barrier，使所有threads运行到同步函数的时候stop and wait, 直到所有threads运行到此处，那么问题来了。


Exercise: 
考虑一个程序，将每个位置i的元素移到i-1的位置，需要多少个syncthreads()？ 
e.g kernel中声明如下：

…
int idx = threadIdx.x;
__shared__ int array[128];
array[idx] = idx;
if (idx<127){
     array[idx + 1] = array[idx];
}
…

Ans: 3个~

…
int idx = threadIdx.x;
__shared__ int array[128];
array[idx] = idx;
__syncthreads(); //如果不加将导致array还没赋值就被操作
if (idx<127){
     int tmp = array[idx];
     __syncthreads();//如不加导致先读后写，数据相关
     array[idx] = tmp;
     __syncthreads(); //如不加不能确保下面的程序访问到正确数据
}
…
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13

Quiz: 看下面这个程序会不会出现collision，哪里会出现collision?

1__global__ void f(){
2    __shared__ int s[1024];
3    int i = threadIdx.x;
4    __syncthreads();
5    s[i] = s[i-1];
6    __syncthreads();
7    if(i%2)   s[i] = s[i-1];
8    __syncthreads();
9    s[i] = (s[i-1]+s[i+1])/2;
10    printf(“%d\n”,s[i]);
11 }

Ans: Collision在 
1. 第5行，如上题，应为int tmp = s[i-1]; __syncthread(); s[i] = tmp; 
2. 第9行，同理 
PS： 第7行是没问题的，模拟一下就知道

3.3 Atomic Memory Operation

这一节中我们将要接触到原子操作。 
首先考虑一个问题：用1000000个threads给一个长为10个元素的array做加法，希望每个thread加100000，这个代码大家先写写看，很简单，依照我们之前的方法有下面的code:

注：这里的gputimer.h请去我的资源页面自行下载。

#include <stdio.h>
#include "gputimer.h"
using namespace Gadgetron;

#define NUM_THREADS 1000000
#define ARRAY_SIZE  10
#define BLOCK_WIDTH 1000
void print_array(int *array, int size)
{
    printf("{ ");
    for (int i = 0; i<size; i++)  { printf("%d ", array[i]); }
    printf("}\n");
}
__global__ void increment_naive(int *g)
{
     // which thread is this?
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     // each thread to increment consecutive elements, wrapping at ARRAY_SIZE
     i = i % ARRAY_SIZE;
     g[i] = g[i] + 1;
}
__global__ void increment_atomic(int *g)
{
     // which thread is this?
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     // each thread to increment consecutive elements, wrapping at ARRAY_SIZE
     i = i % ARRAY_SIZE;
     atomicAdd(&g[i], 1);
}
int main(int argc,char **argv)
{
    GPUTimer timer;
    printf("%d total threads in %d blocks writing into %d array elements\n",
           NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, ARRAY_SIZE);
    // declare and allocate host memory
    int h_array[ARRAY_SIZE];
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);

    // declare, allocate, and zero out GPU memory
    int * d_array;
    cudaMalloc((void **) &d_array, ARRAY_BYTES);
    cudaMemset((void *) d_array, 0, ARRAY_BYTES);
    // launch the kernel - comment out one of these
    timer.start();
    //increment_atomic<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
    increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
    timer.stop();

    // copy back the array of sums from GPU and print
    cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);
    print_array(h_array, ARRAY_SIZE);

    // free GPU memory allocation and exit
    cudaFree(d_array);
    return 0;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19
  • 20
  • 21
  • 22
  • 23
  • 24
  • 25
  • 26
  • 27
  • 28
  • 29
  • 30
  • 31
  • 32
  • 33
  • 34
  • 35
  • 36
  • 37
  • 38
  • 39
  • 40
  • 41
  • 42
  • 43
  • 44
  • 45
  • 46
  • 47
  • 48
  • 49
  • 50
  • 51
  • 52
  • 53
  • 54
  • 55
  • 56
  • 57
  • 58
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19
  • 20
  • 21
  • 22
  • 23
  • 24
  • 25
  • 26
  • 27
  • 28
  • 29
  • 30
  • 31
  • 32
  • 33
  • 34
  • 35
  • 36
  • 37
  • 38
  • 39
  • 40
  • 41
  • 42
  • 43
  • 44
  • 45
  • 46
  • 47
  • 48
  • 49
  • 50
  • 51
  • 52
  • 53
  • 54
  • 55
  • 56
  • 57
  • 58

执行两次的结果： 

可见结果里每个元素都是648/647，不符合预期100000。这是为什么呢？

看我们的kernel部分代码，每次执行g[i] = g[i] + 1, 一个read-modify-write操作，这样会导致许多线程读到g[i]的value，然后慢的线程将快的线程写结果覆盖掉了。如何解决呢？我们引入原子操作（atomic operation）, 更改上面的kernel部分为：

__global__ void increment_atomic(int *g)
{
     // which thread is this?
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     // each thread to increment consecutive elements, wrapping at ARRAY_SIZE
     i = i % ARRAY_SIZE;
     atomicAdd(&g[i], 1);
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8

我们可以得到结果： 

可见，结果正确。那么原子操作atomicAdd用了怎样的机制呢？——原子操作用了GPU built-in的特殊硬件，用以保证原子操作（同一时刻只能有一个thread做read-modify-write操作）


这里来看一下原子操作的limitations: 
1. only certain operations, data type（功能有限） 
2. still no ordering constraints（还是无序执行） 
3. serializes access to memory（所以慢）

原子操作的限制：

1.只有有限操作是允许的 原子加原子减原子异或原子交换 且大部分数据类型为整型 原子加、交换支持浮点运算

但是通过交换和加可以实现任何运算

2浮点运算的运算先后顺序是会影响结果的（a+b）+c !=a+(b+c)

3.后台并没有发生神奇的事情 线程操作序列的顺序会影响计算速度

用X个进程对Y个元素进行修改成功与否 消耗时间排序

第四个情况比第二个情况快一点点的原因是10的6次方个元素量太大以至于会被存到公共内存中，虽然他们都操作了1000000此线程

第三种情况最快是因为没有使用原子操作且操作元素只有100个 放在了快存里，反正发生冲突他也不管的虽然和第一种情况都操作了100000次

 

 

 

避免线程发散

线程发散不好的原因，蓝色的循环一遍后就干坐着等别的循环做完一起抵达barrier再行动 闲等的时间资源被浪费，而最后所用时间就是循环最久的线程用的时间

(四). 总结

本节课介绍了以下内容：

  • communication patterns

    • map
    • gather
    • scatter
    • stencil
    • transpose
  • gpu hardware & programming model
    • SMs, threads, blocks ordering
    • synchronization
    • Memory model - local, global, shared memory
  • efficient GPU programming
    • coalesced memory access
    • faster memory for common used variable

OK~ 第三课就结束了，过两天我把exercise上上来~ 敬请关注~.~

时间： 2024-11-12 00:51:15

优达学城-并行编程-Unit2 通信模块、同步机制、原子操作的相关文章

优达学城-并行编程-Unit2 硬件内存

GPU负责给SM分配wrap,SM以并行方式运行程序 在一个SM上跑的所有线程可能合作解决一个子问题(错的,不一定的) 一个单Kernel程序在多个wrap上运行,包含X线程块和Y线程块,可以确定x y先后跑的顺序或是在哪个SM上跑吗? 答:伐晓得(这是cuda的小秘密= =||) GPU的优越性: 1.快速切换SM运行,无法知其间通信 2.可扩展性强,GPU越大,任务分散越广 CUDA存储器类型: 每个线程拥有自己的register and loacal memory; 每个线程块拥有一块sh

优达学城机器学习工程师纳米学位项目介绍

本文由 meelo 原创,请务必以链接形式注明 本文地址,简书同步更新地址 一对一的项目辅导是优达学城(udacity)纳米学位的一大亮点.本文将简要介绍优达学城机器学习工程师纳米学位的6个项目.项目覆盖了机器学习的主要领域,包括回归分析.分类.聚类.增强学习及深度学习. 项目 0: 预测泰坦尼克号乘客生还率 这个项目需要你手动地实现一个简单的机器学习模型——决策树.1912年泰坦尼克号在第一次航行中与冰山碰撞沉没,泰坦尼克号乘客生还数据集记录了之中891个乘客的性别.年龄.社会阶级.配偶数量等

优达学城自动驾驶课程项目——车道检测

汽车在道路上行驶需要遵循一定的行驶规则,路面的车道则起到规范汽车行驶规则的作用.车道的种类有很多种,如单行线.双行线,虚线.网格线等,不同颜色.形状的车道线代表着不同的行驶规则,汽车和行人可以根据这些规则来使用道路,避免冲突.因此,准确检测并识别车道类型,并按照相应规则正确行驶,是汽车实现自动驾驶的基础. 优达学城的自动驾驶项目课程包含了一个车道线检测项目,其主要目的就是教给无人车如何检测并识别车道,本文档将该项目内容进行总结整理. 车道线检测方法主要分为两类:(1)基于道路特征的车道线检测:(

优达学城数据分析师纳米学位——第一课总结

从1月13号信誓旦旦的付款了第一位的纳米学位到今天已经一周多的时间了,可以发现自己在完成任务的时候更多的在乎的是不是时间上达到了要求,而没有过多的关注于实质的内容.有时候看到课程的小节数很多就有一种畏惧感和烦躁的心情,逐渐的说服自己取放弃,这其实是一种观念上的偏差.可能是因为自己的性格比较急躁,很多事情都想要迅速的完成并且还能够保证质量,但是实际来说,只要方法得当,最后的成果总会和你投入的时间成正比,放平心态,一步一步慢慢来. 下面对第一课所学的内容做一个梳理和总结. 研究方法 数据可视化 集中

优达学城数据分析师纳米学位——P3项目知识点整理及代码分析

P3 OpenStreetMap 项目思路整理 P3项目的核心在于数据的整理 data cleaning 数据清洗,数据来源于开源的OpenStreetMap平台,该平台上的很多数据都是开发者自行输入的,难免会造成数据的混乱和缺失,错误,也就是所说的dirty data 脏数据 human involved data cleaning的关注点 validity completeness accuracy consistency uniformity validity 有效性 数据是否符合常理 人

Udacity(优达学城) 300块红包优惠券

纳米学位:来自硅谷的名企官方课程 7天免费试用结束后,在"我的教室->设置->纳米学位->续费"页面上的优惠码区域,输入AF55BA53,立即减300元:

优达学城数据分析师纳米学位——第二课 jupyter notebook的使用

shift+Enter Control+Enter shift+tab 查看代码功能 tab显示相关包中的命令名称 ESC转换为命令操作 s 保存快捷键 命令行 小键盘图片 保存类型 Markdown html python文件 Markdown 单元格 使得代码更加易读  LaTeX free software可以编辑数学公式和符号 '配置Jupyter notebook' 6/11-markdown 单元格

C#并行编程中的Parallel.Invoke

一.基础知识 并行编程:并行编程是指软件开发的代码,它能在同一时间执行多个计算任务,提高执行效率和性能一种编程方式,属于多线程编程范畴.所以我们在设计过程中一般会将很多任务划分成若干个互相独立子任务,这些任务不考虑互相的依赖和顺序.这样我们就可以使用很好的使用并行编程.但是我们都知道多核处理器的并行设计使用共享内存,如果没有考虑并发问题,就会有很多异常和达不到我们预期的效果.不过还好NET Framework4.0引入了Task Parallel Library(TPL)实现了基于任务设计而不用

第三篇：GPU 并行编程的运算架构

前言 GPU 是如何实现并行的?它实现的方式较之 CPU 的多线程又有什么分别? 本文将做一个较为细致的分析. GPU 并行计算架构 GPU 并行编程的核心在于线程,一个线程就是程序中的一个单一指令流,一个个线程组合在一起就构成了并行计算网格,成为了并行的程序,下图展示了多核 CPU 与 GPU 的计算网格: 二者的区别将在后面探讨. 下图展示了一个更为细致的 GPU 并行计算架构: 该图表示,计算网格由多个流处理器构成,每个流处理器又包含 n 多块. 下面进一步对 GPU 计算网格中的一些概念