当前位置: 首页 > news >正文

20220831-basic-cuda-interface

layouttitlecategoriestagsdatedescription
post
CUDA并行编程学习笔记
blog
GPU编程
2022-08-31 12:39:39 -0700
CUDA并行编程指南

名词

  • SIMD: 单指令多数据,是基于一个处理器核的,128位
  • MMX:多媒体拓展
  • AVX 高级适量拓展, 256位

计算机架构

冯诺依曼计算机架构

  • 内存受限型
  • QPI (quick path interconnect) 快速通道互联

连接机

采用4096个16核的CPU组装到一台机器上,也就是说64K个处理器来完成一个任务。连接机采用SIMD型并行处理,但是处理器之间的同步和通讯是很大的问题

Cell处理器(众核)

用一个常规处理器作为监管处理器(PowerPC),该处理器与大量高速流处理(SPE)相连。

  • 每个流处理单元SPE调用执行一个程序
  • 通过共享的网络,SPE之间和SPE与PowerPC之间进行相互通讯

国产申威 26010 处理器架构图

多点计算

集群,当前最流行的莫过于Hadoop和spark了,一个是分布式文件系统,一个是分布式计算框架,这两个工具使得多点计算的方法充分发挥。

GPU架构

CUDA编程基础知识

学习CUDA C,可以在异构计算平台中实现高性能的应用。CUD的编译原则--基于虚拟指令集的运行时编译。

计算能力—高性能硬件与技术

GPU在高性能计算和深度学习加速中扮演着非常重要的角色, GPU的强大的并行计算能力,大大提升了运算性能。随着运算数据量的不断攀升,GPU间需要大量的交换数据,GPU通信性能成为了非常重要的指标。NVIDIA推出的GPUDirect就是一组提升GPU通信性能的技术。但GPUDirect受限于PCI Expresss总线协议以及拓扑结构的一些限制,无法做到更高的带宽,为了解决这个问题,NVIDIA提出了NVLink总线协议。

GPUDirect P2P

GPUDirect Peer-to-Peer(P2P) 技术主要用于单机GPU间的高速通信,它使得GPU可以通过PCI Express直接访问目标GPU的显存,避免了通过拷贝到CPU host memory作为中转,大大降低了数据交换的延迟。 以深度学习应用为例,主流的开源深度学习框架如TensorFlow、MXNet都提供了对GPUDirect P2P的支持,NVIDIA开发的NCCL(NVIDIA Collective Communications Library)也提供了针对GPUDirect P2P的特别优化。 通过使用GPUDirect P2P技术可以大大提升深度学习应用单机多卡的扩展性,使得深度学习框架可以获得接近线性的训练性能加速比

  • 浅析GPU通信技术(上)--GPUDirect P2P

NVLink 拓扑结构图

首先我们简单看下NVIDIA对NVLink的介绍:NVLink能在多GPU之间和GPU与CPU之间实现非凡的连接带宽。带宽有多大?2016发布的P100是搭载NVLink的第一款产品,单个GPU具有160GB/s的带宽,相当于PCIe Gen3 * 16带宽的5倍。去年GTC 2017上发布的V100搭载的NVLink 2.0更是将GPU带宽提升到了300G/s,差不多是PCIe的10倍了。

NVLINK网络拓扑结构

  • NVIDIA NVLINK
  • 浅析GPU通信技术(中)-NVLink总线协议

RDMA原理介绍

前面介绍的GPUDirect P2P和NVLink技术可以大大提升GPU服务器单机的GPU通信性能,当前深度学习模型越来越复杂,计算数据量暴增,对于大规模深度学习训练任务,单机已经无法满足计算要求,多机多卡的分布式训练成为了必要的需求,这个时候多机间的通信成为了分布式训练性能的重要指标。

多机通讯RMDA架构图 如上图所示,传统的TCP/IP协议,应用程序需要要经过多层复杂的协议栈解析,才能获取到网卡中的数据包,而使用RDMA协议,应用程序可以直接旁路内核获取到网卡中的数据包。RDMA可以简单理解为利用相关的硬件和网络技术,服务器1的网卡可以直接读写服务器2的内存,最终达到高带宽、低延迟和低资源利用率的效果。 应用RMDA技术的应用拓扑图 所谓GPUDirect RDMA,就是计算机1的GPU可以直接访问计算机2的GPU内存。而在没有这项技术之前,GPU需要先将数据从GPU内存搬移到系统内存,然后再利用RDMA传输到计算机2,计算机2的GPU还要做一次数据从系统内存到GPU内存的搬移动作。GPUDirect RDMA技术使得进一步减少了GPU通信的数据复制次数,通信延迟进一步降低。

  • 浅析GPU通信技术(下)-GPUDirect RDMA

CUDA的基础入门

函数的类型

__host__ float HostFunc() 默认情况下,被host函数调用在CPU上执行

__devide__ float DeviceFunc() 被GPU设备执行调用

__global__ void Kernelfunc() 被host函数调用,在设备上执行

Note:
* __global__函数返回值必须为void
* 在设备上执行的函数不能是递归,函数参数是固定的,不能再函数内部使用static变量

变量类型

__shared__ A[4];//在share memory,块内线程共享。 设备上的函数,声明的变量都是存在register上的,存不下的放到local memory; cudaMalloc()的空间是在设备的global memory上的。

CUDA几个头文件

#include<cuda_runtime.h>  // cuda程序运行必须的头文件

CUDA routine

  1. cudaError_t err = cudaSuccess; cudaError_t类型,表示错误类型。cudaSuccess表示成功。一般cuda routine的返回值都是cudaError_t类型,表示函数是否执行成功。

  2. printf("%s\n", cudaGetErrorString(cudaGetLastError())); 输出错误时,使用以上函数转化为string。

  3. err = cudaMalloc((void **)&d_A, size); 动态内存申请函数,在设备的global memory上申请size个字节空间。

  4. err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);or err = cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost); //内存拷贝函数:从cpu上的内存h_A上拷贝size个字节数据到gpu上的内存d_A。反之,一样。

  5. int threadsPerBlock = 256; int blocksPerGrid =(nElements + threadsPerBlock - 1) / threadsPerBlock; vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, nElements); //前2句,表示Grid,block都是1维时,设置网格内的块数,每块内的线程数。 //最后一句,启动kernel(运行在gpu端的函数)函数。 //注意前2句可以改成。dim3 threadsPerBlock(256);这种形式。

  6. err = cudaGetLastError(); //启动kernel函数时,并没有返回值,通过这个调用这个函数,查看kernel函数是否启动成功。

  7. err = cudaFree(d_A); //释放使用cudaMalloc申请的空间。

  8. err = cudaMemset(d_a, 0, size) //类似于memset函数。将d_A的size个字节置0.

/**
 * CUDA device properties
 */
struct __device_builtin__ cudaDeviceProp
{
    char   name[256];                  /**< ASCII string identifying device */
    size_t totalGlobalMem;             /**< Global memory available on device in bytes */
    size_t sharedMemPerBlock;          /**< Shared memory available per block in bytes */
    int    regsPerBlock;               /**< 32-bit registers available per block */
    int    warpSize;                   /**< Warp size in threads */
    size_t memPitch;                   /**< Maximum pitch in bytes allowed by memory copies */
    int    maxThreadsPerBlock;         /**< Maximum number of threads per block */
    int    maxThreadsDim[3];           /**< Maximum size of each dimension of a block */
    int    maxGridSize[3];             /**< Maximum size of each dimension of a grid */
    int    clockRate;                  /**< Clock frequency in kilohertz */
    size_t totalConstMem;              /**< Constant memory available on device in bytes */
    int    major;                      /**< Major compute capability */
    int    minor;                      /**< Minor compute capability */
    size_t textureAlignment;           /**< Alignment requirement for textures */
    size_t texturePitchAlignment;      /**< Pitch alignment requirement for texture references bound to pitched memory */
    int    deviceOverlap;              /**< Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. */
    int    multiProcessorCount;        /**< Number of multiprocessors on device */
    int    kernelExecTimeoutEnabled;   /**< Specified whether there is a run time limit on kernels */
    int    integrated;                 /**< Device is integrated as opposed to discrete */
    int    canMapHostMemory;           /**< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer */
    int    computeMode;                /**< Compute mode (See ::cudaComputeMode) */
    int    maxTexture1D;               /**< Maximum 1D texture size */
    int    maxTexture1DMipmap;         /**< Maximum 1D mipmapped texture size */
    int    maxTexture1DLinear;         /**< Maximum size for 1D textures bound to linear memory */
    int    maxTexture2D[2];            /**< Maximum 2D texture dimensions */
    int    maxTexture2DMipmap[2];      /**< Maximum 2D mipmapped texture dimensions */
    int    maxTexture2DLinear[3];      /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */
    int    maxTexture2DGather[2];      /**< Maximum 2D texture dimensions if texture gather operations have to be performed */
    int    maxTexture3D[3];            /**< Maximum 3D texture dimensions */
    int    maxTexture3DAlt[3];         /**< Maximum alternate 3D texture dimensions */
    int    maxTextureCubemap;          /**< Maximum Cubemap texture dimensions */
    int    maxTexture1DLayered[2];     /**< Maximum 1D layered texture dimensions */
    int    maxTexture2DLayered[3];     /**< Maximum 2D layered texture dimensions */
    int    maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */
    int    maxSurface1D;               /**< Maximum 1D surface size */
    int    maxSurface2D[2];            /**< Maximum 2D surface dimensions */
    int    maxSurface3D[3];            /**< Maximum 3D surface dimensions */
    int    maxSurface1DLayered[2];     /**< Maximum 1D layered surface dimensions */
    int    maxSurface2DLayered[3];     /**< Maximum 2D layered surface dimensions */
    int    maxSurfaceCubemap;          /**< Maximum Cubemap surface dimensions */
    int    maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */
    size_t surfaceAlignment;           /**< Alignment requirements for surfaces */
    int    concurrentKernels;          /**< Device can possibly execute multiple kernels concurrently */
    int    ECCEnabled;                 /**< Device has ECC support enabled */
    int    pciBusID;                   /**< PCI bus ID of the device */
    int    pciDeviceID;                /**< PCI device ID of the device */
    int    pciDomainID;                /**< PCI domain ID of the device */
    int    tccDriver;                  /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */
    int    asyncEngineCount;           /**< Number of asynchronous engines */
    int    unifiedAddressing;          /**< Device shares a unified address space with the host */
    int    memoryClockRate;            /**< Peak memory clock frequency in kilohertz */
    int    memoryBusWidth;             /**< Global memory bus width in bits */
    int    l2CacheSize;                /**< Size of L2 cache in bytes */
    int    maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */
    int    streamPrioritiesSupported;  /**< Device supports stream priorities */
    int    globalL1CacheSupported;     /**< Device supports caching globals in L1 */
    int    localL1CacheSupported;      /**< Device supports caching locals in L1 */
    size_t sharedMemPerMultiprocessor; /**< Shared memory available per multiprocessor in bytes */
    int    regsPerMultiprocessor;      /**< 32-bit registers available per multiprocessor */
    int    managedMemory;              /**< Device supports allocating managed memory on this system */
    int    isMultiGpuBoard;            /**< Device is on a multi-GPU board */
    int    multiGpuBoardGroupID;       /**< Unique identifier for a group of devices on the same multi-GPU board */
};

常见问题

NVCC没有配置,导致undefined reference HEADER DIR没有配置,导致找不到头文件

相关文章:

  • 3.4背景图片位置
  • Android GNSSGPS 基础知识以及代码流程分析
  • 微控制器通信2 -基础(微逻辑电平与I/O电压)
  • 设计模式-享元模式、享元模式示例
  • 过滤器---普通过滤器使用、串联过滤器、局部与全局过滤器
  • Docker安装Redis 7.x单机模式
  • 2022-09-04 C++并发编程(十八)
  • 插入一百万数据的最优解分析和耗时
  • DockerFile的基本知识及利用DockerFile构建镜像
  • Spring(二)
  • 计算机毕业设计ssm+vue基本微信小程序的执法助手平台
  • Java项目--网页版音乐播放器(JQuery前端逻辑)
  • windows service 服务器安装 MySQL
  • springboot+mybatis+mysql+Quartz实现任务调度(定时任务,实现可配置)
  • python简介常考面试题目:python是什么,有什么好处,python2和python3的主要区别
  • $translatePartialLoader加载失败及解决方式
  • [NodeJS] 关于Buffer
  • 【干货分享】SpringCloud微服务架构分布式组件如何共享session对象
  • Android Studio:GIT提交项目到远程仓库
  • codis proxy处理流程
  • extjs4学习之配置
  • HTML中设置input等文本框为不可操作
  • Java基本数据类型之Number
  • MyEclipse 8.0 GA 搭建 Struts2 + Spring2 + Hibernate3 (测试)
  • orm2 中文文档 3.1 模型属性
  • 和 || 运算
  • 解决iview多表头动态更改列元素发生的错误
  • 前端存储 - localStorage
  • 全栈开发——Linux
  • 我的面试准备过程--容器(更新中)
  • 小程序测试方案初探
  • 7行Python代码的人脸识别
  • 正则表达式-基础知识Review
  • ​​​​​​​​​​​​​​Γ函数
  • # 透过事物看本质的能力怎么培养?
  • $ is not function   和JQUERY 命名 冲突的解说 Jquer问题 (
  • (14)学习笔记:动手深度学习(Pytorch神经网络基础)
  • (C#)Windows Shell 外壳编程系列4 - 上下文菜单(iContextMenu)(二)嵌入菜单和执行命令...
  • (附源码)ssm教师工作量核算统计系统 毕业设计 162307
  • (七)Knockout 创建自定义绑定
  • (算法)前K大的和
  • (算法二)滑动窗口
  • (五)IO流之ByteArrayInput/OutputStream
  • (五)关系数据库标准语言SQL
  • (原創) X61用戶,小心你的上蓋!! (NB) (ThinkPad) (X61)
  • (转)Linux整合apache和tomcat构建Web服务器
  • (轉貼) 2008 Altera 亞洲創新大賽 台灣學生成果傲視全球 [照片花絮] (SOC) (News)
  • ./indexer: error while loading shared libraries: libmysqlclient.so.18: cannot open shared object fil
  • .NET Framework 4.6.2改进了WPF和安全性
  • .NET 使用配置文件
  • .NET 中创建支持集合初始化器的类型
  • .NET是什么
  • .NET正则基础之——正则委托
  • /run/containerd/containerd.sock connect: connection refused
  • /ThinkPHP/Library/Think/Storage/Driver/File.class.php  LINE: 48