Hetcompute sdk中的kernel介绍

大家好，今天小白给大家简单分享下Hetcompute sdk中task相关的基础知识，欢迎一起交流学习。

一、Hetcompute sdk中的kernel

HetCompute任务包含可在系统中的任何设备上执行的工作：CPU，GPU或Qualcomm Hexagon DSP。 这使得HetCompute程序员可以编写由各种任务组成的应用程序，充分利用现代计算系统中可用的异构设备的性能和功效。
HetCompute任务使用内核来实现这种计算异构性。 内核包含计算，即实际的设备代码，任务执行。 这可能是CPU代码，GPU代码或Qualcomm Hexagon DSP代码，导致三种不同类型的内核。 在当前的HetCompute版本中，每个任务只包含一个内核，指示此任务在哪个设备上执行。

二、如何创建kernels

目前，HetCompute中有三种类型的内核。 使用以下方法创建每一种内核：

1、创建CPU kernel

要创建CPU内核，使用hetcompute :: create_cpu_kernel，它需要函数或函数对象（即lambda表达式或仿函数）。

2、创建GPU kernel

要创建GPU内核，使用hetcompute :: create_gpu_kernel，这里有两种：一个用于包装OpenCL C内核，另外一个用于OpenGL ES计算着色器。

OpenCL变体的参数是：

a、包含OpenCL GPU设备函数源的字符串；

b、包含设备函数名称的字符串

OpenGL变量将包含计算着色器程序源的字符串作为参数；

3、创建DSP kernel

要创建DSP内核，请使用hetcompute :: create_dsp_kernel。 它需要兼容DSP的C功能

4、使用以上API创建三种kernel

#include <hetcompute/hetcompute.hh>
static int  f1(int x) {
 return x * 2;
}

static auto f2 = [](int x) -> int { return x * 2; };

static struct
{
 int operator()(int x) { return x * 2; }
} f3;

// Source string for an OpenCL C kernel

static std::string const f4_string = "__kernel void f4(__global int *x, __global int *y) {"
" int i = get_global_id(0);"
" y[i] = x[i];"
"}";

static int f5(int* x, int* y, int l)
{
 int i = 0;
 for (i = 0; i < l; i++)
 y[i] = x[i] * 2;
 return 0;
}
int main()
{
 hetcompute::runtime::init();

 // Create a cpu_kernel from a function
 auto k1 = hetcompute::create_cpu_kernel(f1);

 // Create a cpu_kernel from a lambda expression
 auto k2 = hetcompute::create_cpu_kernel(f2);

 // Create a cpu_kernel from a functor
 auto k3 = hetcompute::create_cpu_kernel(f3);

 // Create a gpu_kernel from an OpenCL C GPU function
 auto k4 = hetcompute::create_gpu_kernel<hetcompute::buffer_ptr<int>, hetcompute::buffer_ptr<int>>(f4_string, "f4");

 // Create a hexagon_kernel from a DSP function

 auto k5 = hetcompute::create_dsp_kernel<>(f5);

 hetcompute::runtime::shutdown();
 return 0;
}

创建内核后，hetcompute :: create_task可以使用它来创建任务。 此外，内核可用于创建多个独立任务。

以上展示了从OpenCL C GPU函数创建gpu_kernel，下面介绍使用OpenGL ES创建GPU kernel

#include <hetcompute/hetcompute.hh>
#define LOCAL_SIZE 16 // This should match the local_size_x value in the shader
const char* shader_code = R"GLCODE(
#version 310 es
precision highp float;
layout(local_size_x = 16) in;
layout(std430) buffer;
layout(binding = 2) writeonly buffer Output {
float elements[];
} output_data;
layout(binding = 0) readonly buffer Input0 {
float elements[];
} input_data0;layout(binding = 1) readonly buffer Input1 {
float elements[];
} input_data1;
void main()
{
uint ident = gl_GlobalInvocationID.x;
output_data.elements[ident] = input_data0.elements[ident] + input_data1.elements[ident];
}
)GLCODE";

int main()
{
hetcompute::runtime::init();
auto buf_a = hetcompute::create_buffer<float>(1024);
auto buf_b = hetcompute::create_buffer<float>(buf_a.size());

buf_a.acquire_wi();
buf_b.acquire_wi();
// Initialize the input vectors
for (size_t i = 0; i < buf_a.size(); ++i)
{
buf_a[i] = i;
buf_b[i] = buf_a.size() - i;
}
buf_a.release();
buf_b.release();
auto buf_c = hetcompute::create_buffer<float>(buf_a.size());
auto gl_vadd = hetcompute::beta::create_gpu_kernel<hetcompute::in<hetcompute::buffer_ptr<float>>,
                                                                                      hetcompute::in<hetcompute::buffer_ptr<float>>,
                                                                                      hetcompute::out<hetcompute::buffer_ptr<float>>>(
hetcompute::beta::gl, shader_code);
hetcompute::range<1> global_range(buf_a.size());
hetcompute::range<1> local_range(LOCAL_SIZE);
// Create a task
auto gpu_task = hetcompute::create_task(gl_vadd,  // gpu kernel

global_range,    // global range
local_range,      // local range
buf_a,               //
parameters       // rest correspond to gpu_vadd template parameters
buf_b,          
buf_c);           

gpu_task->launch();

gpu_task->wait_for();
hetcompute::runtime::shutdown();
}

相应地，当传递OpenCL C函数时，用户可以选择通过
hetcompute :: beta :: cl使OpenCL显式，如下图所示。
auto k4 = hetcompute :: beta :: create_gpu_kernel <hetcompute :: buffer_ptr <int>，hetcompute::buffer_ptr<INT>>（hetcompute :: beta :: cl，f4_string，“f4”）;

三、设置kernel属性

从CPU、GPU、DSP函数创建任务之后，Qualcomm HetCompute用户无法将内部函数替换为内核中的另一个函数。 但是，用户可以更改内核的以下属性：

1、阻塞（blocking）：表示是否希望从此内核生成的CPU任务阻止外部事件（如I / O活动（阻塞任务））。 可以使用内核对象的set_blocking和is_blocking方法设置和查询blocking属性。

2、Big：如在需要大量计算的场合下将负载放到大核上以提高效率，当负载较小时可以关闭大核，将任务放到小核上运行以降低功耗。使用内核对象的set_big和is_big方法设置和查询big属性；（使用set_little和is_little设置和查询little）

3、用法示例

#include <hetcompute/hetcompute.hh>
int
main()
{
hetcompute::runtime::init();

auto k1 = hetcompute::create_cpu_kernel([] { HETCOMPUTE_ILOG("big task
executed"); });

// inform the Hetcompute runtime that the kernel is best executed on a big
// core in a big.LITTLE SoC
k1.set_big();
auto t1 = hetcompute::launch(k1);
t1->wait_for();

auto k2 = hetcompute::create_cpu_kernel([] { HETCOMPUTE_ILOG("LITTLE task
executed"); });

// inform the Hetcompute runtime that the kernel is best executed on a
// LITTLE core in a big.LITTLE SoC
k2.set_little();
auto t2 = hetcompute::launch(k2);
t2->wait_for();

hetcompute::runtime::shutdown();
return 0;

 }

四、总结

本篇主要是简单介绍了如何创建CPU、GPU、DSP kernel以及设置属性，欢迎一起交流学习。