CUDA编程学习

仅为自己学习的记录

Reference
CUDA编程基础入门系列（持续更新）

1.GPU硬件平台

1.1什么是GPU

GPU 意为图形处理器，也常被称为显卡，GPU最早主要是进行图形处理的。如今深度学习大火，GPU高效的并行计算能力充分被发掘，GPU在AI应用上大放异彩。GPU拥有更多的运算核心，其特别适合数据并行的计算密集型任务，如大型矩阵运算，与GPU对应的一个概念是CPU，但CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务，CPU更擅长数据缓存和流程控制

1.2GPU性能

主要GPU性能指标：

1.核心数量：为GPU提供计算能力的硬件单元，核心数量越多，可并行运算的线程越多，计算的峰值越高；

2.GPU显存容量：显存容量决定着显存临时存储数据的多少，大显存能减少读取数据的次数，降低延迟，可类比CPU的内存；

3.GPU计算峰值：每块显卡都会给出显卡的GPU计算峰值，这是一个理论值，代表GPU的最大计算能力，一般实际运行是达不到这个数值的；

4.显存带宽：GPU包含运算单元和显存，显存带宽就是运算单元和显存之间的通信速率，显存带宽越大，代表数据交换的速度越快，性能越高。

2.CPU+GPU异构架构

GPU不能单独进行工作，GPU相当于CPU的协处理器，由CPU进行调度。CPU+GPU组成异构计算架构，CPU的特点是更加擅长逻辑处理，而对大量数据的运算就不是那么擅长了，GPU恰好相反，GPU可以并行处理大量数据运算。

由这个图可以看出，一个典型的 CPU 拥 有少数几个快速的计算核心，而一个典型的 GPU 拥有几百到几千个不那么快速的计算核心。CPU 中有更多的晶体管用于数据缓存和流程控制，但 GPU 中有更多的晶体管用于算术逻辑单元。所以，GPU 是靠众多的计算核心来获得相对较高的计算性能的。

CPU和GPU都有自己的DRAM（dynamic random-access memory，动态随机存取内存），它们之间一般由PCIe总线（peripheral component interconnect express bus）连接。这里多说一点，PCIe总线上的数据传输速率相对来说是比较慢的，也就是说一个不那么复杂的任务，CPU和GPU数据传输的时间可能远远大于GPU计算的时间，所以在一些小型任务上，应用GPU也未必会起到加速作用。

在由 CPU 和 GPU 构成的异构计算平台中，通常将起控制作用的 CPU 称为主机（host），将起加速作用的 GPU 称为设备（device）。所以在今后，说到主机就是指CPU，说到设备就是指GPU。

3.CUDA介绍

3.1什么是CUDA

2006年11月，NVIDIA推出了CUDA，这是一种通用的并行计算平台和编程模型，利用NVIDIA GPU中的并行计算引擎，以比CPU更有效的方式解决许多复杂的计算问题。

CUDA是建立在NVIDIA的GPU上的一个通用并行计算平台和编程模型。

基于GPU的并行计算目前已经是深度学习训练的标配。

3.2CUDA编程语言

CUDA旨在支持各种语言和应用程序编程接口；
最初基于C语言，目前越来越多支持C++；
CUDA还支持Python编写；

3.3CUDA运行时API

CUDA提供两层API接口，CUDA驱动(driver)API和CUDA运行时(runtime)API；
两种API调用性能几乎无差异

4.CUDA的Hello World程序

4.1nvcc:

1.安装CUDA即可使用nvcc
2.nvcc支持纯C++代码的编译
3.编译扩展名为.cu的CUDA文件

编译CUDA文件指令：nvcc hello.cu -o hello

#include <stdio.h>

int main(void)
{
    printf("Hello World!\n");
    return 0;
}

5.核函数

核函数在GPU上进行并行执行
注意：
(1) 限定词__global__修饰
(2) 返回值必须是void
(3) 核函数只能访问GPU内存
(4) 核函数不能使用变长参数
(5) 核函数不能使用静态变量
(6) 核函数不能使用函数指针
(7) 核函数具有异步性
(8) 核函数不支持C++中的iostream
(9) 在调用时需要用 <<<grid, block>>> 来分配 block 数与线程数

CUDA程序编写流程

int main(void)
{
主机代码
核函数调用
主机代码
return 0;
}

核函数示例

#include <stdio.h>

__global__ void hello_from_gpu()
{
    printf("Hello World from the the GPU\n");
}


int main(void)
{
//调用核函数
    hello_from_gpu<<<1, 1>>>();
    //同步函数
    cudaDeviceSynchronize();
    return 0;
}

6. CUDA线程模型

6.1 线程模型结构

1、线程模型重要概念：
（1）grid 网格
（2）block 线程块
2、线程分块是逻辑上的划分，物理上线程不分块
3、配置线程：<<<grid_size, block_size>>>
4、最大允许线程块大小：1024
最大允许网格大小：2³² − 1 （针对一维网格）

一维线程模型

1、每个线程在核函数中都有一个唯一的身份标识；
2、每个线程的唯一标识由这两个<<<grid_size, block_size>>>确定；grid_size, block_size保存在内建变量（build-in variable），目前考虑的是一维的情况:
（1）gridDim.x：该变量的数值等于执行配置中变量grid_size的值；
（2）blockDim.x：该变量的数值等于执行配置中变量block_size的值。
3、线程索引保存成内建变量（ build-in variable）：
（1）blockIdx.x：该变量指定一个线程在一个网格中的线程块索引值，范围为0~ gridDim.x-1；
（2）threadIdx.x：该变量指定一个线程在一个线程块中的线程索引值，范围为0~ blockDim.x-1。

例如 kernel_fun<<<2, 4>>>() ;
gridDim.x 的值为2
blockDim.x的值为4
blockIdx.x取值范围为0~1
threadIdx.x取值范围为0~3
线程唯一标识：
Idx = threadIdx.x + blockIdx.x * blockDim.x

多维线程

1、CUDA可以组织三维的网格和线程块；
2、blockIdx和threadIdx是类型为uint3(一种用于表示三维无符号整数的结构，它通常用于处理三维网格、坐标等数据)的变量，该类型是一个结构体，具有x,y,z三个成员（3个成员都为无符号类型的成员构成）：
$$\{\begin{array}{l}blockIdx.x\\ blockIdx.y\\ blockIdx.z\end{array}$$
$$\{\begin{array}{l}threadIdx.x\\ threadIdx.y\\ threadIdx.z\end{array}$$
3、gridDim和blockDim是类型为dim3(一个用于表示三维维度的结构，常用于定义网格和块的尺寸)的变量，该类型是一个结构体，具有x,y,z三个成员
$$\{\begin{array}{l}gridDim.x\\ gridDim.y\\ gridDim.z\end{array}$$
$$\{\begin{array}{l}blockDim.x\\ blockDim.y\\ blockDim.z\end{array}$$
4、取值范围
blockIdx.x 范围------------[0, gridDim.x-1] threadIdx.x 范围------------[0, blockDim.x-1]
blockIdx.y 范围------------[0, gridDim.y-1] threadIdx.y 范围------------[0, blockDim.y-1]
blockIdx.z范围-------------[0, gridDim.z-1] threadIdx.z 范围------------[0, blockDim.z-1]
注意：内建变量只在核函数有效，且无需定义！

定义多维网格和线程块（c++构造函数语法）：
dim3 grid_size(Gx, Gy, Gz);
dim3 block_size(Bx, By, Bz);
注意：若省略一维 比如dim3 grid_size(2, 2); //等价于dim3 grid_size(2, 2, 1)

多维网格和多维线程块本质是一维的，GPU物理上不分块。

每个线程都有唯一标识：
int tid = threadIdx.y * blockDim.x + threadIdx.x; int bid = blockIdx.y * gridDim.x + blockIdx.x;

多维线程块中的线程索引：
int tid = threadIdx.z * blockDim.x * blockDim.y +threadIdx.y * blockDim.x + threadIdx.x
多维网格中的线程块索引：
int bid = blockIdx.z * gridDim.x * gridDim.y +blockIdx.y * gridDin.x + blockIdx.x

6.2网格和线程块的限制条件

网格大小限制：
gridDim.x 最大值------------ 2³¹ − 1
gridDim.y 最大值------------ 2¹⁶ − 1
gridDim.z 最大值------------ 2¹⁶ − 1
线程块大小限制：
blockDim.x 最大值------------ 1024
blockDim.y 最大值------------ 1024
blockDim.z 最大值------------ 64
注意：线程块总的大小最大为1024！

7线程全局索引计算方式

7.1线程全局索引

一个线程块上的线程是放在同一个流式多处理器（SM）上的，但是单个 SM 的资源有限，这就使得 block 中的线程数是有限制的，现代 GPUs 的 block 可支持的线程数可达 1024 个，这就使得有时若想要知道一个线程在所有线程中的全局 ID，就必须要知道相应的组织结构，通过以下两个内置变量，可获得相应的组织结构信息

7.2不同组合形式

一维Grid 一维Block：
int blockId = blockIdx.x;
int id = blockIdx.x *blockDim.x + threadIdx.x;
一维Grid 二维Block:
int blockId = blockIdx.x;
int id = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
一维Grid 三维Block
int blockId = blockIdx.x;
int id = blockIdx.x * blockDim.x * blockDim.y * blockDim.z+ threadIdx.z * blockDim.y * blockDim.x+ threadIdx.y * blockDim.x + threadIdx.x;
二维Grid 一维Block：
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int id = blockId * blockDim.x + threadIdx.x;
二维Grid 二维Block:
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int id = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
二维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int id = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;
三维Grid 一维Block：
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z; int id = blockId * blockDim.x + threadIdx.x;
三维Grid 二维Block:
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z; int id = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
三维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int id = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;

7.3 示例

#include <stdio.h>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;

    const int id = threadIdx.x + blockIdx.x * blockDim.x; 
    printf("Hello World from block %d and thread %d, global id %d\n", bid, tid, id);
}
int main(void)
{
    hello_from_gpu<<<2, 4>>>();
    cudaDeviceSynchronize();

    return 0;
}

原文地址：https://blog.csdn.net/qq_52237775/article/details/142927877

免责声明：本站文章内容转载自网络资源，如本站内容侵犯了原著者的合法权益，可联系本站删除。更多内容请关注自学内容网（zxcms.com）！