当前位置: 首页 > news >正文

移动端异构运算技术 - GPU OpenCL 编程(基础篇)

一、前言

随着移动端芯片性能的不断提升,在移动端上实时进行计算机图形学、深度学习模型推理等计算密集型任务不再是一个奢望。在移动端设备上,GPU 凭借其优秀的浮点运算性能,以及良好的 API 兼容性,成为移动端异构计算中非常重要的计算单元。现阶段,在 Android 设备市场,高通 Adreno 和华为 Mali 已经占据了手机 GPU 芯片的主要份额,二者均提供了强劲的 GPU 运算能力。OpenCL,作为 Android 的系统库,在两个芯片上均得到良好的支持。

目前,百度 APP 已经将 GPU 计算加速手段,应用在深度模型推理及一些计算密集型业务上,本文将介绍 OpenCL 基础概念与简单的 OpenCL 编程。
(注:Apple 对于 GPU 推荐的使用方式是 Metal,此处暂不做展开)

二、基础概念

2.1 异构计算

异构计算(Heterogeneous Computing),主要是指使用不同类型指令集和体系架构的计算单元组成系统的计算方式。常见的计算单元类别包括 CPU、GPU 等协处理器、DSP、ASIC、FPGA 等。

2.2 GPU

GPU(Graphics Processing Unit),图形处理器,又称显示核心、显卡、视觉处理器、显示芯片或绘图芯片,是一种专门在个人电脑、工作站、游戏机和一些移动设备(如平板电脑、智能手机等)上执行绘图运算工作的微处理器。传统方式中提升 CPU 时钟频率和内核数量而提高计算能力的方式已经遇到了散热以及能耗的瓶颈。虽然 GPU 单个计算单元的工作频率较低,却具备更多的内核数及并行计算能力。相比于 CPU,GPU 的总体性能 - 芯片面积比,性能 - 功耗比都更高。

三、OpenCL

OpenCL(Open Computing Language)是一个由非盈利性技术组织 Khronos Group 掌管的异构平台编程框架,支持的异构平台涵盖 CPU、GPU、DSP、FPGA 以及其他类型的处理器与硬件加速器。OpenCL 主要包含两部分,一部分是一种基于 C99 标准用于编写内核的语言,另一部分是定义并控制平台的 API。

OpenCL 类似于另外两个开放的工业标准 OpenGL 和 OpenAL ,二者分别用于三维图形和计算机音频方面。OpenCL 主要扩展了 GPU 图形生成之外的计算能力。

3.1 OpenCL 编程模型

使用 OpenCL 编程需要了解 OpenCL 编程的三个核心模型,OpenCL 平台、执行和内存模型。

平台模型(Platform Model)

Platform 代表 OpenCL 视角上的系统中各计算资源之间的拓扑联系。对于 Android 设备,Host 即是 CPU。每个 GPU 计算设备(Compute Device)均包含了多个计算单元(Compute Unit),每个计算单元包含多个处理元素(Processing Element)。对于 GPU 而言,计算单元和处理元素就是 GPU 内的流式多处理器。

执行模型 (Execution Model)

通过 OpenCL 的 clEnqueueNDRangeKernel 命令,可以启动预编译好的 OpenCL 内核,OpenCL 架构上可以支持 N 维的数据并行处理。以二维图片为例,如果将图片的宽高作为 NDRange,在 OpenCL 的内核中可以把图片的每个像素放在一个处理元素上执行,借此可以达到并行化执行的目地。

从上面平台模型部分可以知道,为了提高执行效率,处理器通常会将处理元素分配到执行单元中。我们可以在 clEnqueueNDRangeKernel 中指定工作组大小。同一个工作组中的工作项可以共享本地内存,可以使用屏障(Barriers)去进行同步,也可以通过特定的工作组函数(比如 async_work_group_copy)来进行协作。

内存模型 (Memory Model)

下图中描述了 OpenCL 的内存结构:

  • 宿主内存(Host Memory):宿主 CPU 可直接访问的内存。

  • 全局 / 常量内存 (Global/Constant Memory):可以用于计算设备中的所有计算单元。

  • 本地内存(Local Memory):对计算单元中的所有处理元素可用。

  • 私有内存(Private Memory):用于单个处理元素。

3.2 OpenCL 编程

OpenCL 的编程实际应用中需要一些工程化的封装,本文仅以两个数组相加作为举例,并提供一个简单的示例代码作为参考 ARRAY_ADD_SAMPLE (https://github.com/xiebaiyuan/opencl_cook/blob/master/array_add/array_add.cpp)。

本文将用此作为示例,来阐述 OpenCL 的工作流程。

OpenCL 整体流程主要分为以下几个步骤:

初始化 OpenCL 相关环境,如 cl_device、cl_context、cl_command_queue 等

 cl_int status;
// init deviceruntime.device = init_device();
// create contextruntime.context = clCreateContext(nullptr, 1, &runtime.device, nullptr, nullptr, &status);
// create queueruntime.queue = clCreateCommandQueue(runtime.context, runtime.device, 0, &status);

初始化程序要执行的 program、kernel

 cl_int status;// init programruntime.program = build_program(runtime.context, runtime.device, PROGRAM_FILE);// create kernelruntime.kernel = clCreateKernel(runtime.program, KERNEL_FUNC, &status);

准备输入输出,设置到 CLKernel

 // init datas float input_data[ARRAY_SIZE];float bias_data[ARRAY_SIZE];float output_data[ARRAY_SIZE];for (int i = 0; i < ARRAY_SIZE; i++) {input_data[i] = 1.f * (float) i;bias_data[i] = 10000.f;}// create buffersruntime.input_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), input_data, &status);runtime.bias_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), bias_data, &status);runtime.output_buffer = clCreateBuffer(runtime.context, CL_MEM_READ_ONLY |CL_MEM_COPY_HOST_PTR, ARRAY_SIZE * sizeof(float), output_data, &status);// config cl argsstatus = clSetKernelArg(runtime.kernel, 0, sizeof(cl_mem), &runtime.input_buffer);status |= clSetKernelArg(runtime.kernel, 1, sizeof(cl_mem), &runtime.bias_buffer);status |= clSetKernelArg(runtime.kernel, 2, sizeof(cl_mem), &runtime.output_buffer);

执行获取结果

 // clEnqueueNDRangeKernelstatus = clEnqueueNDRangeKernel(runtime.queue, runtime.kernel, 1, nullptr, &ARRAY_SIZE,nullptr, 0, nullptr, nullptr);// read from outputstatus = clEnqueueReadBuffer(runtime.queue, runtime.output_buffer, CL_TRUE, 0,sizeof(output_data), output_data, 0, nullptr, nullptr);// do with output_data...

四、总结

随着 CPU 瓶颈的到来,GPU 或者其他专用计算设备的编程将是未来的一个重要的技术方向。

相关文章:

移动端异构运算技术 - GPU OpenCL 编程(基础篇)

一、前言 随着移动端芯片性能的不断提升&#xff0c;在移动端上实时进行计算机图形学、深度学习模型推理等计算密集型任务不再是一个奢望。在移动端设备上&#xff0c;GPU 凭借其优秀的浮点运算性能&#xff0c;以及良好的 API 兼容性&#xff0c;成为移动端异构计算中非常重要…...

QString类方法和变量简介(全)

QString类方法和变量简介 操作字符串(|append|insert|sprintf|QString::arg()|prepend|replace|trimmed|simplified)查询字符串(startsWith|endsWith|contains|localeAwareCompare|compare)字符串转换 标准C提供了两种字符串&#xff1a;一种是C语言风格的以"\0"字符…...

中移链控制台对接4A平台功能验证介绍

中移链控制台具备单独的注册登录页面&#xff0c;用户可通过页面注册或者用户管理功能模块进行添加用户&#xff0c;通过个人中心功能模块进行用户信息的修改和密码修改等操作&#xff0c;因业务要求&#xff0c;需要对中移链控制台的用户账号进行集中管理&#xff0c;统一由 4…...

必知的Facebook广告兴趣定位技巧,更准确地找到目标受众

在Facebook广告投放中&#xff0c;兴趣定位是非常重要的一环。兴趣定位不仅可以帮助我们找到我们想要的目标受众&#xff0c;还可以帮助我们避免一些常见的坑。今天&#xff0c;就让我们一起来看看必知的Facebook广告兴趣定位技巧&#xff0c;更准确地找到目标受众。 1.不要只关…...

【MySQL】慢查询+SQL语句优化 (内容源自ChatGPT)

慢查询SQL语句优化 1.什么是慢查询2.优化慢查询3.插入数据优化5.插入数据底层是什么6.页分裂7.页合并8.主键优化方式10.count 优化11.order by优化12.group by 优化13.limit优化14.update 优化15.innodb 三大特征 1.什么是慢查询 慢查询是指执行SQL查询语句所需要的时间较长&a…...

HashMap底层源码解析及红黑树分析

HashMap线程不安全&#xff0c;底层数组链表红黑树 面试重点是put方法&#xff0c;扩容 总结 put方法 HashMap的put方法&#xff0c;首先通过key去生成一个hash值&#xff0c;第一次进来是null&#xff0c;此时初始化大小为16&#xff0c;i (n - 1) & hash计算下标值&a…...

科技云报道:一路狂飙的ChatGPT,是时候被监管了

科技云报道原创。 即使你过去从不关注科技领域&#xff0c;但近期也会被一个由OpenAI&#xff08;美国的一家人工智能公司&#xff09;开发的人工智能聊天机器人“ChatGPT”刷屏。 与上届“全球网红”元宇宙不同&#xff0c;这位新晋的“全能网友”似乎来势更加凶猛。 互联网…...

第四十四章 管理镜像 - 传入日记传输率

文章目录 第四十四章 管理镜像 - 传入日记传输率传入日记传输率镜像数据库状态 第四十四章 管理镜像 - 传入日记传输率 传入日记传输率 在备份和异步成员的镜像成员状态列表下方&#xff0c;自上次刷新镜像监视器以来日志数据从主服务器到达的速率显示在该成员的传入日志传输…...

加密解密学习笔记

加密种类 对称加密&#xff0c;分组对称加密算法 加密算法 AES&#xff08;Advanced Encryption Standard&#xff09;高级加密标准 DES&#xff08;Data Encryption Standard&#xff09;数据加密标准 3DES/Triple DEA (Triple Data Encryption Algorithm) 三重数据加密算…...

Spring 属性填充源码分析(简单实用版)

属性填充 属性填充只有 3 种方式 根据名称填充 根据类型填充 思考什么时候会出现呢&#xff1f;&#xff1f;&#xff1f; 多见于第三方框架与 Spring集成&#xff0c;举例&#xff1a;Mybatis 与 Spring集成&#xff0c;把 Mapper 接口注册为 BeanDefinition 时候就指定了自…...

【机器学习分支】重要性采样(Importance sampling)学习笔记

重要性采样&#xff08;importance sampling&#xff09;是一种用于估计概率密度函数期望值的常用蒙特卡罗积分方法。其基本思想是利用一个已知的概率密度函数来生成样本&#xff0c;从而近似计算另一个概率密度函数的期望值。 想从复杂概率分布中采样的一个主要原因是能够使用…...

三角回文数+123

三角回文数&#xff1a;用户登录 问题描述 对于正整数 n, 如果存在正整数 k 使得 n123⋯kk(k1)/2​, 则 n 称为三角数。例如, 66066 是一个三角数, 因为 66066123⋯363 。 如果一个整数从左到右读出所有数位上的数字, 与从右到左读出所有数位 上的数字是一样的, 则称这个数为…...

JAVA常用的异步处理方法总结

前言 在java项目开发过程中经常会遇到比较耗时的任务&#xff0c;通常是将这些任务做成异步操作&#xff0c;在java中实现异步操作有很多方法&#xff0c;本文主要总结一些常用的处理方法。为了简化&#xff0c;我们就拿一个实际的案例&#xff0c;再用每种方法去实现&#xf…...

GitLab统计代码量

gitlab官方文档&#xff1a;https://docs.gitlab.com/ee/api/index.html 1、生成密钥 登录gitlab&#xff0c;编辑个人资料&#xff0c;设置访问令牌 2、获取当前用户所有可见的项目 接口地址 GET请求 http://gitlab访问地址/api/v4/projects?private_tokenxxx 返回参数 …...

Linux TCP MIB统计汇总

概述 在 linux > 4.7 才将所有TCP丢包收敛到 函数 tcp_drop 中 指标详解 cat /proc/net/netstat 格式化命令 cat /proc/net/netstat | awk (f0) {name$1; i2; while ( i<NF) {n[i] $i; i }; f1; next} (f1){ i2; while ( i<NF){ printf "%s%s %d\n", …...

记录 docker linux部署jar

第一步 web sso user admin 中yml文件还原到阿里mysql数据库 第二步 各个jar进行打包处理 第三步 正式服务器的Jar备份 第四步 拉取以上jar包 到正式服务器中 第五步 查看 docker images 其中 web_service 1.0.2是上一个版本 上一个版本build 镜像命令是这样的&#xff08;需…...

【Linux】教你用进程替换制作一个简单的Shell解释器

本章的代码可以访问这里获取。 由于程序代码是一体的&#xff0c;本章在分开讲解各部分的实现时&#xff0c;代码可能有些跳跃&#xff0c;建议在讲解各部分实现后看一下源代码方便理解程序。 制作一个简单的Shell解释器 一、观察Shell的运行状态二、简单的Shell解释器制作原理…...

onMeasure里如何重置只有1个子view一行满屏, 若有多个自适应一行

onMeasure里如何重置只有1个子view一行满屏, 若有多个自适应一行 可以尝试在 onMeasure 方法中重写 measureChildWithMargins 或 measureChild 方法来实现这个需求。 对于只有一个字的 View,我们可以把它的宽度设为屏幕宽度,高度设为最大高度,这样这个 View 就会占满一整行…...

Postman创建项目 对接口发起请求处理

查看本文之前 您需要理解了解 Postman 的几个简单工作区 如果还没有掌握 可以先查看我的文章 简单认识 Postman界面操作 那么 掌握之后 我们就可以正式来开启我们的接口测试 我们先选择 Collections 我们点上面这个加号 多拉一个项目出来 然后 我们选我们刚加号点出来的项目…...

在Vue3项目中js-cookie库的使用

文章目录 前言1.安装js-cookie库2.引入、使用js-cookie库 前言 今天分享一下在Vue3项目中引入使用js-cookie。 1.安装js-cookie库 js-cookie官网 安装js-cookie&#xff0c;输入 npm i js-cookie安装完成可以在package.json中看到&#xff1a; 安装以后&#xff0c;就可…...

AI-调查研究-01-正念冥想有用吗?对健康的影响及科学指南

点一下关注吧&#xff01;&#xff01;&#xff01;非常感谢&#xff01;&#xff01;持续更新&#xff01;&#xff01;&#xff01; &#x1f680; AI篇持续更新中&#xff01;&#xff08;长期更新&#xff09; 目前2025年06月05日更新到&#xff1a; AI炼丹日志-28 - Aud…...

VTK如何让部分单位不可见

最近遇到一个需求&#xff0c;需要让一个vtkDataSet中的部分单元不可见&#xff0c;查阅了一些资料大概有以下几种方式 1.通过颜色映射表来进行&#xff0c;是最正规的做法 vtkNew<vtkLookupTable> lut; //值为0不显示&#xff0c;主要是最后一个参数&#xff0c;透明度…...

鸿蒙DevEco Studio HarmonyOS 5跑酷小游戏实现指南

1. 项目概述 本跑酷小游戏基于鸿蒙HarmonyOS 5开发&#xff0c;使用DevEco Studio作为开发工具&#xff0c;采用Java语言实现&#xff0c;包含角色控制、障碍物生成和分数计算系统。 2. 项目结构 /src/main/java/com/example/runner/├── MainAbilitySlice.java // 主界…...

Kafka入门-生产者

生产者 生产者发送流程&#xff1a; 延迟时间为0ms时&#xff0c;也就意味着每当有数据就会直接发送 异步发送API 异步发送和同步发送的不同在于&#xff1a;异步发送不需要等待结果&#xff0c;同步发送必须等待结果才能进行下一步发送。 普通异步发送 首先导入所需的k…...

通过 Ansible 在 Windows 2022 上安装 IIS Web 服务器

拓扑结构 这是一个用于通过 Ansible 部署 IIS Web 服务器的实验室拓扑。 前提条件&#xff1a; 在被管理的节点上安装WinRm 准备一张自签名的证书 开放防火墙入站tcp 5985 5986端口 准备自签名证书 PS C:\Users\azureuser> $cert New-SelfSignedCertificate -DnsName &…...

恶补电源:1.电桥

一、元器件的选择 搜索并选择电桥&#xff0c;再multisim中选择FWB&#xff0c;就有各种型号的电桥: 电桥是用来干嘛的呢&#xff1f; 它是一个由四个二极管搭成的“桥梁”形状的电路&#xff0c;用来把交流电&#xff08;AC&#xff09;变成直流电&#xff08;DC&#xff09;。…...

WEB3全栈开发——面试专业技能点P7前端与链上集成

一、Next.js技术栈 ✅ 概念介绍 Next.js 是一个基于 React 的 服务端渲染&#xff08;SSR&#xff09;与静态网站生成&#xff08;SSG&#xff09; 框架&#xff0c;由 Vercel 开发。它简化了构建生产级 React 应用的过程&#xff0c;并内置了很多特性&#xff1a; ✅ 文件系…...

sshd代码修改banner

sshd服务连接之后会收到字符串&#xff1a; SSH-2.0-OpenSSH_9.5 容易被hacker识别此服务为sshd服务。 是否可以通过修改此banner达到让人无法识别此服务的目的呢&#xff1f; 不能。因为这是写的SSH的协议中的。 也就是协议规定了banner必须这么写。 SSH- 开头&#xff0c…...

小智AI+MCP

什么是小智AI和MCP 如果还不清楚的先看往期文章 手搓小智AI聊天机器人 MCP 深度解析&#xff1a;AI 的USB接口 如何使用小智MCP 1.刷支持mcp的小智固件 2.下载官方MCP的示例代码 Github&#xff1a;https://github.com/78/mcp-calculator 安这个步骤执行 其中MCP_ENDPOI…...

32位寻址与64位寻址

32位寻址与64位寻址 32位寻址是什么&#xff1f; 32位寻址是指计算机的CPU、内存或总线系统使用32位二进制数来标识和访问内存中的存储单元&#xff08;地址&#xff09;&#xff0c;其核心含义与能力如下&#xff1a; 1. 核心定义 地址位宽&#xff1a;CPU或内存控制器用32位…...