自己做网站网页归档,全屏网站模版,自动下单网站开发,企业信息查询软件cuda教程目录
第一章 指针篇 第二章 CUDA原理篇 第三章 CUDA编译器环境配置篇 第四章 kernel函数基础篇 第五章 kernel索引(index)篇 第六章 kenel矩阵计算实战篇 第七章 kenel实战强化篇 第八章 CUDA内存应用与性能优化篇 第九章 CUDA原子(atomic)实战篇 第十章 CUDA流(strea…cuda教程目录
第一章 指针篇 第二章 CUDA原理篇 第三章 CUDA编译器环境配置篇 第四章 kernel函数基础篇 第五章 kernel索引(index)篇 第六章 kenel矩阵计算实战篇 第七章 kenel实战强化篇 第八章 CUDA内存应用与性能优化篇 第九章 CUDA原子(atomic)实战篇 第十章 CUDA流(stream)实战篇 第十一章 CUDA的NMS算子实战篇 第十二章 YOLO的部署实战篇 第十三章 基于CUDA的YOLO部署实战篇
cuda教程背景
随着人工智能的发展与人才的内卷很多企业已将深度学习算法的C部署能力作为基本技能之一。面对诸多arm相关且资源有限的设备往往想更好的提速满足更高时效性必将更多类似矩阵相关运算交给CUDA处理。同时面对市场诸多教程与诸多博客岑子不起的教程或高昂教程费用使读者(特别是小白)容易迷糊无法快速入手CUDA编程实现工程化。 因此我将结合我的工程实战经验我将在本专栏实现CUDA系列教程帮助读者(或小白)实现CUDA工程化掌握CUDA编程能力。学习我的教程专栏你将绝对能实现CUDA工程化完全从环境安装到CUDA核函数编程从核函数到使用相关内存优化从内存优化到深度学习算子开发(如nms)从算子优化到模型(以yolo系列为基准)部署。最重要的是我的教程将简单明了直切主题CUDA理论与实战实例应用并附相关代码可直接上手实战。我的想法是掌握必要CUDA相关理论去除非必须繁杂理论实现CUDA算法应用开发待进一步提高将进一步理解更高深理论。
cuda教程内容
第一章到第三章探索指针在cuda函数中的作用与cuda相关原理及环境配置
第四章初步探索cuda相关函数编写(global、device、__host__等)实现简单入门
第五章探索不同grid与block配置如何计算kernel函数的index以便后续通过index实现各种运算
第六、七章由浅入深探索核函数矩阵计算深入探索grid、block与thread索引对kernel函数编写作用与影响并实战多个应用列子(如kernel函数实现图像颜色空间转换)
第八章探索cuda内存纹理内存、常量内存、全局内存等分配机制与内存实战应用(附代码)通过不同内存的使用来优化cuda计算性能
第九章探索cuda原子(atomic)相关操作并实战应用(如:获得某些自加索引等)
第十章探索cuda流stream相关应用并给出相关实战列子(如:多流操作等)
第十一到十三章探索基于tensorrt部署yolo算法我们首先将给出通用tensorrt的yolo算法部署该部署的前后处理基于C语言的host端实现然后给出基于cuda的前后处理的算子核函数编写最后数据无需在gpu与host间复制操作实现gpu处理提升算法性能。
目前以上为我们的cuda教学全部内容若后续读者有想了解知识可留言我们将根据实际情况更新相关教学内容。
大神忽略
源码链接地址点击这里 文章目录 cuda教程目录cuda教程背景cuda教程内容源码链接地址[点击这里](https://github.com/tangjunjun966/cuda-tutorial-master) 前言一、内存知识回顾二、GPU内存信息查询三、可分页内存与页锁定内存四、cudaMallocHost 和 cudaMalloc(可分页内存与页锁定内存)1、内存分配方式Host端内存分配(Pageable Memory) 2、分配的内存类型3、内存的使用方式4、内存的传输方式5、性能6、总结 八、总结 前言
以上章节中我们已经比较熟练掌握如何使用cuda编写自己想要的计算逻辑已能成功编写cuda代码了。 那么另外一个重要问题值得我们关注如何优化其性能使其工程部署能加速运行了。而这种性能优化与cuda内存密切相关。为此我们在本节中介绍cuda内存相关内容并附其源码。 一、内存知识回顾
我再次简单回顾下相关内存概念详细内容可看我第二章内容(我个人局的还是重要)链接点击这里。
Registers寄存器是GPU中最快的memorykernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话就是内置类型数组也是放在寄存器中。 寄存器变量是每个线程私有的一旦thread执行结束寄存器变量就会失效。
Shared Memory用__shared__修饰符修饰的变量存放在shared memory中。Shared Memory位于GPU芯片上访问延迟仅次于寄存器。所有Thread来进行访问的可以实现Block内的线程间的低开销通信。 要使用__syncthread()同步。
Local Memory本身在硬件中没有特定的存储单元而是从Global Memory虚拟出来的地址空间。是为寄存器无法满足存储需求的情况而设计的主要是用于存放单线程的大型数组和变量。Local Memory是线程私有的线程之间是不可见的。它的访问是比较慢的跟Global Memory的访问速度是接近的。使用情景无法确定其索引是否为常量的数组会消耗太多寄存器空间的大型结构或数组如果内核使用了多于寄存器的任何变量这也称为寄存器溢出
Constant Memory固定内存空间驻留在设备内存中并缓存在固定缓存中constant cache范围是全局的针对所有kernel kernel只能从constant Memory中读取数据因此其初始化必须在host端使用下面的function调用cudaError_t cudaMemcpyToSymbol(const void* symbol,const void* src,size_t count); 当一个warp中所有线程都从同一个Memory地址读取数据时constant Memory表现会非常好会触发广播机制。
Global MemoryGlobal Memory在某种意义上等同于GPU显存kernel函数通过Global Memory来读写显存。Global Memory是kernel函数输入数据和写入结果的唯一来源。
Texture Memory是GPU的重要特性之一也是GPU编程优化的关键。Texture Memory实际上也是Global Memory的一部分但是它有自己专用的只读cache。这个cache在浮点运算很有用Texture Memory是针对2D空间局部性的优化策略所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。从读取性能的角度跟Constant Memory类似。
Host Memory主机端存储器主要是内存可以分为两类可分页内存Pageable和页面 Page-Locked 或 Pinned内存。可分页内存通过操作系统 API(malloc/free) 分配存储器空间该内存是可以换页的即内存页可以被置换到磁盘中。可分页内存是不可用使用DMADirect Memory Acess)来进行访问的普通的C程序使用的内存就是这个内存。
二、GPU内存信息查询
代码如下
int inquire_GPU_info() {int deviceCount;cudaGetDeviceCount(deviceCount);int dev;for (dev 0; dev deviceCount; dev){int driver_version(0), runtime_version(0);cudaDeviceProp deviceProp;cudaGetDeviceProperties(deviceProp, dev);if (dev 0)if (deviceProp.minor 9999 deviceProp.major 9999)printf(\n);printf(\nDevice%d:\%s\\n, dev, deviceProp.name);cudaDriverGetVersion(driver_version);printf(CUDA驱动版本: %d.%d\n, driver_version / 1000, (driver_version % 1000) / 10);cudaRuntimeGetVersion(runtime_version);printf(CUDA运行时版本: %d.%d\n, runtime_version / 1000, (runtime_version % 1000) / 10);printf(设备计算能力: %d.%d\n, deviceProp.major, deviceProp.minor);printf(设备全局内存总量 Global Memory: %u M\n, deviceProp.totalGlobalMem/(1024*1024));printf(Number of SMs: %d\n, deviceProp.multiProcessorCount);printf(常量内存 Constant Memory: %u K\n, deviceProp.totalConstMem/1024);printf(每个block的共享内存 Shared Memory: %u K\n, deviceProp.sharedMemPerBlock/1024);printf(每个block的寄存器 registers : %d\n, deviceProp.regsPerBlock);printf(线程束Warp size: %d\n, deviceProp.warpSize);printf(每个SM的最大线程数 threads per SM: %d\n, deviceProp.maxThreadsPerMultiProcessor);printf(每个block的最大线程数 threads per block: %d\n, deviceProp.maxThreadsPerBlock);printf(每个block的最大维度 each dimension of a block: %d x %d x %d\n, deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]);printf(每个grid的最大维度 dimension of a grid: %d x %d x %d\n, deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);printf(Maximum memory pitch: %u bytes\n, deviceProp.memPitch);printf(Texture alignmemt: %u bytes\n, deviceProp.texturePitchAlignment);printf(Clock rate: %.2f GHz\n, deviceProp.clockRate * 1e-6f);printf(Memory Clock rate: %.0f MHz\n, deviceProp.memoryClockRate * 1e-3f);printf(Memory Bus Width: %d-bit\n, deviceProp.memoryBusWidth);}return 0;
}
}查询结果显示如下
三、可分页内存与页锁定内存
CPU内存称之为Host Memory逻辑上可分为Pageable Memory(可分页内存)、Page Lock Memory(页锁定内存)Page Lock Memory又称为Pinned Memory从字面意思上而言Page Lock Memory是锁定的内存一旦申请后就专供申请者使用Pageable Memory则没有锁定特性申请后可能会被交换。
总结如下 ①、pinned memory具有锁定特性是稳定不会被交换的 pageable memory没有锁定特性对于第三方设备比如GPU去访问时因为无法感知内存是否被交换可能得不到正确的数据 ②、pageable memory的性能比pinned memory差很可能降低你程序的优先级然后把内存交换给别人用 pageable memory策略使用内存假象实际8GB但是可以使用15GB可以提高程序运行数量但运行速度会降低 pinned memory太多会导致操作系统整体性能降低因为程序运行数量减少了 ③、GPU可以直接访问pinned memory而不能访问pageable memory因为第二条。 说明当将pageable host Memory数据送到device时CUDA驱动会首先分配一个临时的page-locked或者pinned host Memory并将host的数据放到这个临时空间里。然后GPU从这个所谓的pinned Memory中获取数据如下图所示 四、cudaMallocHost 和 cudaMalloc(可分页内存与页锁定内存)
之前章节一直以实例介绍cuda代码编写也对host与device端的变量进行了内存分配并未重点说明cudaMallocHost 和 cudaMalloc的使用方法我个人觉得很重要对于不同设备的数据传输(如host与GPU间)均需要使用复制方法而针对GPU内存分配与CPU数据间关系需要我们有更深入了解在此我介绍重点介绍一下。
1、内存分配方式
以上介绍gpu如何访问cpu的内存方式。对于给GPU访问而言距离计算单元越近内存访问效率越高。为此由低到高访问速度为Pinned Memory Global Memory Shared Memory。 重点说明GPU可以直接访问Pinned Memory称之为DMA Direct Memory Access
接下来我将介绍实际内存分配的几种方式
Host端内存分配(Pageable Memory)
之前代码使用cudaMallocHost对内存分配但也可使用new或malloc分配而该分配属于Pageable Memory可分页内存。 其分配内存代码如下 std::cout 设置new(malloc)可分页内存 std::endl;float* memory_device nullptr;float* memory_host new float[100]; // Pageable Memoryfor (int i 0; i 100; i) { memory_host[i] i * 100; }checkRuntime(cudaMemcpy(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice)); // 返回的地址是开辟的device地址存放在memory_deviceshow_value dim3(1), dim3(100) (memory_device);以上直接使用CPU分配内存然后使用cudaMemcpy复制给memory_device中仍然可以实现然这种效率较低。但切记这种memory_host可以直接使用new赋值在核函数中使用。同时这种速度较慢不建议使用。 预测结果显示(仅显示前10个数)如下 2、分配的内存类型
cudaMallocHost 分配的内存是页锁定内存而 cudaMalloc 分配的内存是普通可分页内存。
3、内存的使用方式
cudaMallocHost 分配的内存可以通过主机和设备访问而 cudaMalloc 分配的内存只能通过设备访问。 cudaMalloc分配内存方式为GPU的全局内存代码如下 std::cout 设置全局内存 std::endl;float* memory_device nullptr; // Global MemorycheckRuntime(cudaMalloc(memory_device, 100 * sizeof(float))); // pointer to device以上代码实际在前面章节中已大量使用实际作用为使用cudaMalloc在gpu设备上分配一个全局内存空间便于在kernel计算中存储数据。 cudaMallocHost分配内存方式代码如下 std::cout 设置页锁定内存 std::endl;float* memory_device nullptr;float* memory_page_locked nullptr; // Pinned MemorycheckRuntime(cudaMallocHost(memory_page_locked, 100 * sizeof(float))); // 返回的地址是被开辟的pin memory的地址存放在memory_page_lockedcheckRuntime(cudaMemcpy(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost)); // 将其返回host内存
以上代码实际在前面章节中已大量使用实际作用也就是上面解释即使用cudaMallocHost在gpu设备上分配内存可使主机host和设备device均可访问并使用cudaMemcpy赋值gpu数据。
4、内存的传输方式
由于 cudaMallocHost 分配的内存可以通过主机和设备访问因此可以通过零拷贝技术Zero-Copy将数据直接从主机内存传输到设备内存而 cudaMalloc 分配的内存则需要使用显式的数据传输函数如 cudaMemcpy进行传输。
5、性能
由于 cudaMallocHost 分配的内存是页锁定内存因此可以避免在主机和设备之间进行数据传输时产生额外的复制操作从而提高数据传输的性能。
6、总结
尽量多用Pinned Memory储存host端数据或者显式处理Host到Device时用PinnedMemory做缓存都是提高性能的关键。因此如果需要在主机和设备之间进行频繁的数据传输建议使用 cudaMallocHost 分配内存。如果只需要在设备上进行计算并且不需要频繁地与主机进行数据交换则可以使用 cudaMalloc 分配内存。
八、总结
以上为cuda内存应用与性能优化篇关于内存的相关知识和简单代码说明旨在掌握在host端与device端间内存传输与分配相关原理后面篇章将介绍如何定义常量内存、共享内存、纹理内存等相关方法并附上相应代码。
