Mali软件开发套件(SDK)促进了基于Mali平台的OpenGL ES,Vulkan或OpenCL应用程序的开发。OpenGL和Vulkan主要用于图形处理,而OpenCL提供了在Mali GPU上从各个域执行应用程序的功能,例如图像处理,机器学习(ML),计算机视觉(CV)等。OpenCL的主要概念是在具有高计算能力的设备上加速数据并行计算。如果应用程序适合OpenCL加速,则与传统执行相比,加速的增加可以是几个数量级。任何OpenCL程序都具有主机端,通常在CPU上执行常规计算,而主机端则是从主机发送计算密集型内核以进行加速的设备端。 在主机方面,开发人员可以使用标准的早期版本中的C,C ++甚至Python。但是,必须使用源自C99的OpenCL C编写设备上加速的内核核心。
去年,在上游LLVM中开发了用于OpenCL的新内核语言C ++。这允许在OpenCL内核代码中使用大多数C ++ 17功能-有关此功能的更多详细信息,请参见前面的博客。可以使用开源工具离线编译C ++ for OpenCL语言的内核,而SPIR-V中间格式的二进制文件可以由与现有OpenCL驱动程序一起运行的OpenCL应用程序导入。现在,我们很高兴地宣布,在最新版本的Mali SDK中,开发人员可以使用cl_ext_cxx_for_opencl扩展程序,从在线和离线编译用C ++编写的OpenCL内核中受益。这已在Khronos网站上最近发布。
Arm是第一个在其SDK中提供对此新扩展的支持的供应商。虽然这对开发人员来说是个好消息,但仍处于试验阶段。在发布此博客时,尚不支持某些功能,例如具有非平凡导体或析构函数的程序范围对象。
C ++ for OpenCL
OpenCL的C++提供了在OpenCL内核中使用C++17标准中的大多数现代C++特性的能力。随着GPU上运行的复杂性不断增加,这提高了应用程序的程序员工作效率。另一个巨大的好处是,用于OpenCL的C++向后兼容OpenCL的C版本。这意味着现有的应用程序可以顺利地迁移到C++特性,开发人员可以继续使用熟悉的OpenCL编程概念和工具。总体而言,用C++为OpenCL编写的内核代码看起来与用OpenCL C编写的代码一样,但为了方便起见,还提供了一些额外的C++特性。
用于OpenCL内核的C++可以利用许多C++特定的编译器优化,从而在OpenCL设备上获得具有竞争力的性能。社区定义的语言文档与OpenCL的其他Khronos规范一起托管在GitHub上的OpenCL-Docs存储库中。可在此处找到最新发布的修订版。
您可以从幻灯片PDF和今年在IWOCL上展示的演讲视频中找到更多关于C++for OpenCL编程语言的信息。
使用C++功能实现内核
以下代码段说明了如何使用C ++功能以复数算法实现内核。完整的示例可以在Code Explorer中找到。
// Define a class - Complex, that can perform complex number arithmetic
// with various precision when different types for ‘T’ are used - double, float, half...
template<typename T>
class complex_t {
T m_re; // Real component.
T m_im; // Imaginary component.
public:
complex_t(T re, T im): m_re{re}, m_im{im}{};
complex_t operator*(complex_t &other)
{
return {m_re * other.m_re - m_im * other.m_im,
m_re * other.m_im + m_im * other.m_re};
}
int get_re() { return m_re;}
int get_im() { return m_im;}
};
// A kernel function to compute multiplication over complex numbers read from
// the input buffer and to store the result into the output buffer.
kernel void compute_helper(global float *in, global float *out) {
auto idx = get_global_id(0);
// Every work-item uses 4 consecutive items from the input buffer -
// two for each complex number.
auto offset = idx * 4;
complex_t num1{in[offset], in[offset + 1]};
complex_t num2{in[offset + 2], in[offset + 3]};
// Perform complex number multiplication.
complex_t res = num1 * num2;
// Every work-item writes 2 consecutive items to the output buffer.
out[idx * 2] = res.get_re();
out[idx * 2 + 1] = res.get_im();
}
关于离线编译
从2019年9月开始,用C++for OpenCL编写的内核的离线编译已通过开源工具提供。这是在上游的llvm-project中,在C++for OpenCL的实验性支持下发布Clang 9.0的时候。内核可以按照Khronos网站上这篇博客中解释的流程脱机编译为SPIR-V格式。然后,运行在OpenCL 2.0或更高版本驱动程序上的OpenCL应用程序可以使用常规clCreateProgramWithIL API调用加载SPIR-V二进制文件。
开发人员,自己尝试
我们建议开发人员在其应用程序中尝试使用C ++ for OpenCL内核语言。如果您对新的内核语言,其在工具中的支持或总体上对Mali SDK的支持有任何反馈,请与我们联系。