模型部署——cuda编程入门

CUDA中的线程与线程束

在这里插入图片描述

  • kernel是在device上线程中并行执行的函数,核函数用__global__符号声明,在调用时需要用<<<grid_size, block_size>>>来指定kernel要执行的线程数量。在CUDA中,每一个线程都要执行核函数,并且每个线程会分配一个唯一的线程号thread ID,这个ID值可以通过核函数的内置变量threadIdx来获得。
  • synchronize是同步的意思,有几种synchronize
    cudaDeviceSynchronize: CPU与GPU端完成同步,CPU不执行之后的语句,直到这个语句以前的所有cuda操作结束
    cudaStreamSynchronize: 跟cudaDeviceSynchronize很像,但是这个是针对某一个stream的。只同步指定的stream中的cpu/gpu操作,其他的不管
    cudaThreadSynchronize: 现在已经不被推荐使用的方法
    __syncthreads: 线程块内同步
  • 核函数编写和调用举例
#include <cuda_runtime.h>
#include <stdio.h>// 核函数
__global__ void print_idx_kernel(){printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}void print_one_dim(){int inputSize = 8;int blockDim = 4;int gridDim = inputSize / blockDim;dim3 block(blockDim);dim3 grid(gridDim);// 核函数调用print_idx_kernel<<<grid, block>>>();cudaDeviceSynchronize();
}

.cu与.cpp的相互引用及Makefile

编译器:gcc g++ nvcc
举个例子:

nvcc print_index.cu -o app -I /usr/local/cuda/include

获取编译器选项:

g++ --help
nvcc --help

Makefile编写(是否可以使用CMakeLists.txt?)
.cpp中不能直接调用核函数,需要在.cu中提供调用接口

使用CUDA进行MATMUL计算

host端与device端数据传输

在这里插入图片描述
host端与device端数据传输代码实现:

void MatmulOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize){/* 设置矩阵大小 */int size = width * width * sizeof(float);/* 分配M, N在GPU上的空间*/float *M_device;float *N_device;cudaMalloc(&M_device, size);cudaMalloc(&N_device, size);/* 分配M, N拷贝到GPU上*/cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice);cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice);/* 分配P在GPU上的空间*/float *P_device;cudaMalloc(&P_device, size);/* 调用kernel来进行matmul计算, 在这个例子中我们用的方案是:将一个矩阵切分成多个blockSize * blockSize的大小 */dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);MatmulKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);/* 将结果从device拷贝回host*/cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost);cudaDeviceSynchronize();/* Free */// free与malloc的顺序是反着的cudaFree(P_device);cudaFree(N_device);cudaFree(M_device);
}

cuda core矩阵乘法核函数编写

/* matmul的函数实现*/
__global__ void MatmulKernel(float *M_device, float *N_device, float *P_device, int width){/* 我们设定每一个thread负责P中的一个坐标的matmul所以一共有width * width个thread并行处理P的计算*/// 确定负责计算的结果元素的索引int y = blockIdx.y * blockDim.y + threadIdx.y;int x = blockIdx.x * blockDim.x + threadIdx.x;float P_element = 0;/* 对于每一个P的元素,我们只需要循环遍历width次M和N中的元素就可以了*/for (int k = 0; k < width; k ++){float M_element = M_device[y * width + k];float N_element = N_device[k * width + x];P_element += M_element * N_element;}P_device[y * width + x] = P_element;
}

cuda core 每个线程执行核函数计算一个结果元素
GPU刚开始执行核函数的时候,会存在一个warmup阶段,耗时会比较长
CPU与GPU的浮点运算会存在误差,误差控制在e-4以内是ok的
CUDA中规定,一个block中可以分配的thread的数量最大是1024个线程,如果大于1024会显示配置错误
为什么block size = 1的时候比等于16的时候慢很多?

cuda中的error handler

在这里插入图片描述

获取GPU的硬件信息

利用cuda runtime api打印硬件信息 & LOG
#include <stdio.h>
#include <cuda_runtime.h>
#include <string>#include "utils.hpp"int main(){int count;int index = 0;cudaGetDeviceCount(&count);while (index < count) {cudaSetDevice(index);cudaDeviceProp prop;cudaGetDeviceProperties(&prop, index);LOG("%-40s",             "*********************Architecture related**********************");LOG("%-40s%d%s",         "Device id: ",                   index, "");LOG("%-40s%s%s",         "Device name: ",                 prop.name, "");LOG("%-40s%.1f%s",       "Device compute capability: ",   prop.major + (float)prop.minor / 10, "");LOG("%-40s%.2f%s",       "GPU global meory size: ",       (float)prop.totalGlobalMem / (1<<30), "GB");LOG("%-40s%.2f%s",       "L2 cache size: ",               (float)prop.l2CacheSize / (1<<20), "MB");LOG("%-40s%.2f%s",       "Shared memory per block: ",     (float)prop.sharedMemPerBlock / (1<<10), "KB");LOG("%-40s%.2f%s",       "Shared memory per SM: ",        (float)prop.sharedMemPerMultiprocessor / (1<<10), "KB");LOG("%-40s%.2f%s",       "Device clock rate: ",           prop.clockRate*1E-6, "GHz");LOG("%-40s%.2f%s",       "Device memory clock rate: ",    prop.memoryClockRate*1E-6, "Ghz");LOG("%-40s%d%s",         "Number of SM: ",                prop.multiProcessorCount, "");LOG("%-40s%d%s",         "Warp size: ",                   prop.warpSize, "");LOG("%-40s",             "*********************Parameter related************************");LOG("%-40s%d%s",         "Max block numbers: ",           prop.maxBlocksPerMultiProcessor, "");LOG("%-40s%d%s",         "Max threads per block: ",       prop.maxThreadsPerBlock, "");LOG("%-40s%d:%d:%d%s",   "Max block dimension size:",     prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2], "");LOG("%-40s%d:%d:%d%s",   "Max grid dimension size: ",     prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2], "");index ++;printf("\n");}return 0;
}
Roofline model(待补充)

Nsight system and Nsight compute

谷歌搜索下载:官网链接

Nsight system

参考链接

安装目录:
ls /usr/local/bin |grep nsys
nsys
nsys-ui启动GUI界面
sudo ./nsys-ui(不加sudo会存在权限问题)

举个例子:
配置可执行文件以及感兴趣内容:
在这里插入图片描述可视化分析:
在这里插入图片描述详细使用手册:官网文档

Nsight compute
查看可安装版本:
sudo apt policy nsight-compute-2022.2.1
安装:
sudo apt install nsight-compute-2022.2.1
查看安装位置:
dpkg -L nsight-compute-2022.2.1
路径:/opt/nvidia/nsight-compute/2022.2.1/
文件:ncu ncu-ui等启动:
sudo ./ncu-ui

举个例子:
基本配置:replay mode: application
在这里插入图片描述选择感兴趣内容:
在这里插入图片描述launch即可,第一次运行会比较慢,会重复运行很多次。
结果:
在这里插入图片描述不知道为什么roofline model没有正常显示出来,需要查一查?

扩展知识

共享内存以及BANK CONFLICT

shared memory

硬件结构

SM(Streaming Multiprocessor)
在CUDA编程模型中,线程被组织成线程块(block),多个线程块组成一个网格(grid)。每个线程块被分配到一个SM中执行,而SM内部的warp调度器会将线程块中的线程分成多个warp进行执行。
当一个warp中的线程需要等待某些操作(例如内存访问)完成时,SM可以切换到另一个warp继续执行,从而提高计算效率。
在这里插入图片描述

核函数编写
#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticKernel(float *M_device, float *N_device, float *P_device, int width){__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE];/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素,我们只需要循环遍历width / tile_width 次就okay了,这里有点绕,画图理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)];N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty)* width + x];__syncthreads(); // 上述两句所有thread都会执行,等待所有thread执行完成for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[ty][k] * N_deviceShared[k][tx];}__syncthreads();}P_device[y * width + x] = P_element;
}__global__ void MatmulSharedDynamicKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 声明动态共享变量的时候需要加extern,同时需要是一维的 注意这里有个坑, 不能够像这样定义: __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因为在cuda中定义动态共享变量的话,无论定义多少个他们的地址都是一样的。所以如果想要像上面这样使用的话,需要用两个指针分别指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = blockSize * blockSize;/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素,我们只需要循环遍历width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {deviceShared[ty * blockSize + tx] = M_device[y * width + (m * blockSize + tx)];deviceShared[stride + (ty * blockSize + tx)] = N_device[(m * blockSize + ty)* width + x];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[ty * blockSize + k] * deviceShared[stride + (k * blockSize + tx)];}__syncthreads();}if (y < width && x < width) {P_device[y * width + x] = P_element;}
}

动态共享内存比静态共享内存速度慢,没有特殊情况下,使用静态共享内存。

cuda event进行时间测算

在这里插入图片描述

BANK CONFLICT(存储体冲突)

在shared memory中什么是bank?

在这里插入图片描述

什么时候会发生bank conflict

在这里插入图片描述

按行存储,按列访问的时候,会发生bank conflict:
在这里插入图片描述

如何减缓bank conflict

在这里插入图片描述

代码实现
#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticConflictPadKernel(float *M_device, float *N_device, float *P_device, int width){/* 添加一个padding,可以防止bank conflict发生,结合图理解一下*/__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * BLOCKSIZE + threadIdx.x;int y = blockIdx.y * BLOCKSIZE + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素,我们只需要循环遍历width / tile_width 次就okay了,这里有点绕,画图理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {/* 这里为了实现bank conflict, 把tx与tx的顺序颠倒,同时索引也改变了*/M_deviceShared[tx][ty] = M_device[x * width + (m * BLOCKSIZE + ty)];N_deviceShared[tx][ty] = M_device[(m * BLOCKSIZE + tx)* width + y];__syncthreads();for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[tx][k] * N_deviceShared[k][ty];}__syncthreads();}/* 列优先 */P_device[x * width + y] = P_element;
}__global__ void MatmulSharedDynamicConflictPadKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 声明动态共享变量的时候需要加extern,同时需要是一维的 注意这里有个坑, 不能够像这样定义: __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因为在cuda中定义动态共享变量的话,无论定义多少个他们的地址都是一样的。所以如果想要像上面这样使用的话,需要用两个指针分别指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = (blockSize + 1) * blockSize;/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素,我们只需要循环遍历width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {/* 这里为了实现bank conflict, 把tx与tx的顺序颠倒,同时索引也改变了*/deviceShared[tx * (blockSize + 1) + ty]             = M_device[x * width + (m * blockSize + ty)];deviceShared[stride + (tx * (blockSize + 1) + ty)]  = N_device[(m * blockSize + tx) * width + y];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[tx * (blockSize + 1) + k] * deviceShared[stride + (k * (blockSize + 1 ) + ty)];}__syncthreads();}/* 列优先 */P_device[x * width + y] = P_element;
}

STREAM和EVENT

什么是stream

在这里插入图片描述

参考下述链接,理解cuda编程的一些基础概念:
理解CUDA中的thread,block,grid和warp
cuda stream的使用

多流编程实现

单流:
在这里插入图片描述
多流:
在这里插入图片描述
利用nsight systems进行分析:
在这里插入图片描述

如何利用多流进行隐藏访存和核函数执行延迟的调度

举一个栗子:
在这里插入图片描述

使用CUDA进行预处理/后处理

双线性插值

在这里插入图片描述

双线性插值的cuda实现

在这里插入图片描述

查看图片大小:
identity xx.png
可视化图片:
feh xx.png

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

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

相关文章

WordPress不支持中文TAG标签出现404的解决方法

我们在后台编辑文章时输入中文标签会发现出现404的情况&#xff0c;其实中文TAG标签链接无法打开的原因是WordPress不支持中文的编码。那么解决的方法也很容易&#xff0c;只要改代码让WordPress能支持中文的编码形式&#xff0c;也就是UTF-8和GBK编码即可&#xff0c;无需用到…

金融信贷公司所需的技术和风控体系及其带来的价值

金融信贷公司的技术架构通过集成传统大型机系统与现代数据平台&#xff0c;能够有效支持金融信贷业务的运作&#xff0c;同时通过大数据、ETL、报表开发、数据仓库等技术为公司带来更高效的数据驱动决策、精准的风控分析和更灵活的业务支持。 一、公司技术架构 数据仓库架构&…

《AI大模型应知应会100篇》第43篇:大模型幻觉问题的识别与缓解方法

第43篇&#xff1a;大模型幻觉问题的识别与缓解方法 摘要 当AI系统自信满满地编造"量子计算机使用香蕉皮作为能源"这类荒谬结论时&#xff0c;我们不得不正视大模型的幻觉问题。本文通过15个真实案例解析、6种检测算法实现和3套工业级解决方案&#xff0c;带您掌握…

计算方法实验五 插值多项式的求法

【实验性质】 综合性验 【实验目的】 掌握Lagrange插值算法、Newton插值算法&#xff1b;理解Newton插值算法相对于Lagrange插值算法的优点。 【实验内容】 先用C语言自带的系统函数sin x求出 的值&#xff0c;然后分别用Lagrange、Newton方法求出的值&#xff0c;并与用…

文献总结:TPAMI端到端自动驾驶综述——End-to-End Autonomous Driving: Challenges and Frontiers

端到端自动驾驶综述 1. 文章基本信息2. 背景介绍3. 端到端自动驾驶主要使用方法3. 1 模仿学习3.2 强化学习 4. 测试基准4.1 真实世界评估4.2 在线/闭环仿真测试4.3 离线/开环测试评价 5. 端到端自动驾驶面临的挑战5.1 多模态输入5.2 对视觉表征的依赖5.3 基于模型的强化学习的世…

PostgreSQL:pgAdmin 4 使用教程

pgAdmin 4 是一个用于管理和维护 PostgreSQL 数据库的强大工具。它提供了一个图形化界面&#xff0c;使用户能够轻松地连接到数据库、创建表、运行 SQL 语句以及执行其他数据库管理任务。 安装和使用 安装 pgAdmin 4 安装 pgAdmin 4 非常简单。下载并运行安装程序&#xff0…

Java学习手册:关系型数据库基础

一、关系型数据库概述 关系型数据库是一种基于关系模型的数据库&#xff0c;它将数据组织成一个或多个表&#xff08;或称为关系&#xff09;&#xff0c;每个表由行和列组成。每一列都有一个唯一的名字&#xff0c;称为属性&#xff0c;表中的每一行是一个元组&#xff0c;代…

wpf CommandParameter 传递MouseWheelEventArgs参数

在 WPF 中通过 CommandParameter 传递 MouseWheelEventArgs 参数时&#xff0c;需结合 ‌事件到命令的转换机制‌ 和 ‌参数转换器‌ 来实现。以下是具体实现方案及注意事项&#xff1a; 一、核心实现方法 1. ‌使用 EventToCommand 传递原始事件参数‌ 通过 Interaction.Tr…

八大排序之选择排序

本篇文章将带你详细了解八大基本排序中的选择排序 目录 &#xff08;一&#xff09;选择排序的时间复杂度和空间复杂度及稳定性分析 &#xff08;二&#xff09;代码实现 (三)输出结果 选择排序的基本原理是&#xff1a;每次从待排序的数组中找出最大值和最小值。具体流程是…

【算法学习】哈希表篇:哈希表的使用场景和使用方法

算法学习&#xff1a; https://blog.csdn.net/2301_80220607/category_12922080.html?spm1001.2014.3001.5482 前言&#xff1a; 在之前学习数据结构时我们就学习了哈希表的使用方法&#xff0c;这里我们主要是针对哈希表的做题方法进行讲解&#xff0c;都是leetcode上的经典…

Java 中如何实现自定义类加载器,应用场景是什么?

在 Java 中&#xff0c;可以通过继承 java.lang.ClassLoader 类来实现自定义类加载器。自定义类加载器可以控制类的加载方式&#xff0c;实现一些特殊的应用场景。 实现自定义类加载器的步骤&#xff1a; 继承 java.lang.ClassLoader 类。 重写 findClass(String name) 方法 …

信创开发中跨平台开发框架的选择与实践指南

&#x1f9d1; 博主简介&#xff1a;CSDN博客专家、CSDN平台优质创作者&#xff0c;高级开发工程师&#xff0c;数学专业&#xff0c;10年以上C/C, C#, Java等多种编程语言开发经验&#xff0c;拥有高级工程师证书&#xff1b;擅长C/C、C#等开发语言&#xff0c;熟悉Java常用开…

WebRTC 服务器之Janus架构分析

1. Webrtc三种类型通信架构 1.1 1 对 1 通信 1 对 1 通信模型设计的主要⽬标是尽量让两个终端进⾏直联&#xff0c;这样即可以节省服务器的资源&#xff0c;⼜可以提⾼ ⾳视频的服务质量。WebRTC ⾸先尝试两个终端之间是否可以通过 P2P 直接进⾏通信&#xff0c;如果⽆法直接…

数字化转型进阶:26页华为数字化转型实践分享【附全文阅读】

本文分享了华为数字化转型的实践经验和体会。华为通过数字化变革,致力于在客户服务、供应链、产品管理等方面提高效率,并把数字世界带入每个组织,构建万物互联的智能世界。华为的数字化转型愿景是成为行业标杆,通过推进数字化战略、构建面向业务数字化转型的IT组织阵型、坚…

Hal库下备份寄存器

首先要确保有外部电源给VBAT供电 生成后应该会有这两个文件&#xff08;不知道为什么生成了好几次都没有&#xff0c;复制工程在试一次就有了&#xff09; 可以看到stm32f407有20个备份寄存器 读写函数 void HAL_RTCEx_BKUPWrite(RTC_HandleTypeDef *hrtc, uint32_t Backup…

使用 Vue3 + Webpack 和 Vue3 + Vite 实现微前端架构(基于 Qiankun)

在现代前端开发中&#xff0c;微前端架构逐渐成为一种流行的解决方案&#xff0c;尤其是在大型项目中。通过微前端&#xff0c;我们可以将一个复杂的单体应用拆分为多个独立的小型应用&#xff0c;每个子应用可以独立开发、部署和运行&#xff0c;同时共享主应用的基础设施。本…

【c++】【STL】list详解

目录 list的作用list的接口构造函数赋值运算符重载迭代器相关sizeemptyfrontbackassignpush_frontpop_frontpush_backpop_backinserteraseswapresizeclearspliceremoveremove_ifuniquemergesortreverse关系运算符重载&#xff08;非成员函数&#xff09; list的模拟实现结点类迭…

Redis持久化:

什么是Redis持久化&#xff1a; Redis 持久化是指将 Redis 内存中的数据保存到硬盘等持久化存储介质中&#xff0c;以便在 Redis 服务器重启或出现故障时能够恢复数据&#xff0c;保证数据的可靠性和持续性。Redis 提供了两种主要的持久化方式&#xff1a;RDB&#xff08;Redi…

VBA 64位API声明语句第009讲

跟我学VBA&#xff0c;我这里专注VBA, 授人以渔。我98年开始&#xff0c;从源码接触VBA已经20余年了&#xff0c;随着年龄的增长&#xff0c;越来越觉得有必要把这项技能传递给需要这项技术的职场人员。希望职场和数据打交道的朋友&#xff0c;都来学习VBA,利用VBA,起码可以提高…

在pycharm profession 2020.3将.py程序使用pyinstaller打包成exe

一、安装pyinstaller 在pycharm的项目的Terminal中运行pip3 install pyinstaller即可。 安装后在Terminal中输入pip3 list看一下是否成功 二、务必在在项目的Terminal中输入命令打包&#xff0c;命令如下&#xff1a; python3 -m PyInstaller --noconsole --onefile xxx.py …