啥都吃的豆芽 · 2021年01月11日

Mali SDK支持为OpenCL用C ++编写的内核编译

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的支持有任何反馈,请与我们联系

推荐阅读
关注数
23565
内容数
1018
Arm相关的技术博客,提供最新Arm技术干货,欢迎关注
目录
极术微信服务号
关注极术微信号
实时接收点赞提醒和评论通知
安谋科技学堂公众号
关注安谋科技学堂
实时获取安谋科技及 Arm 教学资源
安谋科技招聘公众号
关注安谋科技招聘
实时获取安谋科技中国职位信息