通过OpenCL的cl_qcom_ml_ops（CLML）扩展，我们为开发人员提供了一套常用的machine learning ops，可以在高通平台的Adreno GPU上获得最佳性能。对于希望创建自定义操作的开发人员，CLML提供了添加自定义操作的接口。

在高通 Adreno CLML SDK 2.1中，CLML_mobilenet_custom_kernel示例演示了如何使用CLML Ops、自定义OpenCL内核以及ML操作来构建和运行mobilenet图像分类网络。

自定义操作首先需要编写OpenCL内核，再将其与CLML操作内联起来，最后将所有操作分发到gpu进行运算。我们将这个OpenCL内核称为自定义内核（custom kernel）。

自定义内核涉及到kernel的输入输出，也就是tensor。这里tensor layout描述了tensor中数据的布局方式。布局可以是CL_TENSOR_LAYOUT_NCHW_QCOM、CL_TENSOR_ LAYOUT _NHWC_QCCOM或CL_TENSOR_LAYOUT_OPTIMAL_QCOM中的一种。NCHW和NHWC是线性排列的，而OPTIMAL方式则是按照块状重新组织数据（例如4x4），这可以节省Adreno GPU硬件访问时间。

自定义内核是内联在CLML ops中的，所以需要接收上一层CLMLops的tensor数据作为输入，还要将计算结果保存至输出tensor。要将tensor数据从CLML Op传递到自定义内核，并从自定义内核传递回CLML Op，有两种方法：

1. 第一种方法是将CLML Op的tensor布局切换为CL_TENSOR_LAYOUT_NCHW_QCOM或CL_TENSOR_LAYOUT_NHWC_QCOM。CLML Op的输出/输入tensor创建时使用CL_TENSOR_LAYOUT_NCHW_QCOM或CL_TENSOR_LAYOUT_NHWC_QCOM之一；与输出/输入tensor关联的OpenCL  buffer对象可以作为参数直接传递给自定义内核。QCOM CLML SDK 1中有一个示例（\src\examples\mobilenet_custom_kernel）就使用了这个方式。
2. 第二种方法是使用CL_TENSOR_LAYOUT_OPTIMAL_QCOM，并利用新添加的内置函数读取和写入tensor。

首先我们为tensor创建内核参数对象，这里使用clCreateKernelArgMLTensorQCOM API：

 // Create tensor memory descriptor for custom layer output
    cl_ml_tensor_memory_desc_qcom customReluOutputTensorMemDesc ={
                customReluOptimalOutputTensor,
                customReluOutputMem
     };
    // API function that generates kernel argument from an existing clml tensor and cl memory pair.
    result=h_ClmlIntf->clCreateKernelArgMLTensorQCOM(context, &customReluInputTensorMemDesc, &tensorArgInput);

创建的ML tensor参数可以直接使用OpenCL clSetKernelArg() API设置，释放的时候使用OpenCL clReleaseMemObject() API。

自定义内核里使用内置函数qcom_read_tensor*/qcom_ write_tensor*来访问每个tensor参数的数据。为了使用这些内置函数，必须启用基础功能cl_qcom_tensor

   #pragma OPENCL EXTENSION cl_qcom_tensor: enable 

在QCOM CLML SDK2.0中，clml_mobilenet_custom_kernel展示了如何使用内置函数：

  "__kernel void custom_relu_with_optimal_activation(\n"
    "                      __read_only  qcom_tensor_t  input,\n"
    "                      __write_only qcom_tensor_t  output)\n"
    "{\n"
    "    short h = get_global_id(0);\n"
    "    short w = get_global_id(1);\n"
    "    short c = get_global_id(2);\n"
    "\n"
    "    int4 coord = (int4)(w, h, c, 0);\n"
    "    half input_value  = qcom_read_tensorh(input, coord); \n"
    "    half output_value = max(0.0h, input_value);"
    "    qcom_write_tensorh(output, coord, output_value); \n"
"}";

在示例中我们自定义创建mobilenet网络中的第三层，这一层是Relu内核。自定义内核中的“coord”是4D  tensor坐标：coord xyzw对应于tensor的W/H/C/N维。在clml_mobilenet_custom_kernel中，mobilenet输入是一个224x224 RGB888图像，因此coord.w是0。

自定义内核作为普通的OpenCL计算内核进行编译，然后用clEnqueueNDRangeKernel分发到gpu任务队列，而ML op则使用clEnqueueMLOpQCOM分发。详细信息可以参考OpenCL ML SDK2.1。现在您可以根据需要定制ML模型，同时更接近硬件，获取性能优势。

其他相关内容：

边缘机器学习训练：移动设备端训练

利用高通openCL ML SDK加速机器学习模型

高通OpenCL ML扩展之可记录队列

作者：Ya Kong