资讯详情

NVIDA CUDA-DirverAPI入门

NVIDA公司发布了CUDA，CUDA是建立在NVIDA的CPUs基于通用并行计算平台和编程模型CUDA可以利用GPUs并行计算引擎可以更有效地解决更复杂的计算问题。近年来，GPU基于深度学习，最成功的应用领域是基于GPU并行计算已成为训练深度学习模型的标准。

GPU它不是一个需要和独立运行的计算平台CPU协同工作，可视为CPU所以当我们谈论协处理器时GPU并行计算实际上是基于CPU GPU异构计算架构。在异构计算架构中，GPU和CPU通过PCle连接总线协同工作，CPU位置叫主机端host，而GPU位置称为设备端device。

CUDA是NVIDA公司开发的GPU它提供了编程模型GPU基于编程的简单接口CUDA编程可以构建基于GPU应用程序的计算。

0x01 CUDA编程基础

（一）CUDA认知

(二)显卡，显卡驱动，NVCC、CUDA driver、CUDAtoolkit、cudnn是什么？

（三）runtime和driver API的区别

0x02 环境配置

0x03 cuInit——驱动初始化

0x04 CUcontext——上下文管理

（一）简介

(二)使用代码

0x05 Memory

（一）Pinned Memory与Pageable Memory

(二)内存分配

（三）代码

0x06 stream - 流

0x07 核函数

(1)编写核函数必须遵循CUDA那么有哪些规范呢？

(二)函数声明，__ global __ 、__ device __ 、__ host __三者有什么区别？

(3)如何在核函数内使用线程？为什么要分为三个层次？

（四）cu文件引入的新符号和语法强调

(五)核函数线程总数

（六）代码

0x08 共享内存

0x09 Warpaffine

(一)点的缩放平移旋转

(二)线性插值

(三)双线性差值

（五）代码

0x0A Yolov5加速后处理

0x0B Thrust and Error

0x01 CUDA编程基础

我认为参考价值较大的文章：

CUDA编程入门极简教程 - 知乎

显卡，显卡驱动,nvcc, cuda driver,cudatoolkit,cudnn到底是什么？ - 知乎

（一）CUDA认知

  • CUDA是一个异构模型，需要CPU和GPU协同工作。

  • 在CUDA中，host指代CPU及其内存，device指代GPU及其内存。

  • CUDA程序中既包含host程序又包含device程序，他们分别在CPU和GPU上运行，同时CPU与GPU之间还可以进行通讯。

  • CUDA Driver是与GPU沟通的驱动级别底层API，对应于cuda.h和libcuda.so文件

  • 典型的CUDA程序的执行流程如下：

    • 分配host内存，并进行数据初始化。

    • 分配device内存，并从host将数据拷贝到device上。

    • 调用CUDA的核函数在device上完成指定的运算。

    • 将device上的运算结果拷贝到host上。

    • 释放device和host上分配的内存。

在上面的过程中，调用了CUDA的核函数来执行并行计算，kernel是CUDA中一个重要的概念，kernel是在device上线程中并执行的函数，核函数用__ global __符号声明，在调用中需要用<<grid,block>>来指定kernel要执行的线程数量。在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号threadID，这个ID值可以通过核函数的内置变量threadIdx来获得。

  • 由于GPU实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词来区分开host和device上的函数，主要函数类型限定词如下：

    __ global __ ：在device上执行，从host中调用（一些特定的GPU也可以从device中调用），返回类型必须是void，不支持可变参数，不能成为类成员函数。注意用 __ global __定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。

    __ device __ ：在device上执行，但仅可以从device中调用，不可以和 __ global __ 同时用。

    __ host __ ：在host上执行，仅可以从host上调用，一般省略不写，不可以和 __ global __同时用，但可和 _ _ device _ _ ，此时函数会在device和host都编译。

（二）显卡、显卡驱动、NVCC、CUDA driver、CUDAtoolkit、cudnn是什么？

  • 显卡： 简单理解这个就是我们前面说的，尤其指NVIDIA公司生产的GPU系列，因为后面介绍的cuda,cudnn都是NVIDIA公司针对自身的GPU独家设计的。

  • 显卡驱动：驱动软件，类比声卡驱动，摄像头驱动。很明显就是字面意思，通常指，其实它就是一个驱动软件，而前面的就是硬件。

  • GPU架构：硬件设计方式。现在有Tesla、Fermi、Kepler、Maxwell、Pascal等。

  • CUDA：其中一种理解是它是一种编程语言（像c++,python等，只不过它是专门用来操控GPU的）

  • cudnn：这个其实就是一个专门为深度学习计算设计的软件库，里面提供了很多专门的计算函数，它里面包含了许多库，例如：cudart, cublas等。

  • NVCC：nvcc其实就是CUDA的编译器,可以从CUDA Toolkit的/bin目录中获取,类似于gcc就是c语言的编译器。由于程序是要经过编译器编程成可执行的二进制文件，而cuda程序有两种代码，一种是运行在cpu上的host代码，一种是运行在gpu上的device代码，所以nvcc编译器要保证两部分代码能够编译成二进制文件在不同的机器上执行。

  • nvidia -smi：nvidia-smi全程是NVIDIA System Management Interface ，它是一个基于前面介绍过的NVIDIA Management Library(NVML)构建的命令行实用工具，旨在帮助管理和监控NVIDIA GPU设备。

    那么NVCC、nvidia -smi的区别在哪：

    我们在windows上输入nvcc --version以及nvidia-smi得到的CUDA版本信息是不一样的。

>nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2021 NVIDIA Corporation
Built on Fri_Dec_17_18:28:54_Pacific_Standard_Time_2021
Cuda compilation tools, release 11.6, V11.6.55
Build cuda_11.6.r11.6/compiler.30794723_0
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 511.23       Driver Version: 511.23       CUDA Version: 11.6     |
|-------------------------------+----------------------+----------------------+
| GPU  Name            TCC/WDDM | Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|                               |                      |               MIG M. |
|===============================+======================+======================|
|   0  NVIDIA GeForce ... WDDM  | 00000000:01:00.0 Off |                  N/A |
| N/A   56C    P8    N/A /  N/A |      0MiB /  4096MiB |      1%      Default |
|                               |                      |                  N/A |
+-------------------------------+----------------------+----------------------+

+-----------------------------------------------------------------------------+
| Processes:                                                                  |
|  GPU   GI   CI        PID   Type   Process name                  GPU Memory |
|        ID   ID                                                   Usage      |
|=============================================================================|
|  No running processes found                                                 |
+-----------------------------------------------------------------------------+

CUDA有两个主要的API：和。这两个API都有对应的CUDA版本（runtime API和driver API），在网上的某个解释是这样的：

来源于：Different CUDA versions shown by nvcc and NVIDIA-smi - Stack Overflow

  • 用于支持的必要文件(如libcuda.so)是由安装的。nvidia-smi就属于这一类API。

  • 用于支持的必要文件(如libcudart.so以及nvcc)是由安装的。（CUDA Toolkit Installer有时可能会集成了GPU driver Installer）。nvcc是与CUDA Toolkit一起安装的CUDA compiler-driver tool，它只知道它自身构建时的CUDA runtime版本。它不知道安装了什么版本的GPU driver，甚至不知道是否安装了GPU driver。

综上，如果driver API和runtime API的CUDA版本不一致可能是因为你使用的是单独的GPU driver installer，而不是CUDA Toolkit installer里的GPU driver installer。

（三）runtime和driver API的区别

runtime和driverAPI在很多情况下非常相似，用起来其实都是差不多的，但是这两个API不可以混用，两者是互斥的，也就是说在开发时我们只能选择其中一种API。

  • CUDA Driver与CUDA Runtime相比更偏底层，就意味着Driver API有着更灵活的控制，也伴随着更复杂的编程。

  • 因此CUDA driver需要做显式的初始化cuInit(0)，否则其他API都会返回CUDA_ERROR_NOT_INITIALIZED

0x02 环境配置

这里网上有很多教程，那么就不过多描述：

最新CUDA环境配置(Win10 + CUDA 11.6 + VS2019)_扫地的小何尚的博客-CSDN博客_cuda环境配置

0x03 cuInit——驱动初始化

cuInit的意义是：初始化驱动API，如果不执行，则所有API都将返回错误，全局执行一次即可。

cuInit没有对应的cuDestroy，不需要释放，程序销毁自动释放。

cuInit的参数以及函数说明可以在NVIDA官网中查看：

CUDA Driver API :: CUDA Toolkit Documentation

说明：

 cuInit(int flags), 这里的flags目前必须给0;
 对于cuda的所有函数，必须先调用cuInit，否则其他API都会返回CUDA_ERROR_NOT_INITIALIZED

那么我们就先实现使用CUDA驱动头文件来进行初始化，获取我们的CUDA版本及其获取当前设备信息吧：

获取驱动版本管理：CUDA 驱动程序 API ：： CUDA Toolkit Documentation (nvidia.com)

获取设备信息管理：CUDA 驱动程序 API ：： CUDA Toolkit Documentation (nvidia.com)

// CUDA驱动头文件cuda.h
#include <cuda.h>
#include <stdio.h> 
#include <string.h>
int main(){

    /* 
    cuInit(int flags), 这里的flags目前必须给0;
    对于cuda的所有函数，必须先调用cuInit，否则其他API都会返CUDA_ERROR_NOT_INITIALIZED
     */
    CUresult code=cuInit(0);  //CUresult 类型：用于接收一些可能的错误代码
    if(code != CUresult::CUDA_SUCCESS){
        const char* err_message = nullptr;
        cuGetErrorString(code, &err_message);    // 获取错误代码的字符串描述
        // cuGetErrorName (code, &err_message);  // 也可以直接获取错误代码的字符串
        printf("Initialize failed. code = %d, message = %s\n", code, err_message);
        return -1;
    }
    /* 
    测试获取当前cuda驱动的版本
    显卡、CUDA、CUDA Toolkit

    1. 显卡驱动版本，比如：Driver Version: 460.84
    2. CUDA驱动版本：比如：CUDA Version: 11.2
    3. CUDA Toolkit版本：比如自行下载时选择的10.2、11.2等；这与前两个不是一回事, CUDA Toolkit的每个版本都需要最低版本的CUDA驱动程序
        
    三者版本之间有依赖关系, 可参照https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html
    nvidia-smi显示的是显卡驱动版本和此驱动最高支持的CUDA驱动版本
        
     */
    int driver_version = 0;
    code = cuDriverGetVersion(&driver_version);  // 获取驱动版本
    printf("CUDA Driver version is %d\n", driver_version); // 若driver_version为11020指的是11.2

    // 测试获取当前设备信息
    char device_name[100]; // char 数组
    CUdevice device = 0;
    code = cuDeviceGetName(device_name, sizeof(device_name), device);  // 获取设备名称、型号如：Tesla V100-SXM2-32GB // 数组名device_name当作指针
    printf("Device %d name is %s\n", device, device_name);
    return 0;
}

 那么NVIDA中如何对没有初始化的错误应该如何检查，NVIDA官网提供的查错函数：

// 使用有参宏定义检查cuda driver是否被正常初始化, 并定位程序出错的文件名、行数和错误信息
// 宏定义中带do...while循环可保证程序的正确性
#define checkDriver(op)    \
    do{                    \
        auto code = (op);  \
        if(code != CUresult::CUDA_SUCCESS){     \
            const char* err_name = nullptr;     \
            const char* err_message = nullptr;  \
            cuGetErrorName(code, &err_name);    \
            cuGetErrorString(code, &err_message);   \
            printf("%s:%d  %s failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, #op, err_name, err_message);   \
            return -1;   \
        }                \
    }while(0)

那么我们首先就来使用一下，相对于是一个helloword吧：

// CUDA驱动头文件cuda.h
#include <cuda.h>

#include <stdio.h>
#include <string.h>

// 使用有参宏定义检查cuda driver是否被正常初始化, 并定位程序出错的文件名、行数和错误信息
// 宏定义中带do...while循环可保证程序的正确性
#define checkDriver(op)    \
    do{                    \
        auto code = (op);  \
        if(code != CUresult::CUDA_SUCCESS){     \
            const char* err_name = nullptr;     \
            const char* err_message = nullptr;  \
            cuGetErrorName(code, &err_name);    \
            cuGetErrorString(code, &err_message);   \
            printf("%s:%d  %s failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, #op, err_name, err_message);   \
            return -1;   \
        }                \
    }while(0)

int main(){

    //检查cuda driver的初始化。虽然不初始化或错误初始化某些API不会报错（不信你试试），但安全起见调用任何API前务必检查cuda driver初始化
    cuInit(2); // 正确的初始化应该给flag = 0
    checkDriver(cuInit(0));

    // 测试获取当前cuda驱动的版本
    int driver_version = 0;
    checkDriver(cuDriverGetVersion(&driver_version));
    printf("Driver version is %d\n", driver_version);

    // 测试获取当前设备信息
    char device_name[100];
    CUdevice device = 0;
    checkDriver(cuDeviceGetName(device_name, sizeof(device_name), device));
    printf("Device %d name is %s\n", device, device_name);
    return 0;
}

其实官方给的查错代码还是有很多值得改进的地方的，首先它使用宏定义，每当我需要再加一行的时候，就需要再加一个“\”，这样很麻烦，而且可以看到do...while(0)的使用与顺序执行...的效果是一样的，但是前者可以保证程序的正确性，下面提供了一种改进的方法：

// 很明显，这种代码封装方式，更加的便于使用
//宏定义 #define <宏名>（<参数表>） <宏体>
#define checkDriver(op)  __check_cuda_driver((op), #op, __FILE__, __LINE__)

bool __check_cuda_driver(CUresult code, const char* op, const char* file, int line){

    if(code != CUresult::CUDA_SUCCESS){    
        const char* err_name = nullptr;    
        const char* err_message = nullptr;  
        cuGetErrorName(code, &err_name);    
        cuGetErrorString(code, &err_message);   
        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

那么我们的查错程序就可以这么修改：

// CUDA驱动头文件cuda.h
#include <cuda.h>
#include <stdio.h>
#include <string.h>

// 很明显，这种代码封装方式，更加的便于使用
//宏定义 #define <宏名>（<参数表>） <宏体>
#define checkDriver(op)  __check_cuda_driver((op), #op, __FILE__, __LINE__)

bool __check_cuda_driver(CUresult code, const char* op, const char* file, int line){

    if(code != CUresult::CUDA_SUCCESS){    
        const char* err_name = nullptr;    
        const char* err_message = nullptr;  
        cuGetErrorName(code, &err_name);    
        cuGetErrorString(code, &err_message);   
        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

int main(){

    // 检查cuda driver的初始化
    // 实际调用的是__check_cuda_driver这个函数
    checkDriver(cuInit(0));

    // 测试获取当前cuda驱动的版本
    int driver_version = 0;
    if(!checkDriver(cuDriverGetVersion(&driver_version))){
        return -1;
    }
    printf("Driver version is %d\n", driver_version);

    // 测试获取当前设备信息
    char device_name[100];
    CUdevice device = 0;
    checkDriver(cuDeviceGetName(device_name, sizeof(device_name), device));
    printf("Device %d name is %s\n", device, device_name);
    return 0;
}

0x04 CUcontext——上下文管理

（一）简介

参考：交互式 GPU 编程 - 第 3 部分 - CUDA 上下文恶作剧 (dragan.rocks)

  • context是一种上下文，。设备与特定进程相关连的所有状态。比如，你写的一段kernel code对GPU的使用会造成不同状态（内存映射、分配、加载的code），Context则保存着所有的管理数据来控制和使用设备。

  • context与一块显卡关联，一个显卡可以被多个context关联。GPU的context相对于CPU的program，一块GPU上可以有多个contexts，但是它们之间是相互隔离的，建议一块设备就一个context。

  • 每个线程都有一个栈结构储存context，栈顶是当前使用的context，对应有push、pop函数操作context的栈，所有api都以当前context为操作目标。

  • 它的作用是什么呢，我们可以想象到我们每次进行任何操作都需要传递一个device决定送到哪个设备执行，很麻烦，可以看看下面这种没有context的代码：

 那么再看看有context的代码：

  • 这样可以说是更安全且更方便了吧

  • context只是为了方便控制device的一种手段而提出来的，栈的存在是为了方便控制多个设备。

  • 由于高频操作，是一个线程基本固定访问一个显卡不变，且只使用一个context，很少会用到多context。

  • CreateContext、PushCurrent、PopCurrent这种多context的管理就显得很麻烦了，还可以再简单，因此推出了cuDevicePrimaryCtxRetain，为设备关联主context，分配、释放、设置、栈都不用你来控制管理了：

    看看有context的代码：

 再看看不用管理context栈的代码：（runtimeAPI自动使用cuDevicePrimaryCtxRetain）

  • primaryContext：一个显卡对应一个primary context，你只需要提供一个设备id，就可以把你的context设置好。

  • 不同线程，只要设备id一样，primary context就一样，context是线程安全的。

上下文可以做什么事情呢：

  1. 持有分配的内存列表

  2. 持有加载进设备的kernel code

  3. CPU与GPU之间的unified memory

  4. ...

如何管理上下文？

  1. 在cuda driver同样需要显示管理上下文

    • 开始时cuCtxCreate()创建上下文，结束时cuCtxDestroy销毁上下文。像文件管理一样须手动开关。

    • 用cuDevicePrimaryCtxRetain()创建上下文更好！

    • cuCtxGetCurrent()获取当前上下文

    • 可以使用堆栈管理多个上下文cuCtxPushCurrent()压入，cuCtxPopCurrent()推出

    • 对ctxA使用cuCtxPushCurrent()和cuCtxCreate()都相当于将ctxA放到栈顶（让它成为current context）

  2. cuda runtime可以自动创建，是基于cuDevicePrimaryCtxRetain()创建的。

（二）使用代码

// CUDA驱动头文件cuda.h
#include <cuda.h>   // include <> 和 "" 的区别    
#include <stdio.h>  // include <> : 标准库文件 
#include <string.h> // include "" : 自定义文件  

#define checkDriver(op)  __check_cuda_driver((op), #op, __FILE__, __LINE__)

bool __check_cuda_driver(CUresult code, const char* op, const char* file, int line){
    if(code != CUresult::CUDA_SUCCESS){    // 如果 成功获取CUDA情况下的返回值 与我们给定的值(0)不相等， 即条件成立， 返回值为flase
        const char* err_name = nullptr;    // 定义了一个字符串常量的空指针
        const char* err_message = nullptr;  
        cuGetErrorName(code, &err_name);    
        cuGetErrorString(code, &err_message);   
        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message); //打印错误信息
        return false;
    }
    return true;
}

int main(){

    // 检查cuda driver的初始化
    checkDriver(cuInit(0));

    // 为设备创建上下文
    CUcontext ctxA = nullptr;                                   // CUcontext 其实是 struct CUctx_st*（是一个指向结构体CUctx_st的指针）
    CUcontext ctxB = nullptr;
    CUdevice device = 0;
    checkDriver(cuCtxCreate(&ctxA, CU_CTX_SCHED_AUTO, device)); // 这一步相当于告知要某一块设备上的某块地方创建 ctxA 管理数据。输入参数 参考 https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDA__CTX_g65dc0012348bc84810e2103a40d8e2cf.html
    checkDriver(cuCtxCreate(&ctxB, CU_CTX_SCHED_AUTO, device)); // 参考 1.ctx-stack.jpg
    printf("ctxA = %p\n", ctxA);
    printf("ctxB = %p\n", ctxB);
    /* 
        contexts 栈：
            ctxB -- top <--- current_context
            ctxA 
            ...
     */

    // 获取当前上下文信息
    CUcontext current_context = nullptr;
    checkDriver(cuCtxGetCurrent(&current_context));             // 这个时候current_context 就是上面创建的context
    printf("current_context = %p\n", current_context);

    // 可以使用上下文堆栈对设备管理多个上下文
    // 压入当前context
    checkDriver(cuCtxPushCurrent(ctxA));                        // 将这个 ctxA 压入CPU调用的thread上。专门用一个thread以栈的方式来管理多个contexts的切换
    checkDriver(cuCtxGetCurrent(&current_context));             // 获取current_context (即栈顶的context)
    printf("after pushing, current_context = %p\n", current_context);
    /* 
        contexts 栈：
            ctxA -- top <--- current_context
            ctxB
            ...
    */
    
    // 弹出当前context
    CUcontext popped_ctx = nullptr;
    checkDriver(cuCtxPopCurrent(&popped_ctx));                   // 将当前的context pop掉，并用popped_ctx承接它pop出来的context
    checkDriver(cuCtxGetCurrent(&current_context));              // 获取current_context(栈顶的)
    printf("after poping, popped_ctx = %p\n", popped_ctx);       // 弹出的是ctxA
    printf("after poping, current_context = %p\n", current_context); // current_context是ctxB

    checkDriver(cuCtxDestroy(ctxA));
    checkDriver(cuCtxDestroy(ctxB));

    // 更推荐使用cuDevicePrimaryCtxRetain获取与设备关联的context
    // 注意这个重点，以后的runtime也是基于此, 自动为设备只关联一个context
    checkDriver(cuDevicePrimaryCtxRetain(&ctxA, device));       // 在 device 上指定一个新地址对ctxA进行管理
    printf("ctxA = %p\n", ctxA);
    checkDriver(cuDevicePrimaryCtxRelease(device));
    return 0;
}

0x05 Memory

首先要学会内存的各个单元，这样才可以使用cuda进行高效的内存分配。

CUDA的内存模型：

 首先可以看看GPU的电路板：

 那么GPU与电脑主板的关系：

（一）Pinned Memory与Pageable Memory

对于整个Host Memory内存条而言，操作系统区分为两个大类（逻辑区分，物理上是同一个东西）：

  • Pageable Memory：可分页内存

  • Page Lock Memory（Pinned Memory）：页锁定内存

怎么理解呢：

Pageable Memory可以理解为是VIP房间，锁定给你一个人用。而Page Lock Memory是普通房间，当酒店的房间不够用的时候，会选择性把你的房间腾出来给其他人用，这就可以容纳更多人了。这其实有点像是虚拟内存的东西，造成了房间很多的这个假象，代价是性能降低。

那么再做一下总结：

  • pinned memory具有锁定特性，是稳定不会被交换的（这很重要，相当于每次去这个房间都一定能找到你）。

  • pageable memory没有锁定特性，对于第三方设备（比如GPU），去访问时，因为无法感知内存是否被交换，可能得不到正确的数据（每次去房间找，说不准你的房间被人交换了）。

  • pageable memory的性能比pinned memory差，很可能降低你程序的优先级然后把内存交换给别人用。

  • pageable memory策略能使用内存假象，实际8GB但是可以使用15GB，提高程序运行数量（不是速度）。

  • pinned memory太多，会导致操作系统整体性能降低（程序运行数量减少），8GB就只能用8GB。注意不是你的应用程序性能降低。

  • 。

那么当数据传到GPU时，路径是这样的：

 当我们使用Pageable Memory时，我们还需要经过Pinned Memory才可以把数据传到DRAM，这当我们数据规模大时，大打折扣；而使用Pinned Memory时，我们是直接传输到DRAM。也就是说当我们使用显卡访问Pinnned Memory轨迹时，是这样的：

通过PICE接口，到主板，再到内存条。

那么我们再详细地概括一下总结：

1.GPU可以直接访问pinned memory，称之为（DMA Direct Memory Access），以前接触单片机的时候觉得DMA只是一个越过CPU访问内存的一个小东西，没想到在这里直接上升为GPU。

2.，所以PinnedMemory<GlobalMemory<SharedMemory

3.代码中，由分配的，是，由分配的是，由分配的是。

4.尽量多用PinnedMemory储存host数据，或者显式处理Host到Device时，用PinnedMemory做缓存，都是提高性能的关键。

（二）内存分配

统一内存（Unified addressing）：

参考博客：初步介绍CUDA中的统一内存_扫地的小何尚的博客-CSDN博客_cuda统一内存

可以非常轻松地分配合访问可由系统中任何处理器、CPU或GPU上运行代码使用的数据。这种硬件/软件技术允许应用程序分配可以从CPU或GPU上运行的代码读取或写入的数据。

  • 分配线性内存cuMemAlloc():

    线性内存：线性内存被组织在单个连续的地址空间中，可以直接以及线性地访问这些内存位置。内存分配空间以字节为大小，并返回所分配的内存地址。

 

  • 分配主机锁页内存cuMemAllocHost():

锁页内存：页面不允许被调入调出的叫锁页内存，反之叫可分页内存。

好处：快。

  • a. 设备可以直接访问内存，与可分页内存相比，它的读写带宽要高得多

  • b. 驱动程序会跟踪使用cuMemAllocHost()分配的虚拟内存范围，并自动加速对cuMemcpy（）等函数的调用。

  • 使用注意：分配过多锁业内存会减少系统可用于分页的内存量，可能会降低系统性能。因此，在主机和设备之间为数据交换分配临时区域时，最好少用此功能。

这里是从主机分配内存，因此不是输入device prt的地址，而是一个。

使用：

内存的初始化cuMemsetD32(CUdeviceptr dstDevice, unsigned int ui, size_t N), 将N个32位值的内存范围设置为指定的值ui。

内存的释放cuMemFreeHost(): 有借有还 再借不难。

选择右边的内存分配模型：

参考博客：用于数据传输的页面锁定主机内存 - 雷毛的日志 (leimao.github.io)

主机和 CUDA 设备之间的数据传输需要主机上的页面锁定内存。如果数据不在主机上的可分页内存中，则在将数据从主机传输到设备期间，数据将从可分页主机内存隐式传输到临时页面锁定的主机内存，然后将数据从页面锁定的主机内存传输到设备内存。

这会带来额外的数据传输开销，并且在某些情况下可能会显著影响计算机程序的性能。因此，问题是为什么主机和CUDA设备之间的数据传输需要在主机上使用页面锁定内存。我认为主要动机是保证数据传输效率。

由于主机和CUDA设备之间的数据传输无论如何都必须使用，因此为了优化数据传输，我们可以直接分配页面锁定内存来存储数据，而不是分配可分页内存来存储数据。

需要注意的是，手动使用大量页面锁定内存可能会导致操作系统性能问题。手动使用页面锁定内存意味着用户负责分配和释放页面锁定内存（内存泄漏）。如果用户未及时释放页面锁定内存，则物理内存上可用于新数据和新应用程序的可用内存将减少。因此，操作系统可能会变得不稳定。如果页面锁定内存仅在可分页内存和 CUDA 内存之间的数据传输期间临时使用，则它将始终及时释放，但代价是在可分页内存和页面锁定内存之间增加数据传输开销。

因此，我们被鼓励使用页面锁定的内存，但我们不应该滥用它。通常，如果我们知道数据将在主机和CUDA设备之间多次传输，则最好将数据放在页面锁定的内存中以避免不必要的开销。

（三）代码

// CUDA驱动头文件cuda.h
#include <cuda.h>
#include <stdio.h>
#include <string.h>

#define checkDriver(op)  __check_cuda_driver((op), #op, __FILE__, __LINE__)

bool __check_cuda_driver(CUresult code, const char* op, const char* file, int line){

    if(code != CUresult::CUDA_SUCCESS){    
        const char* err_name = nullptr;    
        const char* err_message = nullptr;  
        cuGetErrorName(code, &err_name);    
        cuGetErrorString(code, &err_message);   
        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

int main(){

    // 检查cuda driver的初始化
    checkDriver(cuInit(0));

    // 创建上下文
    CUcontext context = nullptr;
    CUdevice device = 0;
    checkDriver(cuCtxCreate(&context, CU_CTX_SCHED_AUTO, device));
    printf("context = %p\n", context);

    // 输入device prt向设备要一个100 byte的线性内存，并返回地址
    CUdeviceptr device_memory_pointer = 0;
    checkDriver(cuMemAlloc(&device_memory_pointer, 100)); // 注意这是指向device的pointer, 
    printf("device_memory_pointer = %p\n", device_memory_pointer);

    // 输入二级指针向host要一个100 byte的锁页内存，专供设备访问。
    float* host_page_locked_memory = nullptr;
    checkDriver(cuMemAllocHost((void**)&host_page_locked_memory, 100));
    printf("host_page_locked_memory = %p\n", host_page_locked_memory);

    // 向page-locked memory 里放数据（仍在CPU上），可以让GPU可快速读取
    host_page_locked_memory[0] = 123;
    printf("host_page_locked_memory[0] = %f\n", host_page_locked_memory[0]);
    /* 
        记住这一点
        host page locked memory 声明的时候为float*型，可以直接转换为device ptr，这才可以送给cuda核函数（利用DMA(Direct Memory Access)技术）
        初始化内存的值: cuMemsetD32 ( CUdeviceptr dstDevice, unsigned int  ui, size_t N )
        初始化值必须是无符号整型，因此需要将new_value进行数据转换：
        但不能直接写为:(int)value，必须写为*(int*)&new_value, 我们来分解一下这条语句的作用：
        1. &new_value获取float new_value的地址
        2.(int*)将地址从float * 转换为int*以避免64位架构上的精度损失
        3.*(int*)取消引用地址，最后获取引用的int值
     */
    
    float new_value = 555;
    checkDriver(cuMemsetD32((CUdeviceptr)host_page_locked_memory, *(int*)&new_value, 1)); 
    printf("host_page_locked_memory[0] = %f\n", host_page_locked_memory[0]);

    // 释放内存
  checkDriver(cuMemFreeHost(host_page_locked_memory));
    return 0;
}

0x06 stream - 流

CUDA流表示一个GPU操作队列，该队列中的操作将以添加到流中的先后顺序而一次执行。可以将一个流看做是GPU上的一个任务，不同的任务可以。使用CUDA流，首先要选择一个支持的设备，支持设备重叠功能的GPU能够在执行一个CUDA核函数的同时，还能在主机和设备之间执行复制数据操作。

支持重叠功能的设备的这一特性很重要，可以在一定程度上提升GPU程序的执行效率。一般情况下，CPU内存远大于GPU内存，对于数据量比较大的情况，不可能把CPU缓冲区中的数据一次性传输给GPU，需要分块传输，，这样的异步操作，就用到设备的重叠功能，能够提高运算性能。说白了，其实就是这个设备具有多个核，可以同时进行工作。

  • 流是一种基于context之上的，一个context可以创建n个流。

  • 流是的主要方式。

  • nullptr表示默认流，。

  • 某一方发出指令后，他就可以做任何事情，无需等待指令执行完毕，指令发出的耗时也是极短的，这是一个异步操作，执行的代码加入流的队列后，立即返回，不耽误时间。

  • 接收指令的一方，根据流的队列，顺序执行；发出指令的一方具有选择性，在需要的时候等待所有的执行结果。

  • 新建一个流，就是新建一个接收指令的一方，可以新建很多个流。

  • 通过可以选择性等待任务队列中的部分任务是否就绪。

需要注意的是：

  1. 要十分注意，指令发出后，流队列中储存的是指令参数，不能加入队列后立即释放参数指针，这会导致流队列执行该指令时指针失效而出错。

  2. 应当在十分肯定流已经不需要这个指针后，才进行修改或者释放，否则会有非预期结果出现。

总结来说：

  1. stream是一个流句柄，可以当做是一个

    • cuda执行器从stream中一条条的读取并执行指令。

    • 例如函数等同于向stream这个队列中加入一个指令并排队。

    • 使用到了stream的函数，便立即向stream中加入指令后立即返回，并不会等待指令执行结束。

    • 通过函数，等待stream中所有指令执行完毕，也就是队列为空。

  2. 当使用stream时，需要注意

    • 由于异步函数会立即返回，因此传递进入的参数要考虑其生命周期，应。

  3. 还可以向stream中加入Event，用以监控是否到达了某个检查点

    • cudaEventCreate，创建事件

    • cudaEventRecord，记录事件，即在stream中加入某个事件，当队列执行到该事件后，修改其状态

    • cudaEventQuery，查询事件当前状态

    • cudaEventElapsedTime，计算两个事件之间经历的时间间隔，若要统计某些核函数执行时间，请使用这个函数，能够得到最准确的统计

    • cudaEventSynchronize，同步某个事件，等待事件到达

    • cudaStreamWaitEvent，等待流中的某个事件

  4. 默认流，对于cudaMemcpy等同步函数，其等价于执行了

    • cudaMemcpyAsync(... 默认流) 加入队列

    • cudaStreamSynchronize(默认流) 等待执行完成

    • 默认流与当前设备上下文类似，是与当前设备进行的关联

    • 因此，如果大量使用默认流，会导致性能低下

如何使用：

// CUDA运行时头文件
#include <cuda_runtime.h>
#include <stdio.h>
#include <string.h>

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

int main(){

    int device_id = 0;
    checkRuntime(cudaSetDevice(device_id));

    cudaStream_t stream = nullptr;
    checkRuntime(cudaStreamCreate(&stream));

    // 在GPU上开辟空间
    float* memory_device = nullptr;
    checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float)));

    // 在CPU上开辟空间并且放数据进去，将数据复制到GPU
    float* memory_host = new float[100];
    memory_host[2] = 520.25;
    checkRuntime(cudaMemcpyAsync(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice, stream)); // 异步复制操作，主线程不需要等待复制结束才继续

    // 在CPU上开辟pin memory,并将GPU上的数据复制回来 
    float* memory_page_locked = nullptr;
    checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float)));
    checkRuntime(cudaMemcpyAsync(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost, stream)); // 异步复制操作，主线程不需要等待复制结束才继续
    checkRuntime(cudaStreamSynchronize(stream));
    
    printf("%f\n", memory_page_locked[2]);
    
    // 释放内存
    checkRuntime(cudaFreeHost(memory_page_locked));
    checkRuntime(cudaFree(memory_device));
    checkRuntime(cudaStreamDestroy(stream));
    delete [] memory_host;
    return 0;
}

0x07 核函数

CUDA核函数是指GPU端运行的代码，核函数内部主要干了什么？简而言之，就是规定GPU的各个线程访问哪个数据并执行什么计算。

通过xxx.cu创建一个cudac程序文件，并把cu交给nvcc编译，才能识别cuda语法。

参考博客：

1.2.CUDA核函数 - 知乎

（一）编写核函数必须遵循CUDA规范，那么有哪些规范？

  1. 必须写在*.cu文件中。

  2. 必须以限定符声明定义（只有他才可以执行核函数）。

  3. 返回类型必须是。

  4. 不支持可变数量的参数。调用核函数是传值的，不能传引用，可以传递类、结构体等，核函数可以是模板。

  5. 核函数内部只能访问设备内存。

  6. 核函数内部不能使用静态变量。

（二）函数声明中，__ global __ 、__ device __ 、__ host __三者区别是什么？

1） __ global __ 修饰的函数是，在设备端执行，可以从主机端调用，也可以在sm3以上的设备端调用（比如动态并行）；只有__ global __ 修饰的函数才可以用<<<>>>的方式调用。 2） __ device __ 修饰的函数是，在设备端执行，只能从设备端调用； 3） __ host __ 修饰的函数是，在主机端执行，只能从主机端调用；

4） __ device __ 和 __ host __ 可以一起使用，来表示该函数可以；

5）__ shared __ 表示变量为共享变量。

6） nvcc编译选项中添加-dc（相当于--relocatable-device-code=true --compile）时，函数可以调用其它文件中的函数，否则只能调用同文件中的函数。

（三）核函数内部怎么使用线程？为什么要分成三个层次？

CUDA从逻辑上将GPU线程分成了三个层次——。

每个核函数对应一个线程格，一个线程格中有一个或多个线程块，一个线程块中有一个或多个线程。

关于GPU的硬件描述可以看这篇：

2.1.GPU硬件架构 - 知乎

CUDA核函数中文说明要将线程分为三个层次，其实是与GPU的硬件组成相关联的。在GPU硬件中本身就存在三个层次——，这是一种类似于计算机集群的层次结构，而我们编写的核函数正是运行在这种层次结构上，所以核函数必须支持这三个层次，否则任务无法顺利分解，也就无法从高层次向低层次传递。

实际设计中，CUDA将这种对应关系规定为：

  1. 分配到上运行。

  2. 分配到上运行。SM可以理解为一个计算机集群。

  3. 分配到上运行。Core为计算核心，负责浮点数和整数计算。

（四）cu文件引入的新的符号和语法强调

  • __global__标记，核函数标记

    • 调用方必须是

    • 返回值必须是

    • 例如：__global__ void kernel(const float* pdata, int ndata)

    • 必须以kernel<<<gridDim, blockDim, bytesSharedMemorySize, stream>>>(pdata, ndata)的方式启动

      • 其参数类型是：<<<dim3 gridDim, dim3 blockDim, size_t bytesSharedMemorySize, cudaStream_t stream>>>

        • dim3有构造函数dim3(int x, int y=1, int z=1)

          • 因此当直接赋值为int时，实则定义了dim.x = value, dim.y = 1, dim.z = 1

      • 其中gridDim, blockDim, bytesSharedMemory, stream是线程layout参数

        • 如果指定了stream，则把

      • pdata和ndata则是核函数的函数调用参数

      • 函数调用参数必须传值，不能传引用等。参数可以是类类型等

    • 核函数的执行无论stream是否为nullptr，都将是

      • 因此在核函数中进行printf操作，你必须进行等待，例如cudaDeviceSynchronize、或者cudaStreamSynchronize，否则你将无法看到打印的信息

  • __device__标记，设备调用的函数

    • 调用方必须是device

  • __host__标记，主机调用函数

    • 调用方必须是主机

  • 也可以__device__ __host__两个标记同时有，表明该函数可以设备也可以主机

  • __constant__标记，定义

  • __shared__标记，定义

  • 可以通过函数，捕获核函数是否出现错误或者异常

（五）核函数的线程总数

host调用核函数：function<<<gridDim, blockDim, sharedMemorySize, stream>>>(args…);

在核函数内部有四个非常有用的内置变量——。线程layout主要用到、。

首先，我们可以先不严谨地认为是一个立方体，这个立方体有很多小方块：

 之后其中的每一个小块都是一个thread，为了方便讨论，我们只考虑2D：

我们关心的是某一个thread的位置，比如图上的黄色方块：

其位置为(blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y) =(1, 0, 1, 1)，如果把这个2D转换为1D，这格黄色的thread的1D位置为13。

我们可以通过blockIdx索引到线程块，通过threadIdx索引到某个块内的线程，通过blockDim得到一个块内线程总数，通过gridDim得到一个格内块总数：

那么我们可以得到在一维的条件下，计算公式为：

在一维的条件下，计算总数为：

那我们就可以暗中维度高低排序来看待这个信息：

 之后计算的话就这么计算：（左乘右加）

（六）代码

kernel.cu：

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void test_print_kernel(const float* pdata, int ndata){

    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    /*    dims                 indexs
        gridDim.z            blockIdx.z
        gridDim.y            blockIdx.y
        gridDim.x            blockIdx.x
        blockDim.z           threadIdx.z
        blockDim.y           threadIdx.y
        blockDim.x           threadIdx.x

        Pseudo code:
        position = 0
        for i in 6:
            position *= dims[i]
            position += indexs[i]
    */
    printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

void test_print(const float* pdata, int ndata){

    // <<<gridDim, blockDim, bytes_of_shared_memory, stream>>>
    test_print_kernel<<<1, ndata, 0, nullptr>>>(pdata, ndata);

    // 在核函数执行结束后，通过cudaPeekAtLastError获取得到的代码，来知道是否出现错误
    // cudaPeekAtLastError和cudaGetLastError都可以获取得到错误代码
    // cudaGetLastError是获取错误代码并清除掉，也就是再一次执行cudaGetLastError获取的会是success
    // 而cudaPeekAtLastError是获取当前错误，但是再一次执行 cudaPeekAtLastError 或者 cudaGetLastError 拿到的还是那个错
    // cuda的错误会传递，如果这里出错了，不移除。那么后续的任意api的返回值都会是这个错误，都会失败
    cudaError_t code = cudaPeekAtLastError();
    if(code != cudaSuccess){    
        const char* err_name    = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("kernel error %s:%d  test_print_kernel failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, err_name, err_message);   
    }
}

main.cpp：


#include <cuda_runtime.h>
#include <stdio.h>

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

void test_print(const float* pdata, int ndata);

int main(){
    float* parray_host = nullptr;
    float* parray_device = nullptr;
    int narray = 10;
    int array_bytes = sizeof(float) * narray;

    parray_host = new float[narray];
    checkRuntime(cudaMalloc(&parray_device, array_bytes));

    for(int i = 0; i < narray; ++i)
        parray_host[i] = i;
    
    checkRuntime(cudaMemcpy(parray_device, parray_host, array_bytes, cudaMemcpyHostToDevice));
    test_print(parray_device, narray);
    checkRuntime(cudaDeviceSynchronize());

    checkRuntime(cudaFree(parray_device));
    delete[] parray_host;
    return 0;
}

0x08 共享内存

共享内存特性：

共享内存的主要特点在于“共享”，也即，所以“共享”是针对同一个线程块所有线程而言的。。

  • 共享内存更靠近计算单元，所以访问速度更快。

  • 共享内存通常可以作为访问全局内存的缓存使用。

  • 可以利用共享内存实现线程间的通信。

  • 通常与同时出现，这个函数是，全部执行到这一行才往下走。

  • 使用方式，通常是在，然后syncthreads，然后再使用。

共享内存在CUDA核函数中定义，通常有两种方式：静态方式、动态方式。

  • 静态方式定义：这种定义的特点是定义的同时指定大小，并且定义内存大小通常地址是不一样的。

const size_t static_shared_memory_num_element = 6 * 1024; // 6KB
__shared__ char static_shared_memory[static_shared_memory_num_element]; 
__shared__ char static_shared_memory2[2];
  •  动态方式定义：这种定义的特点是定义的时候，在调用核函数的时候将共享内存大小以输入参数的形式传入。动态定义共享内存时，是使用进行传入数值的，传入的数值必须以字节byte为单位。如果定义的共享内存不是byte类型，数值必须乘以类型所占用的字节数。
// 表示的是传入12个字节大小的共享内存
demo1_kernel<<<1, 1, 12, nullptr>>>();

  • 如何知道本机共享内存的大小：

    指示了block中最大可用的共享内存，共享内存使得block内存内的threads可以相互通信。

  • 共享内存是，更靠近计算单元，因此比globalMem速度更快，通常可以充当使用。因此在运算的时候，可以数据先读入到sharedMem，做各类计算时，使用sharedMem而非globalMem。

  • 动态共享内存的使用注意：

    • dynamic_shared_memory变量必须使用extern 开头

    • 并且定义为不确定大小的数组[]

    • 变量放在函数外面和里面都一样

    • 其指针由执行时赋值

    • 动态共享内存无论定义多少个，地址都一样

  • 静态共享内存的使用注意：

    • 不加extern，以开头

    • 定义时需要明确数组的大小

    • 静态分配的地址比动态分配的地址低

    • 静态共享内存定义几个地址随之叠加

  • 如果配置的各类共享内存总和大于sharedMemPerBlock，则核函数执行出错，Invalid argument。

    • 不同类型的静态共享变量定义，其内存划分并不一定是连续的。

    • 中间会有策略，使得第一个和第二个变量之间可能存在空隙。

    • 因此你的变量之间如果存在空隙，可能小于全部大小的共享内存就会报错。

代码：

shared-memory.cu：

#include <cuda_runtime.h>
#include <stdio.h>

//demo1 //
/* 
demo1 主要为了展示查看静态和动态共享变量的地址
 */
const size_t static_shared_memory_num_element = 6 * 1024; // 6KB
__shared__ char static_shared_memory[static_shared_memory_num_element]; 
__shared__ char static_shared_memory2[2]; 

__global__ void demo1_kernel(){
    extern __shared__ char dynamic_shared_memory[];      // 静态共享变量和动态共享变量在kernel函数内/外定义都行，没有限制
    extern __shared__ char dynamic_shared_memory2[];
    printf("static_shared_memory = %p\n",   static_shared_memory);   // 静态共享变量，定义几个地址随之叠加
    printf("static_shared_memory2 = %p\n",  static_shared_memory2); 
    printf("dynamic_shared_memory = %p\n",  dynamic_shared_memory);  // 动态共享变量，无论定义多少个，地址都一样
    printf("dynamic_shared_memory2 = %p\n", dynamic_shared_memory2); 

    if(blockIdx.x == 0 && threadIdx.x == 0) // 第一个thread
        printf("Run kernel.\n");
}

/demo2//
/* 
demo2 主要是为了演示的是如何给 共享变量进行赋值
 */
// 定义共享变量，但是不能给初始值，必须由线程或者其他方式赋值
__shared__ int shared_value1;

__global__ void demo2_kernel(){
    
    __shared__ int shared_value2;
    if(threadIdx.x == 0){

        // 在线程索引为0的时候，为shared value赋初始值
        if(blockIdx.x == 0){
            shared_value1 = 123;
            shared_value2 = 55;
        }else{
            shared_value1 = 331;
            shared_value2 = 8;
        }
    }

    // 等待block内的所有线程执行到这一步
    __syncthreads();
    
    printf("%d.%d. shared_value1 = %d[%p], shared_value2 = %d[%p]\n", 
        blockIdx.x, threadIdx.x,
        shared_value1, &shared_value1, 
        shared_value2, &shared_value2
    );
}

void launch(){
    
    demo1_kernel<<<1, 1, 12, nullptr>>>();
    demo2_kernel<<<2, 5, 0, nullptr&g

标签： 100pin矩形连接器7p8pin连接器集成电路x0tps767d318集成电路st层集成电路

锐单商城拥有海量元器件数据手册、 IC替代型号，打造 电子元器件IC百科大全！

锐单商城 - 一站式电子元器件采购平台