大家好，今天小白给大家简单分享下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 kernelglobal_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以及设置属性，欢迎一起交流学习。