深度学习模型部署(十二)CUDA编程-绪

请添加图片描述

CUDA 运行时 API 与 CUDA 驱动 API 速度没有差别,实际中使用运行时 API 较多,运行时 API 是在驱动 API 上的一层封装。​

CUDA 是什么?​

CUDA(Compute Unified Device Architecture) 是 nvidia 推出的一个通用并行技术架构,用它来进行 GPU 编程。CUDA 本身并不是一门语言,而是一个 GPU 编程模型,是对 C++,Python 这种常见 CPU 编程语言的一个补充。​

为什么用 GPU 编程要用 CUDA?​

因为 GPU 的控制硬件少,所以编程模型要求非常严格,最早期 GPU 的唯一交互方式是通过 OpenGL 和 DirectX 这些图形 API,​

基本上所有的编程语言都是在 CPU 上运行的,所以催生了 CUDA 编程框架,为了更方便的与 GPU 交互和调用 GPU 的资源.​

请添加图片描述

CUDA编译运行流程

1、CUDA代码文件后缀.cu,使用nvcc进行编译
nvcc是nvidia基于LLVM开发的专门用于编译cuda代码的编译器,cuda代码有一套完整的工具链,称为cuda-toolkit,包括nvcc编译器,cuda-gdb调试工具等。
2、CUDA中的代码执行设备有两种,一种是device,一种是host,CPU被称为host,普通代码都是在host上执行,GPU被称为device,在device上执行的代码需要添加__global__ 或者__device__前缀
具体CUDA不同前缀意义见
3、CUDA中device使用的是显存,所以在device上执行的函数只能传入device上定义的变量,具体方法为定义好变量,然后使用cudaMalloc函数在显存中给变量分配空间,再使用cudaMemcpy将变量拷贝到显存中(cudaMemcpy是内存拷贝函数,可以根据给的参数将device中的memory拷贝到host中,也可以反过来拷贝),使用完后要使用cudaFree进行释放
4、通过__global__前缀定义的函数执行时需要设置执行的block数量和线程数
5、在host函数中进行完cuda调用后,需要使用cudaDeviceSynchronize()函数,因为启动内核是一个异步操作,只要发布了内核启动命令,不等内核执行完成,控制权就会立刻返回给调用内核的CPU线程。

CUDA kernel和线程管理

在CUDA中,函数称为kernel,每个kernel都有一个前缀,不同的前缀代表了kernel不同的运行要求。

#include<stdio.h>
//在GPU上执行的kernel中不允许使用C++的标准库 iostream
// warning #20096-D: address of a host variable "std::cout" cannot be directly taken in a device function
__global__ void hello_from_gpu()
{printf("hello world from gpu\n");//compute capacity 2.0以后才支持printf,也就是GeForce 830M以后的GPU
}int main(void)
{hello_from_gpu<<<4,4>>>();cudaDeviceSynchronize();return 0;
}

<<<>>>是cuda调用kernel时的语法,<<<numBlocks, threadsPerBlock>>>意思为调用numBlocks个block,每个block中threadsPerBlock个线程。

函数前缀

  1. device
    使用 device 限定符声明的函数具有以下特征:
  • 在设备上执行;
  • 仅可通过设备调用。
  1. global
    使用 global 限定符可将函数声明为内核。此类函数:
  • 在设备上执行;
  • 仅可通过主机调用。
  1. host
    使用 host 限定符声明的函数具有以下特征:
  • 在主机上执行;
  • 仅可通过主机调用。
    仅使用 host 限定符声明函数等同于不使用限定符声明函数,这两种情况下,函数都将仅为主机进行编译。
    函数前缀的一些限制
    deviceglobal 函数不支持递归。
    deviceglobal 函数的函数体内无法声明静态变量。
    deviceglobal 函数不得有数量可变的参数。
    device 函数的地址无法获取,但支持 global 函数的函数指针。
    globalhost 限定符无法一起使用。
    global 函数的返回类型必须为空。
    global 函数的任何调用都必须按规定指定其执行配置。
    global 函数的调用是异步的,也就是说它会在设备执行完成之前返回。
    global 函数参数将同时通过共享存储器传递给设备,且限制为 256 字节。
    线程管理
    CUDA有一套专门用于线程管理的机制,一个kernel调用时的配置<<<numBlocks, threadsPerBlock>>>中的numBlocks和threadPerBlock可以是int变量,也可以是线程管理的结构体,结构体更为常见。
#include<stdio.h>
#include<cuda_runtime.h>__global__ void build_in_variables(void)
{// build-in variables// blockDim:等同于threadsPerBlock// gridDim:等同于numBlocks// blockIdx:一个block在grid中的id// threadIdx:一个thread在block中的idconst int blockId = blockIdx.x + blockIdx.y * gridDim.x;const int threadId = threadIdx.x + blockDim.x * threadIdx.y;printf("blockIdx=(%d,%d) \n",blockIdx.x,blockIdx.y);printf("threadIdx=(%d,%d) \n",threadIdx.x,threadIdx.y);printf("blockid=:%d,threadId=%d \n",blockId,threadId);}int main(void)
{printf("*****device message*******\n");int dev=0;cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp,dev);printf("Using Device %d:%s\n",dev,deviceProp.name);printf("Device %d has compute capability %d.%d.\n",dev,deviceProp.major,deviceProp.minor);printf("Device %d has %d multi-processors.\n",dev,deviceProp.multiProcessorCount);printf("Device %d has %zu byte total global memory.\n",dev,deviceProp.totalGlobalMem);printf("Device %d has %zu byte total constant memory.\n",dev,deviceProp.totalConstMem);printf("Device %d has %zu byte shared memory per block.\n",dev,deviceProp.sharedMemPerBlock);printf("Device %d has %d total registers per block.\n",dev,deviceProp.regsPerBlock);printf("Device %d has %d max threads per block.\n",dev,deviceProp.maxThreadsPerBlock);printf("Device %d has %d max threads dimensions.\n",dev,deviceProp.maxThreadsDim[0]);printf("Device %d has %u max grid size.\n",dev,deviceProp.maxGridSize[0]);printf("Device %d has %d warp size.\n",dev,deviceProp.warpSize);printf("Device %d has %d clock rate.\n",dev,deviceProp.clockRate);printf("Device %d has %d max threads per multi-processor.\n",dev,deviceProp.maxThreadsPerMultiProcessor);dim3 numBlocks(2,2);// 2*2个block per grid// dim3,是一个包含xyz三个无符号整型数的结构体,默认值为1//三个维度,x变化最快,然后是y,最后是zdim3 threadsPerBlock(2,2);// 2*2个thread per blockbuild_in_variables<<<numBlocks, threadsPerBlock>>>();cudaDeviceReset();return 0;
}

在这里插入图片描述

在计算能力9.0以前的架构,thread的Hierarchy是二维的,只有两个层次,一个grid,一个block。在9.0以后的架构,thread的Hierarchy是三维的,新引入了一个可选层次:Cluster集群,每个Cluster中的block可以确保是在同一个GPC(GPU Processing Cluster)GPU集群中运行的。

#include<stdio.h>__global__ void __cluster_dims__(2,1,1) hello_from_gpu()
{printf("Hello World from GPU!\n");
}int main()
{hello_from_gpu<<<1,1>>>();cudaDeviceSynchronize();return 0;
}

内存管理

内存分为device memory和host memory,二者之间通过cudaMemcpy来进行管理。
[图片]

GPU每个线程有自己单独的寄存器和内存,同一个block中有所有thread都能访问的shared memory,
在有cluster的架构中,同一个cluster中的block的shared memory组成了distributed shared memory,可以相互访问。除此之外还有专门的只读内存,用于存放texture(纹理)

#include <stdio.h>__global__ void sharedMemoryExample(int* input)
{// Define shared memory array__shared__ int sharedArray[256];// Get the thread indexint tid = threadIdx.x;// Load data from global memory to shared memorysharedArray[tid] = input[tid];// Synchronize threads to ensure all data is loaded__syncthreads();// Perform some computation using shared memory datasharedArray[tid] = sharedArray[tid] * 2;// Synchronize threads again before writing back to global memory__syncthreads();// Write the result back to global memoryinput[tid] = sharedArray[tid];
}int main()
{// Define input dataint input[256];// Initialize input datafor (int i = 0; i < 256; i++){input[i] = i;}// Allocate memory on the GPUint* d_input;cudaMalloc((void**)&d_input, sizeof(int) * 256);// Copy input data from host to devicecudaMemcpy(d_input, input, sizeof(int) * 256, cudaMemcpyHostToDevice);// Launch the kernelsharedMemoryExample<<<1, 256>>>(d_input);// Copy the result back from device to hostcudaMemcpy(input, d_input, sizeof(int) * 256, cudaMemcpyDeviceToHost);// Print the resultfor (int i = 0; i < 256; i++){printf("%d ", input[i]);}// Free memory on the GPUcudaFree(d_input);return 0;
}

变量前缀:

1.device
device 限定符声明位于设备上的变量。
在接下来的三节中介绍的其他类型限定符中,最多只能有一种可与 device 限定符一起使用,以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符,则变量具有以下特征:

  • 位于全局存储器空间中;
  • 与应用程序具有相同的生命周期;
    可通过网格内的所有线程访问,也可通过运行时库从主机访问。
    2.constant
    constant 限定符可选择与 device 限定符一起使用,所声明的变量具有以下特征:
  • 位于固定存储器空间中;
  • 与应用程序具有相同的生命周期;
    可通过网格内的所有线程访问,也可通过运行时库从主机访问。
    3.shared
    shared 限定符可选择与 device 限定符一起使用,所声明的变量具有以下特征:
  • 位于线程块的共享存储器空间中;
  • 与块具有相同的生命周期;
  • 尽可通过块内的所有线程访问。
    只有在 syncthreads()(参见第 4.4.2 节)的执行写入之后,才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量,否则只要之前的语句完成,编译器即可随意优化共享存储器的读写操作。
    限制:
    不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。
    shared 和 constant 变量具有隐含的静态存储。
    devicesharedconstant 变量无法使用 extern 关键字定义为外部变量。
    deviceconstant 变量仅允许在文件作用域内使用。
    不可为设备或从设备指派 constant 变量,仅可通过主机运行时函数从主机指派(参见第 4.5.2.3 节和第 4.5.3.6 节)。
    shared 变量的声明中不可包含初始化。
    访问速度顺序为:register>shared>constant>local>device
    下面是具体的一个应用:
    将共享存储器中的变量声明为外部数组时,例如:
extern __shared__ float shared[];

数组的大小将在启动时确定(参见第 4.2.3 节)。所有变量均以这种形式声明,在存储器中的同一地址开始,因此数组中的变量布局必须通过偏移显式管理。
例如:

// 如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容:
// short array0[128];
// float array1[64];
// int array2[256];
// 则应通过以下方法声明和初始化数组:
extern __shared__ char array[];
__device__ void func() // device or global function
{short* array0 = (short*)array;float* array1 = (float*)&array0[128];int* array2 = (int*)&array1[64];
}

在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下,编译器可能选择将其置于本地存储器中。
只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间,此类指针即受支持,否则将仅限于指向在全局存储器空间中分配或声明的存储器。
通过获取 devicesharedconstant 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress() 获取的 deviceconstant 变量的地址仅可在主机代码中使用。

对 global 函数进行配置

global 函数的任何调用都必须指定该调用的执行配置。
执行配置定义将用于在该设备上执行函数的网格和块的维度,以及相关的流。可通过在函数名称和括号参数列表之间插入 <<<Dg, Db, Ns, s>>> 形式的表达式来指定,其中:
Dg 的类型为 dim3,指定网格的维度和大小,Dg.x * Dg.y 等于所启动的块数量,Dg.z 无用;
Db 的类型为 dim3,指定各块的维度和大小,Db.x * Db.y * Db.z 等于各块的线程数量;
Ns 的类型为 size_t,指定各块为此调用动态分配的共享存储器(除静态分配的存储器之外),这些动态分配的存储器可供声明为外部数组的其他任何变量使用,Ns 是一个可选参数,默认值为 0;
S 的类型为 cudaStream_t,指定相关流;S 是一个可选参数,默认值为 0。
举例来说,

//一个函数的声明如下:
__global__ void Func(float* parameter);
//必须通过如下方法来调用此函数:
Func<<<Dg, Db, Ns>>>(parameter);

执行配置的参数将在实际函数参数之前被评估,与函数参数相同,通过共享存储器同时传递给设备。

如果 Dg 或 Db 大于设备允许的最大大小,或 Ns 大于设备上可用的共享存储器最大值,或者小于静态分配、函数参数和执行配置所需的共享存储器数量,则函数将失败。

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

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

相关文章

【蓝桥杯】填空题技巧|巧用编译器|用Python处理大数和字符|心算手数|思维题

目录 一、填空题 1.巧用编译器 2.巧用Excel 3. 用Python处理大数 4.用Python处理字符 5.心算手数 二、思维题 推荐 前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。【点击跳转到网站】 一、填空题 …

【考研数学】张宇全程学习包

可以全程张宇老师的高等数学&#xff0c;张宇老师的拿手绝活是高数 但是其他科目&#xff0c;还有更好的选择&#xff0c;比如线性代数&#xff0c;汤家凤老师还有李永乐老师讲的都不错&#xff0c;概率论&#xff0c;余丙森老师还有方浩老师讲的很好。下面我就讲清楚&#xf…

AI 视频 | 火爆全网的真人转动漫工具 DomoAI,又上新功能了!(三)

DomoAI 又又又上线新功能了&#xff01; 上传一张静态人像图片 一个人像动作视频&#xff0c;就可以生成两者融合的动态视频。 啥都不说&#xff0c;直接看官方的示例视频&#xff1a; DomoAI 新功能 move 官方示例视频 使用非常简单&#xff0c;在 Discord 中通过 /move 指…

C++ 动态规划

文章目录 一、简介二、举个栗子2.1斐波那契数列2.2最短路径&#xff08;DFS&#xff09; 参考资料 一、简介 感觉动态规划非常的实用&#xff0c;因此这里整理一下相关资料。动态规划&#xff08;Dynamic Programming&#xff09;&#xff1a;简称 DP&#xff0c;是一种优化算法…

【JavaEE初阶系列】——多线程案例一——单例模式 (“饿汉模式“和“懒汉模式“以及解决线程安全问题)

目录 &#x1f6a9;单例模式 &#x1f388;饿汉模式 &#x1f388;懒汉模式 ❗线程安全问题 &#x1f4dd;加锁 &#x1f4dd;执行效率提高 &#x1f4dd;指令重排序 &#x1f36d;总结 单例模式&#xff0c;非常经典的设计模式&#xff0c;也是一个重要的学科&#x…

摆扫式(whisk broom)和推扫式(push broom)卫星传感器介绍

目前&#xff0c;我们卫星传感器主要有两大类型&#xff1a;摆扫式&#xff08;whisk broom&#xff09;和推扫式&#xff08;push broom&#xff09;。为了更好的理解和使用卫星影像数据&#xff0c;我们需要简单了解下这两种传感器工作原理。 摆扫式&#xff1a;Whisk Broom…

搭建Hadoop HA

目录 前言 搭建前准备 搭建 前言 Hadoop是一个由Apache基金会所开发的分布式系统基础架构&#xff0c;它允许用户在不了解分布式底层细节的情况下开发分布式程序&#xff0c;充分利用集群的威力进行高速运算和存储。Hadoop主要解决大数据存储和大数据分析两大核心问题&…

Phoenix概念篇

文章目录 前言Phoenix的web层概念PlugEndpointRouterScopePipeline ControllerAction Component 一次请求 前言 Elixir和Phoenix的作者也是Rails社区的核心开发者&#xff0c;如果是之前接触过Ruby on Rails的开发者&#xff0c;对Phoenix也许不会感到太陌生。笔者没有接触过R…

【报错】使用gradio渲染html页面无法加载本地图片

【报错】使用gradio渲染html页面无法加载本地图片 【报错】使用gradio渲染html页面无法加载本地图片[HTML] how to load local image by html output #884成功解决 【报错】使用gradio渲染html页面无法加载本地图片 在使用gradio框架渲染html页面&#xff0c;使用绝对路径&quo…

BUUCTF-Misc14

[WUSTCTF2020]find_me1 1.打开附件 是一个学校的校徽 2.盲文解密 发现图片属性里的备注是一串盲文 用在线盲文解密 3.得到flag

C语言笔记:重学输入和输出

ACM金牌带你零基础直达C语言精通-课程资料 本笔记属于船说系列课程之一&#xff0c;课程链接&#xff1a;ACM金牌带你零基础直达C语言精通https://www.bilibili.com/cheese/play/ep159068?csourceprivate_space_class_null&spm_id_from333.999.0.0 你也可以选择购买『船说…

AI新工具 视频迁移升级中国水墨画风格2.0;新颖的视频编辑框架提示编辑,风格转移,身份操控都不在话下;提取多种风格人脸草图

✨ 1: DomoAI 升级中国水墨画风格2.0 DomoAI是一个多功能的AI视频处理工具&#xff0c;可以将视频转换成多种风格&#xff0c;包括日本动漫、3D卡通、漫画和像素风格等。用户只需上传原始视频&#xff0c;通过简单的操作就能实现风格转换&#xff0c;制作出具有个性的高质量视…

“架构(Architecture)” 一词的定义演变历史(依据国际标准)

深入理解“架构”的客观含义&#xff0c;不仅能使IT行业的系统架构设计师提升思想境界&#xff0c;对每一个积极的社会行动者而言&#xff0c;也具有长远的现实意义&#xff0c;因为&#xff0c;“架构”一词&#xff0c;不只限于IT系统&#xff0c;而是指各类系统(包括社会系统…

蓝鹏智能测量仪应用于这些方面!助力发展新质生产力!

新质生产力是未来几年着重发展的方向&#xff0c;关于如何实现产业化升级&#xff0c;各厂家会在自身的基础上进行产业化调整升级&#xff0c;利用新工具、新手段&#xff0c;大幅缩短研发设计周期&#xff0c;从而让产品迭代速度不断加快&#xff1b;提升产品品质&#xff0c;…

堆排序(六大排序)

前面博客已经分享过堆的知识了&#xff0c;今天我们来分享堆排序。 堆排序 堆排序(Heapsort)是指利用堆积树&#xff08;堆&#xff09;这种数据结构所设计的一种排序算法&#xff0c;它是选择排序的一种。它是通过堆来进行选择数据。 ★★★需要注意的是排升序要建大堆&#…

3、创建项目,什么是路由

一、创建项目 第一次全局安装脚手架 npm install -g vue/clivue create 项目名 二、什么是路由&#xff1f; 路由就是一组 key-value 的对应关系多个路由&#xff0c;需要经过路由器的管理 1、后端路由&#xff1a; 每个url地址都对应着不同的静态资源对于普通的网站。所有…

24计算机考研调剂 | 【官方】湘潭大学

湘潭大学 考研调剂要求 招生专业&#xff1a; 调剂基本要求&#xff1a; &#xff08;1&#xff09;基本要求同《湘潭大学2024年硕士研究生复试录取工作方案》。 &#xff08;2&#xff09;初试成绩要求&#xff1a; 初试成绩各单科均须达到A类考生进入复试的初试成绩基本要…

007 日期类型相关工具类

推荐一篇文章 http://t.csdnimg.cn/72F7Jhttp://t.csdnimg.cn/72F7J

golang+vue微服务电商系统

golangvue微服务电商系统 文章目录 golangvue微服务电商系统一、项目前置准备二、项目简介三、代码GItee地址 golang、vue redis、mysql、gin、nacos、es、kibana、jwt 一、项目前置准备 环境的搭建 官方go开发工程师参考地址&#xff1a;https://blog.csdn.net/qq23001186/cat…

刷题记录:最长公共前缀

编写一个函数来查找字符串数组中的最长公共前缀。 如果不存在公共前缀&#xff0c;返回空字符串 ""。 示例 1&#xff1a; 输入&#xff1a;strs ["flower","flow","flight"] 输出&#xff1a;"fl"示例 2&#xff1a; 输…