Heterogeneous Parallel Programming 异构并行编程 - UIUC伊利诺伊大学(持续更新)

news2024/12/23 23:24:45

  • Lecture 1
    • 1.2 Introduction to Heterogeneous异构
    • 1.3 Portability and Scalability
    • 1.4 Introduction to CUDA 数据并行化 and 执行模型
    • 1.5 Introduction to CUDA 内存模型 and 基本函数API
    • 1.6 Introduction to CUDA Kernel-based SPMD
    • 1.7 更高维的Grid的Kernel-based SPMD例子
    • 1.8 Matrix Multiplication矩阵乘法
  • Lecture 2
    • 2.1 Kernel-Based Parallel Programming

Heterogeneous Programming是采用不同类型的计算节点协同进行计算。而Heterogeneous Parallel Programming则是建立在这种机制上的并行计算。这里使用的是的CUDA。CUDA是NVIDIA推出的建立在C语言和GPU基础上的计算框架。

Lecture 1

1.2 Introduction to Heterogeneous异构

请添加图片描述
GPU与CPU的设计理念不同:GPU旨在提供高吞吐量throughput,而CPU旨在提供低延迟latency,CPU需要降低指令的执行时间,所以有很大的缓存,而GPU则不然。单一的GPU线程执行时间相当长,因此总是多线程并行,这样提高了吞吐量。所以在串行逻辑计算部分应该使用CPU,而并行数据计算部分则应使用GPU

1.3 Portability and Scalability

降低软件开发成本的关键:
可扩展性Scalability:细粒度编程风格,只需调整参数,便可扩展到更多相同的硬件上运行。
请添加图片描述
可移植性Portability: 可扩展到更多不同的硬件上运行
请添加图片描述

1.4 Introduction to CUDA 数据并行化 and 执行模型

数据并行化Data Parallelism:将向量或矩阵对应位置的数据分别交给不同的线程并行处理。
在这里插入图片描述CUDA/OpenCL执行模型:分为两部分,Host 和 Device,串行部分在Host(CPU)上执行,而并行部分在Device(GPU)上执行。相比传统的C语言,CUDA增加了一些扩展,包括了库和关键字。CUDA代码提交给NVCC编译器,该编译器将代码分为Host代码和Device代码两部分。Host代码即为原本的C语言,交由GCC或其他的编译器处理;Device代码部分交给一个称为实时(Just in time)编译器的组件,在给代码运行之前编译。

CUDA架构执行程序时在CPU和GPU上来回切换,CPU程序穿行执行,GPU程序使用核函数Kernel,每个核函数被大量(一组grid)CUDA线程并行执行,Kernel可以使用<<< >>>接受配置参数,也可以使用( )接受调用参数。
在这里插入图片描述CUDA线程: 每个CUDA线程实际上都是一个虚拟的处理器,去执行核函数,CUDA kernel函数被一网格Gird的线程并行执行,每个网格里有很多块Block,每个块里有很多线程Thread。

如下为单Block线程执行核函数:
在Block中,每个Thread都去执行核函数,去承担运算中的不同部分,但每个Thread都有一个唯一的threadidx
在这里插入图片描述
如下为多Block线程执行核函数:
在Grid中,每个Block都去执行核函数,去承担运算中的不同部分,但每个Block都有一个唯一的blockIdx
在这里插入图片描述
并行线程阵列Grid——Block——Thread三级结构组成,每个Thread有唯一定位ID
其中DimGrid是Grid的形状和DimBlock三Block的形状 ,它们都是调用核函数的配置参数,用<<<DimGrid,DimBlock>>>设置。

gridDim.x/y/z为每个Grid在x/y/z维度上的大小,blockDim.x/y/z为每个Block在x/y/z维度上的大小。blockIdx标识Grid中的Block,threadIdx标识Block中的Thread。

CUDA Thread 定位公式
1.先找到当前线程位于Grid中的哪一个Block线程块的blockId:

blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

2.找到当前线程位于Block中的哪一个线程threadId:

threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

3.计算一个Block中一共有多少个线程M:

M = blockDim.x*blockDim.y*blockDim.z

4.求得当前的线程序列号idx:

idx = M*blockId +  threadId ;

blockIdx 和 threadIdx 和 gridDim 和 blockDim都是Dim3类型(x,y,z)的变量每个xyz默认值为1,因此它们可以是1D、2D、3D的。因此每个i也可以是1D、2D、3D的。
在这里插入图片描述
同一Block的Thread通过shared memory 和 atomic operations 和 barrier synchronization进行合作,而不同Block的Thread无法合作。

1.5 Introduction to CUDA 内存模型 and 基本函数API

CUDA程序执行流程: 使用host的数据在device上进行运算,再送回host:

#include <cuda.h>
void vecAdd(float* A, float* B, float* C, int n)
{
   int size = n* sizeof(float); 
   float* A_d, B_d, C_d;1. // Allocate device memory for A, B, and C
   // copy A and B to device memory 
    
2. // Kernel launch code – to have the device
   // to perform the actual vector addition

3. // copy C from the device memory to host memory
   // Free device vectors
}

CUDA 内存模型:每个Grid有一个共享的Global Memory,用于Host和Device交互,其中每个线程有自己私有的寄存器。Host代码负责分配Grid中的共享内存空间,以及数据在Host、Device之间的传输。Device代码则只与共享内存、本地寄存器交互。
在这里插入图片描述

CUDA 内存分配 和 数据传输 API:与C语言中的函数想对应,CUDA有以下几个基本函数:
cudaMalloc()、cudaFree()、cudaMemcpy(),其作用等同于C中的对应函数,不同之处在于这些函数操作的是Device中的共享内存。

  • cudaMalloc(Pointer , Size of bytes ): 给一个Pointer指针 在device上分配global memory。
  • cudaFree(Pointer ):释放这个Pointer指针的global memory
  • cudaMemcpy(目的Pointer, 源Pointer, 要Copy的bytes数, 传输的类型/方向): 则用于Host内存与Device内存异步传输数据。

CUDA程序执行流程: 使用host的数据在device上进行运算,再送回host:

//host函数
void vecAdd(float* A, float* B, float* C, int n)//传入host指针,向量长度为n
{
   int size = n * sizeof(float); 
    float* A_d, B_d, C_d;//定义device指针

1. // Transfer A and B to device memory 
    cudaMalloc((void **) &A_d, size);//申请全局device memory 
    cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);//转移host数据到device
    cudaMalloc((void **) &B_d, size);//申请全局device memory 
    cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);//转移host数据到device

   // Allocate device memory for C
    cudaMalloc((void **) &C_d, size);//申请全局device memory 

2. // Kernel invocation code – to be shown later3. // Transfer C from device to host
    cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost); //转移device数据到host
   // Free device memory for A, B, C
    cudaFree(A_d); cudaFree(B_d); cudaFree (C_d);//释放device memory 
}

1.6 Introduction to CUDA Kernel-based SPMD

SPMD:即Single Program Multiple Data,指相同的程序处理不同的数据。在Device端执行的Threads即属于此类型,每个Grid中的所有线程执行相同的Kernel函数(共享PC和IR指针),但是这些Threads需要从共享的Global Memory中取得自身的数据,这样就需要一种数据定位机制。

CUDA的函数分类
在这里插入图片描述
注意都是双下划线。其中的__global__函数即为C代码中调用Device上计算的入口,必须返回void类型__host__ 函数为传统的C函数,也是默认类型。之所以增加这一标识的原因是有时候可能__device__和__host__共同使用,这时可以让编译器知道,需要编译两个版本的函数。而__device__函数只能由__global__函数或__device__函数调用。

如计算向量加法:

// Compute vector sum C = A+B
// Each thread performs one pair-wise addition

//device code
__global__ void vecAddkernel(float* A_d, float* B_d, float* C_d, int n)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;//device函数要计算thread的索引,以便分配各自的数据
    if(i<n) C_d[i] = A_d[i] + B_d[i];//if做boundry check
}

// host code 
__host__ int vecAdd(float* A, float* B, float* C, int n)
{
 //此处省略了 A_d, B_d, C_d allocations and copies
 // Run ceil(n/256) blocks of 256 threads each
  dim3 DimGrid(n/256, 1, 1);
  if (n%256) DimGrid.x++;
  dim3 DimBlock(256, 1, 1);

  vecAddKernel<<<DimGrid,DimBlock>>>(A_d, B_d, C_d, n);//host函数中调用device函数执行
}

Any call to a kernel function is asynchronous from CUDA 1.0 on, explicit synch needed for blocking

CUDA程序工作流程:host函数 -> device函数(__global __) -> host函数
在这里插入图片描述

CUDA程序异构编译过程:编译器nvcc识别Host函数和Device函数,分别让它们与各自的编译器进一步编译,如gcc和intel icc,最后在异构计算平台上运行。
在这里插入图片描述

1.7 更高维的Grid的Kernel-based SPMD例子

前面我们讨论了向量的加法,这是一维数据的计算,下面我们讲讲二维数据和内存模型。

  • 一维数据:向量。
  • 二维数据:矩阵/图片。对于二维数据,C/C++是为主,Fortran以为主。
__global__ void PictureKernel(float* Pin, float* Pout, int n, int m) 
{

  // Calculate the row # of the Pin and Pout element to process
  int Row = blockIdx.y*blockDim.y + threadIdx.y;
  
  // Calculate the column # of the Pin and Pout element to process
  int Col = blockIdx.x*blockDim.x + threadIdx.x;

  // each thread computes one element of Pout if in range
  if ((Row < m) && (Col < n)) {
    Pout[Row*n+Col] = 2*Pin[Row*n+Col];
  }
  
}

假定需要处理一张76x62像素的图片,采用16x16的Block,如图:
在这里插入图片描述为什么要边界查询,判断(Row < m) && (Col < n):
因为Threads数量>矩阵单元数,并不是所有的Thread都要参与运算。其中1区域的Block每一个Thread都对应有像素,而2、3、4则没有。

1.8 Matrix Multiplication矩阵乘法

矩阵乘法规则:前行成后列,在实际使用中,我们常常用行主导的一维数组代替二维数组,即,twoDim[2, 3] = oneDim[2 * numCol + 3]。
在这里插入图片描述
CPU的Host代码: 两个循环遍历整个将结果矩阵

// Matrix multiplication on the (CPU) host in double precision
void MatrixMulOnHost(float* M, float* N, float* P, int Width)
{   
    for (int i = 0; i < Width; ++i)
        for (int j = 0; j < Width; ++j) {
            double sum = 0;
            for (int k = 0; k < Width; ++k) {
                double a = M[i * Width + k];
                double b = N[k * Width + j];
                sum += a * b;
            }
            P[i * Width + j] = sum;
        }
}

在这里插入图片描述

GPU的Device代码: 每个Thread有Row和Col两个量来索引

__global__ void MatrixMulKernel(float* M, float* N, float* P, int Width)
{
 // Calculate the row index of the P element and M
 int Row = blockIdx.y * blockDim.y + threadIdx.y;
 // Calculate the column index of P and N
 int Col = blockIdx.x * blockDim.x + threadIdx.x;

 if ((Row < Width) && (Col < Width)) {
	float Pvalue = 0;
    
    // each thread computes one element of the block sub-matrix
	for (int k = 0; k < Width; ++k)
      Pvalue += M[Row*Width+k] * N[k*Width+Col];
   
    P[Row*Width+Col] = Pvalue;
  }
}

Host代码调用核函数:

    // Setup the execution configuration
  // TILE_WIDTH is a #define constant
     dim3 dimGrid(Width/TILE_WIDTH, Width/TILE_WIDTH, 1);
     dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);


  // Launch the device computation threads!
     MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);

Lecture 2

2.1 Kernel-Based Parallel Programming

线程调度

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.coloradmin.cn/o/32273.html

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈,一经查实,立即删除!

相关文章

Linux的基本协议与他的堂兄堂弟

14天学习训练营导师课程&#xff1a; 互联网老辛《 符合学习规律的超详细linux实战快速入门》 努力是为了不平庸~ 学习有些时候是枯燥的&#xff0c;但收获的快乐是加倍的&#xff0c;欢迎记录下你的那些努力时刻&#xff08;学习知识点/题解/项目实操/遇到的bug/等等&#xf…

教程一 在Go使用JavaScript、HTML和CSS构建Windows、Linux、MacOSX跨平台的桌面应用

Energy是Go语言使用JavaScript、HTML和CSS构建跨平台的桌面应用程序可用于构建跨平台的桌面应用内嵌 Chromium CEF 二进制 环境安装 Energy 命令行工具 使用命令行工具自动安装Energy框架的所有依赖(CEF)&#xff0c;支持Window、Linux、MacOSX 安装过程从网络下载CEF和Energy…

二、vue基础入门

一、vue简介 1.1、什么是vue 官方给出的概念&#xff1a;Vue (读音 /vjuː/&#xff0c;类似于 view) 是一套用于构建用户界面的前端框架。 1.2、vue的特性 vue框架的特性&#xff0c;主要体现在如下两方面&#xff1a; 数据驱动视图双向数据绑定 1.2.1、数据驱动视图 在…

高灵敏度艾美捷小鼠肿瘤坏死因子α-ELISpot试剂盒

肿瘤坏死因子-a&#xff08;TNF-a&#xff09;由许多不同的细胞类型产生&#xff0c;例如单核细胞&#xff0c;巨噬细胞&#xff0c;T细胞和B细胞。在TNF-a的许多作用中&#xff0c;有针对细菌感染&#xff0c;细胞生长调节&#xff0c;免疫系统调节和参与败血症性休克的保护。…

现代气象仪器 | 太阳辐射测量

南京信息工程大学 实验&#xff08;实习&#xff09;报告 实验&#xff08;实习&#xff09;名称 现代气象仪器 实验&#xff08;实习&#xff09;日期 10.28 得分 指导老师 学院 电信院 专业 电子信息工程 年级 2020 班次 4 姓名 学号 20208327 实验…

万字博客带你全面剖析Spring的依赖注入

1.写在前面 前面的博客我们已经写了Spring的依赖查找&#xff0c;这篇博客我们来了解写Spring的依赖注入。 2.依赖注入的模式和类型 手动模式 - 配置或者编程的方式&#xff0c; 提前安排注入规则 XML 资源配置元信息Java 注解配置元信息API 配置元信息 自动模式 - 实现方…

华为机试 - 最大括号深度

目录 题目描述 输入描述 输出描述 用例 题目解析 算法源码 题目描述 现有一字符串仅由 ‘(‘&#xff0c;’)’&#xff0c;{‘&#xff0c;’}’&#xff0c;[‘&#xff0c;’]’六种括号组成。 若字符串满足以下条件之一&#xff0c;则为无效字符串&#xff1a; ①…

【MySQL】拿来即用 —— MySQL中的数据类型

个人简介&#xff1a;Java领域新星创作者&#xff1b;阿里云技术博主、星级博主、专家博主&#xff1b;正在Java学习的路上摸爬滚打&#xff0c;记录学习的过程~ 个人主页&#xff1a;.29.的博客 学习社区&#xff1a;进去逛一逛~ MySQL数据类型⚪熟悉SQL一、MySQL数据类型总结…

设备树和设备树语法

设备树 驱动代码只负责处理驱动的逻辑&#xff0c;而关于设备的具体信息存放到设备树文件中。许多硬件设备信息可以直 接通过它传递给 Linux&#xff0c;而不需要在内核中堆积大量的冗余代码。 设备树&#xff0c;将这个词分开就是“设备”和“树”&#xff0c;描述设备树的文…

【计算机毕业设计】22.毕业设计选题系统ssm源码

一、系统截图&#xff08;需要演示视频可以私聊&#xff09; 引言 近年来&#xff0c;电子商务发展的愈趋成熟使得人们的消费方式以及消费观念发生巨大改变&#xff0c;网上竞拍的拍卖模式随之发展起来。大学拍卖网旨在为湘大学生提供一个线上拍卖的交易平台。平台展示的商品大…

【American English】美式发音,英语发音,美国音音标列表及发音

首先声明&#xff0c;网上各种英式发音和美式发音的教程&#xff0c;而我的目的是寻找美式发音。但是自己现在也是在不断地找寻中&#xff0c;所以资料找错了请莫怪。另外&#xff0c;资料顺序采用部分倒叙&#xff0c;不喜请勿吐槽。 文章目录发音示意图49. [](https://www.bi…

百度地图有感

以前总认为坚持会让我们变强大&#xff0c;但是长大后发现&#xff0c;让我们强大的&#xff0c;是放下。 生活也许就是这样&#xff0c;多一分经验便少一分幻想&#xff0c;以实际的愉快平衡现实的痛苦。 百度地图开放平台 百度地图入门指南 百度地图开发指南 百度地图API文…

性早熟和微生物群:性激素-肠道菌群轴的作用

谷禾健康 肠道菌群 & 性激素 青春期是生命的一个关键阶段&#xff0c;与性成熟相关的生理变化有关&#xff0c;是一个受多种内分泌和遗传控制调控的复杂过程。 青春期发育可以在适当的时候&#xff0c;早熟或延迟。 未经治疗的性早熟的孩子通常不会达到成年身高的全部潜力。…

Activity的最佳实践

文章目录Activity的最佳实践知晓当前是在哪一个Activiy随时随地退出程序启动Activity的最佳写法Activity的最佳实践 知晓当前是在哪一个Activiy 创建一个BaseActivity类,继承AppCompatActivity类.重写onCreate方法 open class BaseActivity : AppCompatActivity() {override…

xilinx PL测 DP 点屏 /接收(二)--RX

环境&#xff1a; a)硬件&#xff1a;官方ZCU106开发板 , tb-fmch-vfmc-dp子卡。 b)软件&#xff1a;vivado2021.1&#xff0c;vitis2021.1&#xff0c;裸机程序。 1、官方例程&#xff1a; 2、DP RX IP &#xff1a; 3、DP RX寄存器&#xff1a; 4、时钟&#xff1a; 5、像素&…

CentOS 6.6系统怎么安装?CentOS Linux系统安装配置图解教程

服务器相关设置如下&#xff1a; 操作系统&#xff1a;CentOS 6.6 64位 IP地址&#xff1a;192.168.21.129 网关&#xff1a;192.168.21.2 DNS&#xff1a;8.8.8.8 8.8.4.4 备注&#xff1a; CentOS 6.6系统镜像有32位和64位两个版本&#xff0c;并且还有专门针对服务器优化过的…

【端到端存储解决方案】Weka,让企业【文件存储】速度飞起来!

一、HK-Weka概述 虹科WekaIO&#xff08;简称HK-Weka&#xff09;是一个可共享、可扩展的文件存储系统解决方案&#xff0c;其并行文件系统WekaFS支持NVMeoF的flash-native并行文件系统、比传统的NAS存储及本地存储更快。 HK-Weka后端主机被配置为集群&#xff0c;它与安装在应…

在Mysql中新建序列Sequence

在Oracle数据库中想要一个连续的自增数据类型的值&#xff0c;可以通过创建一个sequence来实现。而在Mysql数据库中并没有sequence&#xff0c;如想要在Mysql中像Oracle那样使用序列&#xff0c;该如何操作呢&#xff1f;&#xff08;可以使用mysql中的自增主键&#xff09; 1、…

哪个牌子的led灯质量好?2022LED护眼台灯最好的品牌有哪些

谈及led灯的品牌&#xff0c;就不得不提一些比较专业的厂商了&#xff0c;特别是在护眼照明领域&#xff0c;明基、南卡、飞利浦、松下等品牌都有不俗的实力&#xff0c;出产的led护眼台灯在业内都有广泛的知名度&#xff0c;在消费者领域也是好评连连。那么它们到底好在哪儿呢…

蓝牙协议栈分层

一、分层 BLE协议栈主要用来对你的应用数据进行层层封包&#xff0c;以生成一个满足BLE协议的空中数据包&#xff0c;也就是说&#xff0c;把应用数据包裹在一系列的帧头&#xff08;header&#xff09;和帧尾&#xff08;tail&#xff09;中。 BLE协议栈主要由如下几部分组成…