CUDA 并行编程学习笔记

发表于 2022-08-31 更新于 2022-11-17 分类于 blog 阅读数： 88 评论数：本文字数： 7.3k 阅读时长 ≈ 12 分钟

笔者关于 CUDA 并行编程学习中的记录

技术名词

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

计算机架构

冯诺依曼计算机架构

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

连接机

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

Cell 处理器 (众核)

用一个常规处理器作为监管处理器 (PowerPC)，该处理器与大量高速流处理 (SPE) 相连。 * 每个流处理单元 SPE 调用执行一个程序

通过共享的网络，SPE 之间和 SPE 与 PowerPC 之间进行相互通讯

多点计算

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

GPU 架构

nvidia device arch nvidia-gpu-sm-arch

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网络拓扑结构

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 几个头文件

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

CUDA routine

cudaError_t err = cudaSuccess; cudaError_t 类型，表示错误类型。cudaSuccess 表示成功。一般 cuda routine 的返回值都是 cudaError_t 类型，表示函数是否执行成功。
printf("%s\n", cudaGetErrorString(cudaGetLastError())); 输出错误时，使用以上函数转化为 string。
err = cudaMalloc((void **)&d_A, size); 动态内存申请函数，在设备的 global memory 上申请 size 个字节空间。
err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);or err = cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost); // 内存拷贝函数：从 cpu 上的内存 h_A 上拷贝 size 个字节数据到 gpu 上的内存 d_A。反之，一样。
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); 这种形式。
err = cudaGetLastError(); // 启动 kernel 函数时，并没有返回值，通过这个调用这个函数，查看 kernel 函数是否启动成功。
err = cudaFree(d_A); // 释放使用 cudaMalloc 申请的空间。
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 没有配置，导致找不到头文件