网站推广优化方案,html登录页面制作,网站建设新闻稿,网站优化推广服务主要整理了N多年前#xff08;2013年#xff09;学习CUDA的时候开始总结的知识点#xff0c;好长时间不写CUDA代码了#xff0c;现在LLM推理需要重新学习CUDA编程#xff0c;看来出来混迟早要还的。
1.CUDA 解析#xff1a;2007年#xff0c;NVIDIA推出CUDA#xff08… 主要整理了N多年前2013年学习CUDA的时候开始总结的知识点好长时间不写CUDA代码了现在LLM推理需要重新学习CUDA编程看来出来混迟早要还的。
1.CUDA 解析2007年NVIDIA推出CUDACompute Unified Device Architecture统一计算设备架构这个编程模型目的是为了在应用程序中充分利用CPU和GPU各自的优点实现CPU/GPU联合执行。这种联合执行的需要已经在最新的集中编程模型OpenCLOpenACCC AMP中体现出来了。
2.并行编程语言和模型 解析使用比较广泛的是为可扩展的集群计算设计的消息传递接口Message Passing InterfaceMPI和为共享存储器的多处理器系统设计的OpenMP。目前很多HPCHigh-Performance Computing集群采用的都是异构的CPU/GPU节点模型也就是MPI和CUDA的混合编程来实现多机多卡模型。目前支持CUDA的编程语言有CCFortranPythonJava [2]。CUDA采用的是SPMDSingle-Program Multiple-Data单程序多数据的并行编程风格。
3.数据并行性任务并行性 解析任务并行性通常对应用进行任务分解得到。例如对一个需要做向量加法和矩阵-向量乘法的简单应用来说每个操作可以看作一个任务。如果这两个任务可以独立地执行那么就能得到任务并行性。
4.CUDA对C中函数声明的扩展 解析
1__device__ float DeviceFunc()在设备上执行并且只能从设备上调用。
2__global__ float KernelFunc()在设备上执行并且只能从主机上调用。
3__host__ float HostFunc()在主机上执行并且只能从主机上调用。说明如果在函数声明时没有指定CUDA扩展关键字则默认的函数是主机函数。
5.threadblockgridwarpspsm 解析 1grid、block、thread在利用CUDA进行编程时一个grid分为多个block而一个block分为多个thread。 2sp最基本的处理单元最后具体的指令和任务都是在sp上处理的。 3sm多个sp加上其它的一些资源比如存储资源、共享内存、寄储器等组成一个sm。 4warpGPU执行程序时的调度单位。目前CUDA的warp大小32同在一个warp的线程以不同数据资源执行相同指令。
6.CUDA核函数 解析kernel函数完整的执行配置参数形式是Dg, Db, Ns, S如下所示 1参数Dg用于定义整个grid的维度和尺寸即一个grid有多少个block。 2参数Db用于定义一个block的维度和尺寸即一个block有多少个thread。 3参数Ns是一个可选参数用于设置每个block除了静态分配的shared Memory以外最多能动态分配的shared memory大小单位为byte。不需要动态分配时该值为0或省略不写。 4参数S是一个cudaStream_t类型的可选参数初始值为零表示该核函数处在哪个流之中。
7.CUDA存储系统 解析每个线程拥有独立的寄存器(register)和局部存储器(local memory)每个线程块拥有独立的共享存储器(shared memory)所有线程都可以访问全局存储器(global memory)以及只读存储器常量存储器(constant memory)和纹理存储器(texture memory)。如下所示 1寄存器(register) 2局部存储器(local memory) 3共享存储器(shared memory) eg__shared__ 数据类型 变量名; 4全局存储器(global memory) 5常量存储器(constant memory) eg__constant__ 数据类型 变量名; 6纹理存储器(texture memory)
8.CUDA线程组织 解析CUDA通过内置变量threadIdx.x、threadIdx.y、threadIdx.z表示当前线程所处的线程块的位置blockIdx.x、blockIdx.y、blockIdx.z表示当前线程所处的线程块在整个网格中所处的位置blockDim.x、blockDim.y、blockDim.z表示块的维度gridDim.x、gridDim.y、gridDim.z表示网格的维度。对网格中的任意点(i, j)表征CUDA代码中的一个线程该线程对应的网格中的索引如下所示
i threadIdx.xblockIdx.x*blockDim.x
j threadIdx.yblockIdx.y*blockDim.y9.变量类型 解析
1__device__GPU的global memory空间grid中所有线程可访问。
2__constant__GPU的constant memory空间grid中所有线程可访问。
3__shared__GPU上的thread block空间block中所有线程可访问。
4local位于SM内仅本thread可访问。10.CUDA函数库 解析 1Thrust一个C STL实现的函数库。 2NVPPNVIDIA性能原语和Intel的MKK类似。 3CuBLASBLAS基本线性代数函数库的GPU版本。 4cuFFTGPU加速的快速傅里叶变换函数库。 5cuSparse稀疏矩阵数据的线性代数和矩阵操作。 6MagmaLAPACK和BLAS函数库。 7GPU AI基于GPU的路径规划和碰撞避免。 8CUDA Math Lib支持C99标准的数学函数。 9Jacket对.m代码可选的、基于GPU的Matlab引擎。 10Array Fire类似于IPP、MKL和Eigen的矩阵、信号和图像处理库。 11CULA工具线性代数库。 12IMSLFortran IMSL数值函数库的实现。 13NPPNVIDIA Performance Primitives提供了一系列图像和通用信号处理的函数并且支持所有的CUDA平台。
11.Ubuntu 14.04安装CUDA 7.5 解析由于无论是CentOS还是Ubuntu都预装了开源的nouveau显卡驱动SUSE没有这种问题如果不禁用则CUDA驱动可能不能正确安装。处理方法如下所示
sudo vim /etc/modprobe.d/blacklist.conf
blacklist nouveau # 增加一行
sudo apt-get --purge remove xserver-xorg-video-nouveau # 把官方驱动彻底卸载
sudo apt-get --purge remove nvidia-* # 清除之前安装的任何NVIDIA驱动
sudo service lightdm stop # 进命令行关闭Xserver
sudo kill all Xorg # 杀死所有XorgXorg软件包是包括显卡驱动、图形环境库等一系列软件包说明gksudo nvidia-settings可以进行NVIDIA X Server Settings的设置。 1安装依赖类库
sudo apt-get update
sudo apt-get upgrade
sudo apt-get install build-essential cmake g gfortran git pkg-config python-dev software-properties-common wget
sudo apt-get autoremove
sudo rm -rf /var/lib/apt/lists/*2安装Nvidia驱动
lspci | grep -i nvidiasudo add-apt-repository ppa:graphics-drivers/ppa
sudo apt-get update
sudo apt-get install nvidia-352sudo shutdown -r now
cat /proc/driver/nvidia/version3安装CUDA
sudo dpkg -i cuda-repo-ubuntu1404*amd64.deb
# sudo apt-get update
sudo apt-get install cudaecho export PATH/usr/local/cuda/bin:$PATH ~/.bashrc
echo export LD_LIBRARY_PATH/usr/local/cuda/lib64:$LD_LIBRARY_PATH ~/.bashrc
source ~/.bashrcnvcc -V
sudo shutdown -r now4安装CUDA例子和测试
/usr/local/cuda/bin/cuda-install-samples-7.5.sh ~/cuda-samples
cd ~/cuda-samples/NVIDIA*Samples
make -j $(($(nproc) 1))bin/x86_64/linux/release/deviceQuery说明至此安装完毕就可以使用Nsight Eclipse Edition愉快地进行CUDA异构计算开发了。
12.二维数组使用
#include iostream
#includecstdlib
using namespace std;static const int ROW 10;
static const int COL 5;int main() {int** array (int**)malloc(ROW*sizeof(int*));int* data (int*)malloc(ROW*COL*sizeof(int));// initialize the datafor (int i0; iROW*COL; i) {data[i] i;}// initialize the arrayfor (int i0; iROW; i) {array[i] data i*COL;}// output the arrayfor (int i0; iROW; i) for (int j0; jCOL; j) {cout array[i][j] endl;}free(array);free(data);return 0;
}13.查看GPU信息的代码 如果找不到#include helper_cuda.h那么将$cuda-samples/NVIDIA_CUDA-7.5_Samples/common/inc添加到NVCC Compiler中的Includes中即可。 解析 1#include #includecstdlib提供数据类型size_twchar_tdiv_tldiv_tlldiv_t提供常量NULLEXIT_FAILUREEXIT_SUCCESSRAND_MAXMB_CUR_MAX提供函数atofatoiatol strtod strtof strtols strtol strtoll strtoul strtoull rand srand calloc free malloc realloc abort atexit exit getenv system bsearch qsort abs div labs ldiv llabs tlldiv mblen mbtowc wctomb mbstowcs wcstombs。 2CUDART_VERSION CUDA Runtime API Version在#include cuda_runtime.h中#define CUDART_VERSION 7050。 3stdinstdoutstderr 进程将从标准输入文件中得到输入数据将正常输出数据输出到标准输出文件而将错误信息送到标准错误文件中。stdin的文件描述符为0stdout的文件描述符为1stderr的文件描述符为2。 4host cudaError_t cudaSetDevice ( int device ) 设置device用于GPU扩展。 5cudaDeviceProp数据结构 cudaDeviceProp数据类型针对函式cudaGetDeviceProperties定义的cudaGetDeviceProperties函数的功能是取得支持GPU计算装置的相关属性比如支持CUDA版本号装置的名称、内存的大小、最大的thread数目、执行单元的频率等。如下所示
struct cudaDeviceProp {char name[256]; // 识别设备的ASCII字符串比如GeForce GTX 940Msize_t totalGlobalMem; // 全局内存大小size_t sharedMemPerBlock; // 每个block内共享内存的大小int regsPerBlock; // 每个block 32位寄存器的个数int warpSize; // warp大小size_t memPitch; // 内存中允许的最大间距字节数int maxThreadsPerBlock; // 每个Block中最大的线程数是多少int maxThreadsDim[3]; // 一个块中每个维度的最大线程数int maxGridSize[3]; // 一个网格的每个维度的块数量size_t totalConstMem; // 可用恒定内存量int major; // 该设备计算能力的主要修订版号int minor; // 设备计算能力的小修订版本号int clockRate; // 时钟速率size_t textureAlignment; // 该设备对纹理对齐的要求int deviceOverlap; // 一个布尔值表示该装置是否能够同时进行cudamemcpy()和内核执行int multiProcessorCount; // 设备上的处理器的数量int kernelExecTimeoutEnabled; // 一个布尔值该值表示在该设备上执行的内核是否有运行时的限制int integrated; // 返回一个布尔值表示设备是否是一个集成的GPU即部分的芯片组、没有独立显卡等int canMapHostMemory; // 表示设备是否可以映射到CUDA设备主机内存地址空间的布尔值int computeMode; // 一个值该值表示该设备的计算模式默认值专有的或禁止的int maxTexture1D; // 一维纹理内存最大值int maxTexture2D[2]; // 二维纹理内存最大值int maxTexture3D[3]; // 三维纹理内存最大值int maxTexture2DArray[3]; // 二维纹理阵列支持的最大尺寸int concurrentKernels; // 一个布尔值该值表示该设备是否支持在同一上下文中同时执行多个内核6host cudaError_t cudaDriverGetVersion ( int* driverVersion ) 返回CUDA驱动版本。 7host __device__ cudaError_t cudaRuntimeGetVersion ( int* runtimeVersion ) 返回CUDA运行时版本。 说明__host__和__device__同时使用时触发编译系统生成同一函数的两个不同的版本。它支持一种常见的应用即只需要重编译同一函数的源代码就可以生成一个在设备上运行的版本。 8#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) WIN32_WIN32WIN64_WIN64是Windows操作系统预定义的宏。这句话的目的是C/C编程通过宏定义来判断操作系统的类型。
14.#includedevice_launch_parameters.h 解析#includedevice_launch_parameters.h头文件包含了内核函数的5个变量threadIdx、blockDim、blockIdx、gridDim和wrapSize。
15.事件管理 解析常用函数如下所示 1cudaEventCreate()事件的创建。 2cudaEventDestroy()事件的销毁。 3cudaEventRecord()记录事件。 4cudaEventSynchronize()事件同步。 5cudaEventElapsedTime()计算两事件的时间差。 利用CUDA提供的事件管理API实现计时功能如下所示
float time;
cudaEventCreate(start);
cudaEventCreate(stop);
cudaEventRecord(start, 0);
{
// 需要计时的代码
}
cudaEventRecord(stop,0);
cudaEventElapsedTime(time, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);16.deviceQuery 解析输出结果如下所示
rootubuntu:~/cuda-samples/NVIDIA_CUDA-7.5_Samples/1_Utilities/deviceQuery# ./deviceQuery
./deviceQuery Starting...CUDA Device Query (Runtime API) version (CUDART static linking)Detected 1 CUDA Capable device(s)Device 0: GeForce 940MCUDA Driver Version / Runtime Version 7.5 / 7.5CUDA Capability Major/Minor version number: 5.0Total amount of global memory: 1024 MBytes (1073610752 bytes)( 3) Multiprocessors, (128) CUDA Cores/MP: 384 CUDA CoresGPU Max Clock rate: 980 MHz (0.98 GHz)Memory Clock rate: 1001 MhzMemory Bus Width: 64-bitL2 Cache Size: 1048576 bytesMaximum Texture Dimension Size (x,y,z) 1D(65536), 2D(65536, 65536), 3D(4096, 4096, 4096)Maximum Layered 1D Texture Size, (num) layers 1D(16384), 2048 layersMaximum Layered 2D Texture Size, (num) layers 2D(16384, 16384), 2048 layersTotal amount of constant memory: 65536 bytesTotal amount of shared memory per block: 49152 bytesTotal number of registers available per block: 65536Warp size: 32Maximum number of threads per multiprocessor: 2048Maximum number of threads per block: 1024Max dimension size of a thread block (x,y,z): (1024, 1024, 64)Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)Maximum memory pitch: 2147483647 bytesTexture alignment: 512 bytesConcurrent copy and kernel execution: Yes with 1 copy engine(s)Run time limit on kernels: YesIntegrated GPU sharing Host Memory: NoSupport host page-locked memory mapping: YesAlignment requirement for Surfaces: YesDevice has ECC support: DisabledDevice supports Unified Addressing (UVA): YesDevice PCI Domain ID / Bus ID / location ID: 0 / 4 / 0Compute Mode: Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) deviceQuery, CUDA Driver CUDART, CUDA Driver Version 7.5, CUDA Runtime Version 7.5, NumDevs 1, Device0 GeForce 940M
Result PASS说明3个SM每个SM包含128个SP总共384个SP。
17.矩阵分块乘法[15] 解析
18.一维卷积 解析卷积的物理意义就是加权叠加常见的操作有一维卷积和二维卷积。需要说明的是“幽灵元素通常补0。因为全局存储器和常数存储器变量都是存储在DRAM中的从DRAM中访问一个变量需要数百甚至数千个时钟周期从DRAM访问变量的速度通常要比处理器完成算术操作的 速度慢很多。因此DRAM的长延迟和有限带宽已经成为几乎所有现代处理器的性能瓶颈通常被称为存储墙问题。为了缓解这个问题现代处理器通常引进高速 缓存存储器或多级高速缓存以减少访问DRAM的次数。 1一维卷积 该算法的缺点是受限于访问输入数组d_N的DRAM带宽掩码数组d_M已经放在L1高速缓存。 2分块一维卷积1 规定一个线程块处理的输出元素的集合成为输出块”计算不涉及幽灵元素的分块称为中间块输出数据最左边的分块称为左边界分块输出数据最右边的分块称为右边界分块涉及多个分块、被多个线程块加载的数据称为光环元素或者边缘元素。 该算法的优点是将输入数组d_N分块后放入共享存储器减少了输入数组d_N的DRAM带宽掩码数组d_M已经放在L1高速缓存。分块一维卷积1对一维卷积的改进 3分块一维卷积2 该算法的优点是充分利用了L2高速缓存掩码数组d_M已经放在L1高速缓存。需要说明的是最近的GPU比如Fermi提供了通用L1和L2高速缓存L1缓存是每个SM私有的而L2缓存是所有SM共享的。这样的话线程块中的光环元素可能存放在L2高速缓存中我们直接利用光环元素即可而不再需要将光环元素加载到线程块的共享存储器中。分块一维卷积2对分块一维卷积1的改进
19.cudaMemcpyToSymbol
解析在device中cudaMemcpyToSymbol进行赋值和读取devicesharedconstant而在host中cudaMemcpy进行赋值和读取。需要说明的是在函数体外声明的变量默认为__device__类型即全局变量类型。cudaMemcpyToSymbol(d_M, h_M, MaskLen*sizeof(float));和cudaMemcpyFromSymbol(h_M, d_M,MaskLen*sizeof(float));功能相反。常量存储器中的内容通常放在L1高速缓存中因为常量存储器变量在kernel函数执行期间不会改变刚好与GPU不提供缓存一致性相符合主要是为了最大化利用硬件资源提高处理器的算术运算吞吐率等。需要说明的是现代CPU通常都支持处理器核心之间的缓存一致性。
20.shared memory与bank conflict 解析 1什么是bank conflict 在实际中shared memory被分割成32个等大小的存储体比如Maxwell架构即bank每个bank拥有每周期32bit4byte4char1int1float的宽度。因为一个warp中有32个线程相当于一个线程对应一个bank。不同的设备存储体的数目也不相同比如设备Tesla架构为16个存储体后面我们以32个存储体进行讲解。 对于计算能力1.0设备前个half-warp和后个half-warp不存在bank conflict对于计算能力2.0设备前个half-warp和后个half-warp可能存在bank conflict因为shared memory可以同时让32个bank响应请求。 2bank conflict发生的原因 理想情况下就是不同的线程访问不同的bank可能是规则的访问比如线程0读写bank0线程1读写bank1也可能是不规则的比如线程0读写bank1线程1读写bank0。这种同一时刻每个bank只被最多1个线程访问的情况下不会出现bank conflict。特殊情况如果有多个线程同时访问同一个bank的同一个地址的时候也不会产生bank conflict即broadcast。但当多个线程同时访问同一个bank不同的地址时那么bank conflict就发生了。比如线程0访问地址0而线程1访问地址32归一化后为0由于它们在同一个bank中所以就导致了冲突。因为bank conflict发生后同一个bank的内存读写将被串行化会导致程序性能大大降低。 3什么时候会发生bank conflict呢 bank conflict主要出现在global memory与shared memory数据交换及设备函数对shared memory操作中。 4如何避免bank conflict呢 很多时候shared memory的bank conflict可以通过修改数据存储的方式来解决。
21.Thrust库 解析Thrust是一款基于GPU CUDA的C库其中包含诸多并行算法和数据结构。Thrust主要通过管理系统底层的功能比如memory access内存获取和memory allocation内存分配来实现加速使得工程师们在GPU编程的环境下能更focus在算法的设计上。 1容器 host_vector为主机提供的向量类并且驻留在主机内存中device_vector为设备提供的向量类并且驻留在设备全局内存中。 说明fill()copy()sequence()。 2算法 转换(transformation)规约(reduction)前缀求和(prefix sum)再排序(reordering)排序(sorting)。 3迭代器 constant_iteratorcounting_iteratortransform_iteratorpermutation_iteratorzip_iterator。
22.事件流纹理内存 解析 1事件CUDA中的事件本质上是一个GPU时间戳两个事件的时间差就是算法执行时间。 2流CUDA流表示一个GPU操作队列可以将一个流看做是GPU上的一个任务不同任务可以并行执行。 3纹理内存专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的只读内存。
参考文献 [1] Java bindings for CUDAhttp://jcuda.org/ [2] CNN之Caffe配置http://www.cnblogs.com/alfredtofu/p/3577241.html [3] Setting up a Deep Learning Machine from Scratchhttps://github.com/saiprashanths/dl-setup [4] CUDA初学大全http://www.cnblogs.com/yangs/archive/2012/07/28/2613269.html [5] 有哪些优秀的CUDA开源代码https://www.zhihu.com/question/29036289/answer/42971562 [6] CUDA一维矩阵的加法http://tech.it168.com/a2009/1112/807/000000807771.shtml [7] CUDA二维矩阵加法http://www.cnblogs.com/jugg1024/p/4349243.html [8] NVIDIA CUDA Runtime APIhttp://docs.nvidia.com/cuda/cuda-runtime-api/index.html#axzz4G8M3LWlq [9] C/C是如何通过宏定义来判断操作系统的http://www.myexception.cn/operating-system/1981774.html [10] CUDA编程其实写个矩阵相乘并不是那么难http://www.cnblogs.com/yusenwu/p/5300956.html [11] CUDA实例矩阵乘法http://wenku.baidu.com/link?urlXCOgGQqpPUns-cifgGm1tbfqmY-5wWTwkXHh1_i_5ZZX6vFmbFu22r67fWMpcs-GxsH9thzOjVeNCpKIjGjdx2SYhq7bW4qfIquRTM0AAW_ [12] 华科并行计算上机作业http://wenku.baidu.com/link?url1tWvUvW0t7BnFChxetS_Mr5_pCF_LZHQGLWxN-ArVVPccOM_VmoTx9IUD76l_rVMP-iPKWI97vn7wa5ZChz59rr4wlur3rL6k3MGB15qF4W [13] CUDA编程http://www.cnblogs.com/stewart/archive/2013/01/05/2846860.html [14] NVIDIA DockerGPU Server Application Deployment Made Easyhttps://devblogs.nvidia.com/parallelforall/nvidia-docker-gpu-server-application-deployment-made-easy/ [15] CUDA矩阵乘法——利用共享存储器http://blog.csdn.net/augusdi/article/details/12614247 [16] 华科并行计算上机作业http://wenku.baidu.com/link?url1tWvUvW0t7BnFChxetS_Mr5_pCF_LZHQGLWxN-ArVVPccOM_VmoTx9IUD76l_rVMP-iPKWI97vn7wa5ZChz59rr4wlur3rL6k3MGB15qF4W [17] GPUWattch Energy Model Manualhttp://www.gpgpu-sim.org/gpuwattch/ [18] 图形学领域的关键算法及源码链接http://blog.csdn.net/u013476464/article/details/40857873 [19] 拷贝global memorycudaMemcpyToSymbol和cudaMemcpy函数是否有区别http://blog.csdn.net/litdaguang/article/details/45047015 [20] CUDA GPU编程如何避免bank conflicthttp://www.th7.cn/Program/c/201512/719448.shtml [21] CUDA共享内存bank conflicthttp://blog.csdn.net/endlch/article/details/47043069 [22] CUDA bank conflict in shared memoryhttp://bbs.csdn.net/topics/390836540 [23] CUDA Programming Guide之shared memory的Bank Conficthttp://blog.csdn.net/o_oxo_o/article/details/4296281 [24]Parallel_programming_week3.mdhttps://github.com/mebusy/notes/blob/c278e037aa8a59aa139fc722d01ed41cf978921d/dev_notes/Parallel_programming_week3.md [25] Thrusthttp://docs.nvidia.com/cuda/thrust/index.html#axzz4H6gsFZs3 [26] Thrust File Listhttp://thrust.github.io/doc/files.html
