Nvidia CUDA初级教程6 CUDA编程一

视频：https://www.bilibili.com/video/BV1kx411m7Fk?p=7
讲师：周斌

GPU架构概览

GPU特别使用于：
- 密集计算，高度可并行计算
- 图形学
晶体管主要被用于：
- 执行计算
- 而不是
  - 缓存数据
  - 控制指令流

图中分别是CPU、GPU各个部件所占的芯片面积。可以看到，CPU芯片中大量部分是缓存和控制逻辑，而GPU中则绝大部分都是计算单元。

CUDA编程相关简介

CUDA的一些信息

层次化线程集合
共享存储
同步

CUDA术语

主机端和设备端

HOST - 主机端，通常指CPU
- 采用ANSI标准C语言编程
Device - 设备端，通常指GPU（数据可并行）
- 采用ANSI标准C的扩展语言编程（CUDA C）
HOST 和 Device 拥有各自的存储器
CUDA编程
- 包括主机端和设备端两部分代码

核

Kernel 数据并行处理函数
- 通过调用 Kernel 函数在设备端创建轻量级的线程，线程由硬件负责创建并调度

类似于 OpenCL 的 shader？

核函数会在 N 个不同的 CUDA 线程上并行执行

// 定义核函数
__global__ void VecAdd(float* a, float* B, float* C) {int i = threadIdx.x;C[i] = A[i] + B[i];
}int main() {// ...// 在N个线程上调用核函数VecAdd<<<1, N>>>(A, B, C);
}

CUDA程序的执行

CUDA程序执行的流程大体上是这样的：当我们在CPU端的代码是串行执行的（这里简单地认为指令在CPU上串行执行），当遇到需要并行大量处理数据时，会调用核函数在GPU上进行计算，计算完成后将结果返回给CPU。

线程层次

Grid - 一维或多维线程块（block）
- 一维或二维
Block - 一维线程
- 一维，二维或三维
- 一个 Grid 中的每个 Block 的线程数是一样的
- Block 内部的每个线程可以：
  - 同步 Synchronize
  - 访问共享存储器 shared memory

一个线程可以类比为一个员工，一个 block 是一个科室，grid 是整个公司。

线程ID

每一个线程都有一个索引：threadIdx

一维 Block Thread ID == Thread Index
二维 Block (Dx, Dy)
- 索引为 (x, y) 的 Thread ID == x + yDy
三维 Block (Dx, Dy, Dz)
- 索引为 (x, y) 的 Thread ID == x + yDy + zDxDy

代码实例

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = threadIdx.x;int j = threadIdx.y;C[i][j] = A[i][j] + B[i][j];
}int main() {int numBlocks = 1;dim3 threadsPerBlock(N, N);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

每个 Block 中的线程的索引是二维的，这在我们处理二维数据（矩阵）时可以很方便地进行对应。

线程数

Thread Block 线程块

线程的的集合
- G80 和 GT200：多达512个线程
- Fermi：多达1024个线程
位于相同的处理器核（相同的SM）
共享所在核的存储器

块索引

块索引：blockIdx
维度：blockDim
- 一维，二维或三维

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y * blockDim.y + threadIdx.y;if (i < N && j < N)C[i][j] = A[i][j] + B[i][j];
}int main() {// ...dim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

例如 N = 32

每个块有16x16个线程（跟N无关）
- threadIdx([0, 15], [0, 15])
Grid 里面有 2x2 个线程块 block
- blockIdx([0, 1], [0, 1])
- blockDim = 16

int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;

i = [0, 1] * 16 + [0, 15]

线程块之间

线程块彼此之间独立执行

任意顺序：并行或串行
被任意数量的处理器以任意顺序调度
- 处理器的数量具有可扩展性

一个块内部的线程

一个块内部的线程有一些重要的特性：

共享容量有限的低延迟存储器 (shared memory)
同步执行
- 合并访存
- __syncThreads()
  - barrier - 块内线程一起等待所有的线程都
  - 轻量级线程

CUDA内存传输

主机端与设备端

CUDA内存传输

device 端代码可以：
- 读写该线程的 registers
- 读写该线程的local memory
- 读写该线程所属的块的 shared memory
- 读写grid的 global memory
- 只读grid的 constant memory
host 端代码可以：
- 读写grid的 global memory 和 constant memory
host 可以从 device 往返传输数据
- global memory 全局存储器
- constant memory 常量存储器

CUDA内存传输函数

在设备端分配 global memory：cudaMalloc()
释放存储空间 cudaFree()

float* Md;
int size = Width * Width * sizof(float);
cudaMalloc((void**)&Md, size);
//...
cudaFree(Md);

注意这里的指针 Md 是指向 device（GPU）上的存储空间。

内存传输：cudaMemcpy()
- host to host
- host to device
- device to host
- device to device

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

示例：矩阵相乘 Matrix Multiply

矩阵相乘简介

向量
点乘
行优先或列优先
每次点乘结果输出一个元素

矩阵相乘CPU实现

void MatrixMulOnHost(float* M, float* N, float* P, int width) {for (int i=0; i<width; ++i) {for (int j=0; j<width; ++j) {float sum = 0;for (int k=0; k<width; ++k) {float a = M[i * width + k];float b = N[k * width + j]:sum += a  * b;}P[i * width + j] = sum;}}
}

CUDA算法框架

三步走：

int main(void) {// 1 分配device空间// 2 在GPU上，并行计算MatrixMulOnDevice(M, N, P, width);// 3 将结果拷贝回CPU，并释放device空间return 0;
}

伪代码如下：

void MatrixMulOnDevice(float* M, float* N, float* P, int Width) {int size = Width * Width * sizeof(float);// 1cudaMalloc(Md, size);cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMalloc(Nd, size);cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);cudaMalloc(Pd, size);// 2 调用cuda核函数，并行计算// 3cudaMemcpy(P. Pd, size, cudaMemcpyDeivceToHost)cudaFree(Md); cudaFree(Nd); cudaFree(Pd);
}

CUDA C 实现

矩阵相乘样例

目前版本矩阵相乘的问题

在上述算法实现中最主要的性能问题是什么？
- 访存
主要的限制是什么？
- 访存带宽

Nvidia CUDA初级教程6 CUDA编程一相关推荐

Nvidia CUDA初级教程4 GPU体系架构概述
Nvidia CUDA初级教程4 GPU体系架构概述视频:https://www.bilibili.com/video/BV1kx411m7Fk?p=5 讲师:周斌本节内容: 为什么需要GPU 三 ...
Nvidia CUDA初级教程2 并行程序设计概述
Nvidia CUDA初级教程2 并行程序设计概述视频:https://www.bilibili.com/video/BV1kx411m7Fk?p=3 讲师:周斌本节内容: 为什么需要? 怎么做? ...
NVIDIA CUDA初级教程（P11）CUDA程序基本优化
文章目录 1.Parallel Reduction并行规约 2.Warp分割 1.Parallel Reduction并行规约最优性能=有效的数据并行算法+针对GPU架构特性的优化 eg:Paral ...
NVIDIA CUDA初级教程（P2-P3）CPU体系架构概述、并行程序设计概述
文章目录 1.CPU体系架构概述 2.并行程序设计概述 1.CPU体系架构概述现代CPU架构 CPU的定义 (1)执行指令.处理数据的器件:完成基本的逻辑和算术指令 (2)内存接口.外部设备接口 ( ...
Cuda 学习教程：Cuda 程序初始化
Cuda程序初始化目前,cuda里面没有对设备的初始化函数InitDevice(),只能每次调用的api函数的时候,加载设备的上下文,自动进行初始化,这将带来问题: First函数调用的时候,需要自 ...
Ubuntu + nvidia驱动+ cuda安装教程以及重装问题
操作系统:ubuntu 18 nvidia驱动版本: 515 cuda: 11.7 在此方面掉过坑,写此篇文章mark一下.如果遇到同样问题的小伙伴可以参考一下. 第一次在ubuntu上尝试安装nvi ...
最简单、实用的cuda安装教程！！！（nvidia官方渠道下载）
网上教程一大推,讲了一大堆,也没解释原理,实用的没几个,自己总结的,比较简单 note:无需卸载原机器驱动,无需禁用nouveau驱动 attention:以下内容为有sudo权限安装教程,没有sud ...
NVIDIA GeForce GTX 850M win10 cuda配置及安装教程
NVIDIA GeForce GTX 850M win10 cuda配置 cuda10.0+cudnn7.6.4+tensorflow2.0.0+keras2.3.0 ---------------- ...
最全与最好的——CUDA入门教程
开篇一张图,后面听我编 1. 知识准备 1.1 中央处理器(CPU) 中央处理器(CPU,Central Processing Unit)是一块超大规模的集成电路,是一台计算机的运算核心(Core)和 ...

Nvidia CUDA初级教程6 CUDA编程一

Nvidia CUDA初级教程6 CUDA编程一

GPU架构概览

CUDA编程相关简介

CUDA的一些信息

CUDA术语

主机端和设备端

核

CUDA程序的执行

线程层次

线程ID

代码实例

线程数

块索引

线程块之间

一个块内部的线程

CUDA内存传输

主机端与设备端

CUDA内存传输函数

示例：矩阵相乘 Matrix Multiply

矩阵相乘简介

矩阵相乘CPU实现

CUDA算法框架

CUDA C 实现

矩阵相乘样例

目前版本矩阵相乘的问题

Nvidia CUDA初级教程6 CUDA编程一相关推荐

最新文章

热门文章