cudagpu入门介绍.ppt

An introduction to beginners,GPU & CUDA,内容,1 GPU高性能计算 2 CUDA架构 3 CUDA环境搭建 4 CUDA简单例子,1. GPU高性能计算,GPU: 从图形处理到通用计算 GPU: Graphic Processing Unit GPU为高度并行的实时3D渲染计算而设计, 高GFLOPS , 高带宽 3D渲染技术及3D API的发展, 促进GPU向通用计算处理器发展 NVIDIA GPU为通用计算专门优化设计, 于2007年推出CUDA,1. GPU高性能计算,GPU VS CPU: 计算能力,8x double precision ECC L1, L2 Caches,1 TF Single Precision 4GB Memory,1. GPU高性能计算,GPU VS CPU: 硬件体系,目标: 最小的指令延迟 巨大的缓存 复杂的控制逻辑,目标: 最大吞吐量 更多寄存器、高带宽 通过线程轮换隐藏延迟 多个线程共享控制逻辑,1. GPU高性能计算,GPU应用领域 用同一计算方法对很大量的数据进行并行计算 数据相关度低 计算密度大，计算所花的时间比数据存储的时间大得多,1. GPU高性能计算,GPU关键概念 SM: Stream Multiprocessor SP: Stream Processor 每个SM包含8个SP, 由SM取指, 解码, 发射到各个SP, GPU可看作是一组SM SMEM: shared memory, GPU片内每个SM所占有的高速存储器 Global memory: GPU板载显存,2. CUDA架构,CUDA– Compute Unified Device Architecture 由NVIDIA 2007年推出的通用并行计算架构 第一种不需要借助图形学API就可以使用类C语言极性通用计算的开发环境和软件体系 在采用了统一架构的GPU上运行 NVIDIA提供完全免费的CUDA开发工具包和SDK 对标准C语言的简单扩展, 以最小的扩展集实现CUDA架构,2. CUDA架构,存储器模型 Thread私有register (RW) Thread私有local memory (RW) Block的shared memory (RW) Grid的global memory (RW) Grid的constant memory (RO) Grid的texture memory (RO) Host可以通过API完成显存 与主存交互,2. CUDA架构,线程组织模型 每个内核程序 (Kernel) 包含非常多的线程 所有线程执行相同的顺序程序 层次结构：Grid  Block  Thread 线程被划分成线程块 (Block) 同一线程块内的线程可以通过共享SM资源相互协作 每个线程和线程块拥有唯一ID 通过内建变量读取,2. CUDA架构,线程映射 Thread  SP Block  SM Grid  GPU 把GPU看作一个包含很多个核心的多核处理器 有自己的本地存储器 大量的线程并行执行 每个线程拥有私有变量/存储区域 线程之间拥有共享的存储区域 与CPU的差别: GPU的线程是轻量级的（代码少，硬件轮换，资源负担小）；需要很大量的线程才能使GPU满载,2. CUDA架构,CPU/GPU异构并行模型 由CPU端 (主机端) 传送参数及配置Kernel的执行 (Block, Thread数量, 尺寸等) Kernel在GPU (设备端) 上执行 显存分配, 显存与PC主存间的交互通过API完成,2. CUDA架构,CUDA C扩展: 编译器nvcc CUDA源程序 -- CU文件 编译器驱动 分离主机端源代码和设备端源代码 编译设备端源代码到GPU汇编代码（PTX）或GPU二进制代码（CUBIN） 把最终GPU代码包裹在主机端源代码中 CUDA C扩展: API Driver API 调用Kernel较繁琐 Runtime API 较方便,2. CUDA架构,CUDA C扩展: 函数限定符 __global__ GPU Kernel代码, 由CPU发起, 返回void, 不能由其它Kernel调用 __device__ 由GPU Kernel调用的函数, 不能由CPU发起 __host__ 在CPU上执行的函数 CUDA C扩展: 变量限定符 __constant__ 变量位于常数存储器 __shared__ 变量位于共享存储器,2. CUDA架构,CUDA C扩展: 内建变量 无需定义/声明, 直接在kernel中调用 dim3 gridDim; // Grid尺寸 dim3 blockDim; // Block尺寸 dim3 blockIdx; // Block ID dim3 threadIdx; // Thread 点 CUDA C扩展: 执行配置 调用: my_kernel (.); 实际是为gridDim和blockDim赋值 决定当前线程需要处理的数据和结果存储的位置,2. CUDA架构,CUDA存储器管理 GPU与CPU独立编址 GPU与CPU数据交互通过PCIe总线 在主机端通过API去分配/复制/销毁GPU上的存储器空间 GPU存储器操作API (详见《CUDA参考手册》) 存储空间开辟: cudaMalloc等 存储空间初始化: cudaMemset等 存储空间: cudaFree等 存储空间拷贝cudaMemcpy等,2. CUDA架构,CUDA程序结构,3. CUDA环境搭建,所需组件 CUDA完全免费, 各种操作系统下的CUDA安装包均可以在 CUDA环境搭建,1. Windows 环境要求: Microsoft Visual Stadio 2005 及其以后版本 (1) 依次安装: 显卡驱动ToolkitSDK (2) 语法高亮 到SDK安装目录下, doc  syntax_highlighting  visual_studio_8, 将其中的usertype.dat文件复制到Microsoft Visual Studio 8\Common7\IDE下. 启动VS2005, 工具选项文本编辑器文件扩展名. 在右边的“扩展名”对应的栏里输入“cu”; “编辑器”选择 “Microsoft Visual C++”, “添加”  “确定”  重启VS2005 (3) 设置VS2005环境 打开VS2005，工具选项项目和解决方案VC++ 目录 “显示以下内容的目录” “包含文件”: 分别添加Toolkit和SDK的文件包含目录 “库文件”: 分别添加Toolkit和SDK的库目录 “源文件”: 添加SDK源文件目录,3. CUDA环境搭建,1. Windows (4) 编译定义 工程源文件 属性  常规  工具  自定义生产工具, 在以下项添加相应内容: “命令行”: “$(CUDA_BIN_PATH)\nvcc.exe“ -ccbin “$(VCInstallDir)bin“ -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I“$(CUDA_INC_PATH)“ -o $(ConfigurationName)\$(InputName).obj $(InputFileName) “输出”: $(ConfigurationName)\$(InputName).obj “附加依赖项”: $(InputName)_kernel.cu 右键工程名  属性, 打开链接器, 在以下项做相应修改 “启用增量链接” : 否(/INCREMENTAL:NO) “附加库目录” :改为 “$(CUDA_LIB_PATH)” “链接器” 子项 “输入”: “附加依赖项” 中输入cudart.lib,3. CUDA环境搭建,2. Linux (Fedora, Redhat, Ubuntu…) 环境要求: gcc编译器 (1) 进入Linux文本界面 (2) 安装三个组件(具体版本具体输入) sh NVIDIA-Linux-x86_64-180.22-pkg2.run sh cudatoolkit_2.1_linux64_rhel5.2.run sh cuda-sdk-linux-2.10.1215.2015-3233425.run (3) 在.bash_profile中添加: PATH=$PATH:/bin LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/lib64 export PATH export LD_LIBRARY_PATH 其中: 为实际安装路径,4. 例子—GPU并行矢量求和,算法描述 将两个N维矢量相加, 即对应元素相加: ci=ai+bi i∈[0, N) 并行算法设计 每个GPU线程负责一个维度的相加 GPU线程数目 M ≥ N,a[0],b[0],c[0],Th[0],a[N-1],b[N-1],c[N-1],Th[N-1],a[1],b[1],c[1],Th[1],a[2],b[2],c[2],Th[0],a[…],b[…],c[…],Th[…],Th[…],Th[M],a,b,c,,,,,,,,,,,,,,,4. 例子—GPU并行矢量求和,预处理部分 cuda_runtime.h: CUDA运行时库,4. 例子—GPU并行矢量求和,内核函数 返回值类型必须包括 __global__ void 内核函数定义的是每一个线程的行为, 所有线程并发执行 操作目标通过当前线程程结构中的位置而确定 该问题中所有线程均只执行一次加法, 总计算时间为执行67~70行代码一次所需时间,4. 例子—GPU并行矢量求和,主函数(part-1),使用线程结构的x维度, 8×8一共64个线程,h_前缀表示变量位于Host端存储器 d_前缀表示变量位于Device端存储器,还有多种其他开辟方式,4. 例子—GPU并行矢量求和,主函数(part-2),推荐参考,书籍/资料: NVIDIA CUDA Programming Guide The CUDA Compiler Driver NVCC Get Started NVIDIA CUDA C Installation and Verification on Linux CUDA by example GPU高性能计算之CUDA 深入浅出谈CUDA 网站/论坛 。