当前位置: 首页 > 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;就可…...

【论文笔记】Attention和Visual Transformer

Attention和Visual Transformer Attention和Transformer为什么需要AttentionAttention机制Multi-head AttentionSelf Multi-head Attention&#xff0c;SMA TransformerVisual Transformer&#xff0c;ViT Attention和Transformer Attention机制在相当早的时间就已经被提出了&…...

独立IP服务器和共享IP服务器有什么区别

在选择一个合适的服务器时&#xff0c;最常见的选择是共享IP服务器和独立IP服务器。尽管两者看起来很相似&#xff0c;但它们有着很大的不同。本文将详细介绍共享IP服务器和独立IP服务器的不同之处&#xff0c;以及如何选择适合您需求的服务器。 一、什么是共享IP服务器? 共享…...

Java8

Java8 &#xff08;一&#xff09;、双列集合&#xff08;二&#xff09;、Map集合常用api&#xff08;三&#xff09;、Map集合的遍历方式&#xff08;四&#xff09;、HashMap&#xff08;五&#xff09;、LinkedHashMap&#xff08;六&#xff09;、TreeMap&#xff08;七&a…...

nn.conv1d的输入问题

Conv1d(in_channels, out_channels, kernel_size, stride1, padding0, dilation1, groups1, biasTrue) in_channels(int) – 输入信号的通道。在文本分类中&#xff0c;即为词向量的维度out_channels(int) – 卷积产生的通道。有多少个out_channels&#xff0c;就需要多少个1维…...

js判断是否为null,undefined,NaN,空串或者空对象

js判断是否为null&#xff0c;undefined&#xff0c;NaN&#xff0c;空串或者空对象 这里写目录标题 js判断是否为null&#xff0c;undefined&#xff0c;NaN&#xff0c;空串或者空对象特殊值nullundefinedNaN空字符串&#xff08;""&#xff09;空对象&#xff08;…...

Java每日一练(20230501)

目录 1. 路径交叉 &#x1f31f;&#x1f31f; 2. 环形链表 &#x1f31f;&#x1f31f; 3. 被围绕的区域 &#x1f31f;&#x1f31f; &#x1f31f; 每日一练刷题专栏 &#x1f31f; Golang每日一练 专栏 Python每日一练 专栏 C/C每日一练 专栏 Java每日一练 专栏…...

从零开始学习Web自动化测试:如何使用Selenium和Python提高效率?

B站首推&#xff01;2023最详细自动化测试合集&#xff0c;小白皆可掌握&#xff0c;让测试变得简单、快捷、可靠https://www.bilibili.com/video/BV1ua4y1V7Db 目录 引言&#xff1a; 一、了解Web自动化测试的基本概念 二、选择Web自动化测试工具 三、学习Web自动化测试的…...

fastdfs环境搭建

安装包下载路径 libfastcommon下载地址&#xff1a;https://github.com/happyfish100/libfastcommon/releasesFastDFS下载地址&#xff1a;https://github.com/happyfish100/fastdfs/releasesfastdfs-nginx-module下载地址&#xff1a;https://github.com/happyfish100/fastdf…...

有什么牌子台灯性价比高?性价比最高的护眼台灯

由心感叹现在的孩子真不容易&#xff0c;学习压力比我们小时候大太多&#xff0c;特别是数学&#xff0c;不再是简单的计算&#xff0c;而更多的是培养学生其他思维方式&#xff0c;有时候我都觉得一年级数学题是不是超纲了。我女儿现在基本上都是晚上9点30左右上床睡觉&#x…...

信息系统项目管理师 第9章 项目范围管理

1.管理基础 1.产品范围和项目范围 产品范围:某项产品、服务或成果所具有的特征和功能。根据产品需求来衡量。 项目范围:包括产品范围&#xff0c;是为交付具有规定特性与功能的产品、服务或成果而必须完成的工作。项目管理计划来衡量 2.管理新实践 更加注重与商业分析师一起…...