13-孔荔.pdf

1、乘影开源GPGPU软件工具链介绍清华大学-软件工程师孔荔1内容2 开源GPGPU 软件工具链设计整体架构设计运行时环境设计驱动程序设计测试套搭建与使用 开源GPGPU 开源社区软件工具链 整体架构设计3开源GPGPU的软件工具链整体架构 实现支持硬件验证和仿真的完整工具链 运行时环境（Runtime）实现OpenCL API接口编译、运行内核程序生成metadta.设备驱动程序（Driver）定义通用接口，兼容不同类型的设备实现对GPGPU硬件资源的控制*OpenCL Overview-The Khronos Group IncRuntime:GPGPU编程模型映射4 GPGPU硬件需要映射到

2、OpenCL编程模型NDRangeKernel Workgroup CTA/BlockWorkitemThread WorkgroupWorkitemOpenCL中一个矩阵乘法的例子计算机体系结构 量化研究方法Runtime:OpenCL API实现5运行时环境需要实现的关键回调函数OpenCL关键API操作API 获取平台中的计算设备信息clGetDeviceIDs创建OpenCL上下文clCreateContext.*创建设备和命令队列clCreateCommandQueue创建和构建程序对象clCreateProgreamWith.创建OpenCL内核clCreateKernel创建Op

3、enCL缓冲区clCreateBuffer设置内核函数参数clSetKernelArg执行内核clEnqueueNDRangeKernel 程序通过ocl-icd调用API实现 生成动态库libpocl.soRuntime:metadata实现6OpenCL程序处理kernel参数提取编译信息.meta.dataCTA硬件开始执行读取runtime 程序执行时，硬件需要一些动态的执行信息，由Runtime生成并传递。指针数组的形式索引kernel参数metadata在内存中的布局Driver：统一设备接口7为上层软件提供统一接口，屏蔽下层差异生成动态库libspike_driver.so驱动程

4、序整体架构统一设备接口描述vt_dev_open打开并连接一个 GPGPU 设备，调用该方法会初始化一个设备类，完成设备预留内存空间的初始化。vt_dev_close关闭一个 GPGPU 设备，调用该方法会等待设备执行完成目前所有任务，释放内存空间。vt_root_mem_alloc 对于一个新的内核函数，创建一个根页表。vt_buf_alloc分配一块内存空间，返回一个设备端内存的虚拟地址vt_buf_free释放传入指针所指的内存空间，需要指定所属的内核函数vt_copy_to_dev将数据从主机端搬移到设备端vt_copy_from_dev将数据从设备端搬移到主机端vt_start启动一

5、个内核函数，需要传入内核函数的 id 和 metadata，该方法会将内核函数拆分为多个线程块，然后调度到 GPGPU 设备上执行，任务调度模块需要在此实现8Driver：内存分配机制 OpenCL内存模型实现了将设备端内存映射到OpenCL内存模型的定义中 global memory local memory private memory支持基于虚拟地址的内存空间分配驱动程序的内存分配算法分配一个新的内存块内存空间分配和映射spike指令级仿真器9添加“乘影”的自定义指令“乘影”指令级仿真器架构“乘影”自定义指令 实现GPGPU组件 simt-stack用于处理分支 GPGPU顶层模块 线程

6、束调度器 实现与driver的接口支持内存系统及管理Driver Interfacemem_tsim_tspike_deviceSimt stackwarp_scheprocprocprocprocTask allocatorMemory manager测试套 OpenCL-CTS 用于 OpenCL 兼容性测试 总体通过率约80%10测试项目用例数量通过数通过率basic11510793%compiler662639%api1059691%computeinfo55100%commonfns181478%rela