从零开始学习OpenCL开发(三)深入API

简介:

 

欢迎关注,转载引用请注明 http://blog.csdn.net/leonwei/article/details/8909897

 

这里将更深入的说明一些OpenCL API的功能

1. 创建buffer

涉及到内存与显存的操作总是复杂麻烦的,这个函数也一样。。。

cl_memclCreateBuffer ( cl_context context,
  cl_mem_flags flags,
  size_t size,
  void *host_ptr,
  cl_int *errcode_ret)

 

函数将创建(或分配)一片buffer,并返回。这里创建的mem可以是globla也可以是local或private,具体要看kernal中怎样声明限定符。cl会根据执行情况自动管理global到更进一层如private的copy。这里的buffer概念是用于kernal函数计算的(或者说是用于device访问的,什么是device?host是C++写的那段控制程序,一定运行在CPU,device就是执行kernal计算的,运行在所有有计算能力的处理器上,有时你的CPU同时扮演host与device,有时用GPU做device),这里模糊了host与device的内存,也就是说根据flag的不同,可以是在host上的,也可以是在device上的,反正只有这里分配的内存可以用于kernal函数的执行。

 

主要的参数在 flags,这些参数可以|

1  CL_MEM_READ_WRITE:在device上开辟一段kernal可读可写的内存,这是默认

2  CL_MEM_WRITE_ONLY:在device上开辟一段kernal只可以写的内存

3  CL_MEM_READ_ONLY:在device上开辟一段kernal只可以读的内存

 

4  CL_MEM_USE_HOST_PTR:直接使用host上一段已经分配的mem供device使用,注意:这里虽然是用了host上已经存在的内存,但是这个内存的值不一定会和经过kernal函数计算后的实际的值,即使用clEnqueueReadBuffer函数拷贝回的内存和原本的内存是不一样的,或者可以认为opencl虽然借用了这块内存作为cl_mem,但是并不保证同步的,不过初始的值是一样的,(可以使用mapmem等方式来同步)

5  CL_MEM_ALLOC_HOST_PTR:在host上新开辟一段内存供device使用

6  CL_MEM_COPY_HOST_PTR:在device上开辟一段内存供device使用,并赋值为host上一段已经存在的mem

 

7  CL_MEM_HOST_WRITE_ONLY:这块内存是host只可写的

8  CL_MEM_HOST_READ_ONLY:这块内存是host只可读的

9  CL_MEM_HOST_NO_ACCESS:这块内存是host可读可写的

 

谈谈这些flag,这些flag看起来行为比较复杂和乱,因为Opencl是一个跨硬件平台的框架,所以要照顾到方方面面,更统一就要更抽象。

首先456的区别,他们都是跟host上内存有关,区别是,4是直接使用已有的,5是新开辟,6是在device上开内存,但是初值与host相同(45都是在host上开内存)

然后看看123 和789,123是针对kernal函数的访问说的,而789是针对host的访问说的,kernal函数是device的访问,而除了kernal函数的访问基本都是host的访问(如enqueueRead/write这些操作)

通常使用host上的内存计算的效率是没有使用device上的效率高的,而创建只读内存比创建可写内存又更加高效(我们都知道GPU上分很多种内存区块,最快的是constant区域,那里通常用于创建只读device内存)

通常用各种方式开内存你的程序都work,但这里就要考验不同情况下优化的功力了

size参数:要开的内存的大小

host_ptr参数:只有在4.6两种情况用到,其他都为NULL

 当然这些内存都要使用clReleaseMemObject释放

 

内存的call_back:

有些方式 ,如CL_MEM_USE_HOST_PTR,cl_mem使用的存储空间实际就在host mem上,所以我们要小心处理这块主存,比如你删了它,但是cl_mem还在用呢,就会出现问题,而clReleaseMemObject并不一定会马上删除这个Cl_mem,它只是一个引用计数的消减,这里需要一个回调,告诉我们什么时候这块主存可以被放心的清理掉,就是clSetMemObjectDestructorCallback

CL的规范中特别说明最好不要在这个callback里面加入耗时的系统和cl API。

 

2.内存操作

1 从Cl_mem读回host mem(就算Cl_mem是直接使用host mem实现的,想读它的内容,还是要这样读回来,可以看做cl_mem是更高一层封装)

  clEnqueueReadBuffer

2 使用host_mem的值写cl_mem

  clEnqueueWriteBuffer

3 在Cl_mem和host mem之间做映射

 clEnqueueMapBuffer

这个函数比较特殊,回顾上面一节在创建buf时有一种方法CL_MEM_USE_HOST_PTR,是直接让device使用host上已有的一块的mem(p1)做buf,但是这个产生的CL_mem(p2)经过计算后值会改变,p2改变后通常p1不会被改变,因为虽然用的一块物理空间,但是cl_mem是高层封装,和host上的mem还是不一样的,要想使p1同步到p2的最新值,就要调用这句map

MAP与CopyBack的性能对比

后来我想了想,这和使用clEnqueueReadBuffer从p2read到p1有什么区别呢?map的方法按道理更快,因为p1p2毕竟一块物理地址吗,map是不是就做个转换,而read则多一遍copy的操作。而且应该在CPU做device时map速度更快,但是事实是这样的吗?本着刨根问题的精神,我真的做了一下实验,

我的实验结果是这样的,如果使用CPU做host,GPU做device,那么CopyBack反而更快,但是如果使用CPU做host,CPU也做device,那么MAP更快(不跨越硬件),而且总体上CPU+GPU的方式更快。

这个实验结果彻底颠覆了我最初的一些想法,实验数据说明1.不考虑硬件差异,MAP确实比CopyBack更快,跟我理解一样,从CPU做device的两组数据就可看出。2.至少在我的这个实验中,主存与显存间的数据copy比主存到主存自己的数据copy更快,所以在CPU+GPU的架构中,由于CopyBack方式采用的是主存显存拷贝,而MAP值涉及主存上的操作,所以CopyBack更快。不过这里我仍存在疑虑,我的分析很可能不对,或存在其他因素没考虑,关于这点,要再继续查查关于pinned memory和内存显存传递数据的一些知识。

所以在这种异构计算领域,性能和你的硬件架构、性能、组合有着非常重要的关联,所以最好的方法就是实际做实验对比。

4  在Cl_mem直接做copy

 clEnqueueCopyBuffer

 

这些函数都跟执行kernal一样是投入到device的command queue里的,但是他们又都带有一个参数blocking_read,可以指定函数是否在执行完毕后返回。

 

3.Program

3.1.compile build link

有两种从文本创建program的方式

  • 直接build:clBuildProgram
  • 先complie好,根据情况动态的link,即把上面的过程拆分为两个步骤

   clCompileProgram   clLinkProgram

但是1.2的方式不保险,这是CL1.2中加入的,而目前不是所有的platform都支持到1.2,NVIDIA好像就才1.1

opencl实际上会根据不同的硬件把通样一份代码编译成不同的机器语言,如CPu汇编或GPU汇编

 

4.Kernal的执行

这里是精华

1.设置kernal的参数

  clSetKernelArg

2.执行kernal

  clEnqueueNDRangeKernel

 

先给一段kernal代码,方便下边参数的解释,另外这里需要一些空间想象能力~

    kernal代码

 __kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
 int idx = get_global_id(0);//得到当前单元格的0维度上的序号
 result[idx] = a[idx] + b[idx];
}

参数说明:

command_queue :执行那个device的命令序列

kernel:待执行的kernal obj

work_dim:我们知道CL的执行是放在一个个独立的compute unit中进行的,你可以想像这些unit是排成一条线的,或是一个二维方阵,甚至是一个立体魔方,或着更高维,这里参数就描述这个执行的维度,从1到CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS之间

global_work_size :每个维度的unit的数量,这样总共拥有的计算单元的数量将是global_work_size[0]*global_work_size[1]...

global_work_offset :这里就是规定上面代码里每个维度上第一个get_global_id()0得到的id,默认为0,例如计算一个一维的长度为255work_size的工作,cl会自动虚拟出255个计算单元,每个单元各自计算0-254位置的数相加,而如果你把他设为3,那么cl会从3开始算,也就是说3-254位置的unit会计算出结果,而0 -2这些unit根本不去参与计算。

local_work_size :前面介绍过CL的的unit是可以组合成组的(同组内可以互相通信)这个参数就决定了Cl的每个组的各维度的大小,NULL时CL会自动给你找个合适的,这里贴下我试着用不同大小的group做数组相加的效率,

这里其实看不太出什么,直觉对这个应用实例是组越少越快,但是其中也不是严格的线性关系,无论在CPU还是GPU上这个关系都是近似的,所以在实际开发中,我们选择什么维度?选择什么样的组大小?我的答案是:多做实验吧,或者要偷懒的话就置0吧,交给CL为你做(实时上Cl中很多函数都有这个NULL的自适应选项。。)

关于维度、偏移、worksize这里有个原版的图,说明的更加形象

 

后面几个参数就跟同步有关了

event_wait_list和num_events_in_wait_list:说明这个command的执行要等这些event执行了之后

event:将返回这个command相关联的event

很明显,通过这几个参数的event可以控制command之间的执行顺序。

 

5.指令执行顺序和同步

command的执行默认都是异步的,这才有利于并行度提高效率,在并行的问题中我们有时经常要做些同步的事情,或者等待某个异步的操作完成,这里有两种方法:

  • 使用EnqueueRead/write这些操作可以指定他们为同步的(即执行完毕才在host上返回)
  • 使用event来跟踪,像clEnqueueNDRangeKernel这样的操作都会关联一个event

event:

  • clEnqueue这样的操作都会关联返回一个event
  • 用户可以自己创建一个自定义的event clCreateUserEvent,要使用clReleaseEvent释放

关于event的操作:

        正是通过event同步不同的command:

  •  设置event状态:     

       设置用户自定义event的状态,clSetUserEventStatus 状态只可以被设定一次,只可以为CL_COMPLETE或者一个负值,CL_COMPLETE代表这个event完成了,等待它的那些command得以执行,而负值表示引起错误,所有等待他的那些command都被取消执行。其实event的状态还有CL_RUNNING  CL_SUBMITTED   CL_QUEUED,只是不能在这里设置。

  • 等待event

        clWaitForEvents;可以在host中等待某些event的结束,如clEnqueueNDRangeKernel这样的异步操作,你可以等待他的event结束,就标志着它执行完了

  • 查询event信息:clGetEventInfo clGetEventProfilingInfo
  • 设置回调:clSetEventCallback

不同device上的event:

      clEnqueueNDRangeKernel这样的操作等待的只能是处于相同queue里面的event(也就是同一个device上的),而同步不同queue上的event则只能用显示的方法,如clWaitForEvents等。

marker:

      marker是这样一个object,它可以看做是一个投入queue的空指令,专门用于同步,它可以向其他comman一样设定需要等待的event,操作有clEnqueueMarkerWithWaitList

barrier:

   barrier和marker十分类似,但是从名字上就可以看出最大的不同点是:marker在等待到它的依赖event之后会自动执行完毕,让后续指令执行,而barrier会阻塞在这里,直到他关联的event被显示的设置成完成状态

marker和barrier的实现在1.1和1.2版本上存在着较大的不同

同步是CL的大问题,关于同步,原版overview上也有一个非常生动的图,贴在这里吧:

在同一个device上同步

在多个device间同步

 

  

 

相关实践学习
在云上部署ChatGLM2-6B大模型(GPU版)
ChatGLM2-6B是由智谱AI及清华KEG实验室于2023年6月发布的中英双语对话开源大模型。通过本实验,可以学习如何配置AIGC开发环境,如何部署ChatGLM2-6B大模型。
目录
相关文章
|
8天前
|
数据可视化 测试技术 API
从接口性能到稳定性:这些API调试工具,让你的开发过程事半功倍
在软件开发中,接口调试与测试对接口性能、稳定性、准确性及团队协作至关重要。随着开发节奏加快,传统方式已难满足需求,专业API工具成为首选。本文介绍了Apifox、Postman、YApi、SoapUI、JMeter、Swagger等主流工具,对比其功能与适用场景,并推荐Apifox作为集成度高、支持中文、可视化强的一体化解决方案,助力提升API开发与测试效率。
|
23天前
|
人工智能 自然语言处理 机器人
使用 API 编程开发扣子应用
扣子(Coze)应用支持通过 API 编程,将 AI 聊天、内容生成、工作流自动化等功能集成至自有系统。主要 API 包括 Bot API(用于消息交互与会话管理)及插件与知识库 API(扩展功能与数据管理)。开发流程包括创建应用、获取密钥、调用 API 并处理响应,支持 Python 等语言。建议加强错误处理、密钥安全与会话管理,提升集成灵活性与应用扩展性。
281 0
|
1月前
|
前端开发 Java API
利用 Spring WebFlux 技术打造高效非阻塞 API 的完整开发方案与实践技巧
本文介绍了如何使用Spring WebFlux构建高效、可扩展的非阻塞API,涵盖响应式编程核心概念、技术方案设计及具体实现示例,适用于高并发场景下的API开发。
206 0
|
1月前
|
存储 监控 算法
淘宝买家秀 API开发实录Python(2025)
本文讲述了作者在电商开发领域,尤其是对接淘宝买家秀 API 接口过程中所经历的挑战与收获。从申请接入、签名验证、频率限制到数据处理和实时监控,作者分享了多个实战经验与代码示例,帮助开发者更高效地获取和处理买家秀数据,提升开发效率。
|
1月前
|
测试技术 API 开发工具
API文档该怎么写,开发效率能翻几倍?
API文档是提升开发效率与协作的关键因素,本文探讨了API文档的核心要素、常见类型及编写规范,并介绍了如何借助现代化工具如Apifox实现高效管理与维护,助力团队打造高质量的API文档体系。
|
1月前
|
算法 前端开发 API
京东比价项目开发实录:京东API接口(2025)
本文分享了作者在电商开发中对接京东商品详情API的实战经验,涵盖了申请权限、签名算法、限流控制、数据解析等常见问题,并提供了亲测有效的Python代码示例,帮助开发者避坑。
|
1月前
|
人工智能 API 定位技术
MCP 开发实战:手把手教你封装高德地图与 arXiv API
本教程为 MCP(Model Context Protocol)开发实战第二阶段,带你从零封装第三方 API 为 AI 模型可用工具。通过高德地图地理编码与 arXiv 论文检索两个实例,涵盖项目搭建、工具声明、资源定义、错误处理等核心内容,助你快速上手 MCP 开发并集成至 Claude 使用。
|
1月前
|
缓存 Java API
Spring WebFlux 2025 实操指南详解高性能非阻塞 API 开发全流程核心技巧
本指南基于Spring WebFlux 2025最新技术栈,详解如何构建高性能非阻塞API。涵盖环境搭建、响应式数据访问、注解与函数式两种API开发模式、响应式客户端使用、测试方法及性能优化技巧,助你掌握Spring WebFlux全流程开发核心实践。
372 0
|
2月前
|
缓存 安全 测试技术
精选API实战问答,解决开发中的疑难杂症
这是一份精选的 API 开发高频问题与解决方案合集,涵盖基础概念、设计规范、安全认证、调试测试、性能优化等十大主题。每日学习 1-2 个问题,结合实战应用,快速掌握 API 开发核心技能,高效解决开发中的各类疑难杂症,适合开发者构建系统化知识体系。
|
2月前
|
人工智能 安全 API
探讨电商API开发全系列
本文系统讲解电商API开发全流程,涵盖基础入门、设计实现、测试调试、性能优化、部署运维等内容,帮助开发者从零成长为电商API架构师。