You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/08/10 18:02:25 UTC
[GitHub] [tvm] csullivan commented on a diff in pull request #12204: [Hexagon] Add Hand written HVX conv2d
csullivan commented on code in PR #12204:
URL: https://github.com/apache/tvm/pull/12204#discussion_r942612681
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
Review Comment:
Here also as below
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
Review Comment:
"splatted" == broadcast?
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
Review Comment:
`blockized` means the data has been packed into an array of non-contiguous allocations
`chunkified` means weights are in a contiguous packed layout
For chunkified maybe we just refer to `packed` instead, and define the meaning of `blocked` somewhere in the docstring.
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
+ * @param zero_block A block filled with zeros
+ *
+ * @return
+ */
+void conv_layer_fp16_hvx(DLTensor& cr_out, const DLTensor& cr_act, // NOLINT(*)
+ const DLTensor& cr_filt, const DLTensor& out_shape,
+ const DLTensor& act_shape, const DLTensor& bias_flat,
+ const DLTensor& filt_shape, const DLTensor& pad_shape, bool relu,
+ int stride_h, int stride_w, uintptr_t zero_block) {
+ int64_t filt_height = filt_shape.shape[0];
+ int64_t filt_width = filt_shape.shape[1];
+ int64_t filt_idepth = filt_shape.shape[2];
+ (void)filt_idepth;
+
+ DEBUG_BLOCK(int pad_top = pad_shape.shape[0]; int pad_left = pad_shape.shape[1];)
+
+ debug("filt_height=%" PRId64 ", filt_width=%" PRId64 ", filt_idepth=%" PRId64
Review Comment:
Use the runtime logging system, LOG(INFO).
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
+ * @param zero_block A block filled with zeros
+ *
+ * @return
+ */
+void conv_layer_fp16_hvx(DLTensor& cr_out, const DLTensor& cr_act, // NOLINT(*)
+ const DLTensor& cr_filt, const DLTensor& out_shape,
+ const DLTensor& act_shape, const DLTensor& bias_flat,
+ const DLTensor& filt_shape, const DLTensor& pad_shape, bool relu,
+ int stride_h, int stride_w, uintptr_t zero_block) {
+ int64_t filt_height = filt_shape.shape[0];
+ int64_t filt_width = filt_shape.shape[1];
+ int64_t filt_idepth = filt_shape.shape[2];
+ (void)filt_idepth;
Review Comment:
```suggestion
```
##########
src/runtime/hexagon/ops/conv_utils.cc:
##########
@@ -0,0 +1,191 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include "conv2d.h"
+
+namespace detail {
+
+void blockize_hwc_16b(void* out, void* inp_flat, int height, int width, int depth) {
+ auto inp_data = static_cast<uint16_t*>(inp_flat);
+ auto out_data = static_cast<uintptr_t*>(out);
+ const int stride_x = depth;
+ const int stride_y = stride_x * width;
+
+ for (int cy = 0; cy < height; cy += 8) {
+ for (int cx = 0; cx < width; cx += 4) {
+ for (int cc = 0; cc < depth; cc += 32) {
+ auto block = reinterpret_cast<uint16_t*>(*out_data++);
+ int max_y = std::min(8, height - cy);
+ int max_x = std::min(4, width - cx);
+ int max_c = std::min(32, depth - cc);
+ for (int y = 0; y < max_y; ++y) {
+ for (int x = 0; x < max_x; ++x) {
+ for (int c = 0; c < max_c; ++c) {
+ block[xyc_to_sm_16b(y, x, c)] =
+ inp_data[(cy + y) * stride_y + (cx + x) * stride_x + (cc + c)];
+ }
+ for (int c = max_c; c < 32; ++c) block[xyc_to_sm_16b(y, x, c)] = 0;
+ }
+ for (int x = max_x; x < 4; ++x) {
+ for (int c = 0; c < 32; ++c) block[xyc_to_sm_16b(y, x, c)] = 0;
+ }
+ }
+
+ for (int y = max_y; y < 8; ++y)
+ for (int x = 0; x < 4; ++x)
+ for (int c = 0; c < 32; ++c) block[xyc_to_sm_16b(y, x, c)] = 0;
+ } // cc
+ } // cx
+ } // cy
+}
+
+void deblockize_hwc_16b(void* out_flat, void* inp, int height, int width, int depth) {
+ uintptr_t* inp_data = static_cast<uintptr_t*>(inp);
+ uint16_t* out_data = static_cast<uint16_t*>(out_flat);
+ const int stride_x = depth;
+ const int stride_y = stride_x * width;
+
+ for (int cy = 0; cy < height; cy += 8) {
+ for (int cx = 0; cx < width; cx += 4) {
+ for (int cc = 0; cc < depth; cc += 32) {
+ auto block = reinterpret_cast<uint16_t*>(*inp_data);
+ int max_y = std::min(8, height - cy);
+ int max_x = std::min(4, width - cx);
+ int max_c = std::min(32, depth - cc);
+ for (int y = 0; y < max_y; ++y) {
+ for (int x = 0; x < max_x; ++x) {
+ for (int c = 0; c < max_c; ++c) {
+ out_data[(cy + y) * stride_y + (cx + x) * stride_x + (cc + c)] =
+ block[xyc_to_sm_16b(y, x, c)];
+ }
+ }
+ }
+
+ inp_data++;
+ }
+ }
+ }
+}
+
+void chunkify_hwio_16b(void** out_ptr, int out_ptr_size, void* out, void* inp, int height,
+ int width, int idepth, int odepth) {
+ auto inp_data = static_cast<uint16_t*>(inp);
+ auto out_data = static_cast<uintptr_t*>(out);
+ const int stride_i = odepth;
+ const int stride_x = stride_i * idepth;
+ const int stride_y = stride_x * width;
+
+ for (int cy = 0; cy < height; cy += 8) {
+ // In the chunkified tensor, the chunks are ordered in increasing
+ // x order, but they start from the thin one.
+ for (int cx = width - round_up(width, 4); cx < width; cx += 4) {
+ int cx0 = std::max(0, cx);
+ for (int ci = 0; ci < idepth; ci += 32) {
+ for (int co = 0; co < odepth; co += 32) {
+ int max_y = std::min(8, height - cy);
+ int max_x = std::min(4, cx + 4 - cx0);
+ int max_i = std::min(32, idepth - ci);
+ int max_o = std::min(32, odepth - co);
+
+ auto chunk = reinterpret_cast<uint16_t*>(out_data);
+ for (int y = 0; y < max_y; ++y) {
+ for (int x = max_x - 1; x >= 0; --x) {
+ for (int i = 0; i < max_i; ++i) {
+ for (int o = 0; o < max_o; ++o) {
+ debug(
+ "cy: %d, cx: %d, cx0: %d, ci: %d, co: %d, max_x: %d, y: %d, x: %d, i: %d, o: "
+ "%d, index: %d",
+ cy, cx, cx0, ci, co, max_x, y, x, i, o, hwio_to_sm_16b(max_x, y, x, i, o));
+ chunk[hwio_to_sm_16b(max_x, y, x, i, o)] =
+ inp_data[(cy + y) * stride_y + (cx0 + x) * stride_x + (ci + i) * stride_i +
+ (co + o)];
+ }
+ for (int o = max_o; o < 32; ++o) chunk[hwio_to_sm_16b(max_x, y, x, i, o)] = 0;
+ }
+ for (int i = max_i; i < 32; ++i)
+ for (int o = 0; o < 32; ++o) chunk[hwio_to_sm_16b(max_x, y, x, i, o)] = 0;
+ }
+ }
+
+ *out_ptr++ = chunk;
+ out_data += max_y * max_x * 32 * 32;
+ out_ptr_size--;
+ assert(out_ptr_size >= 0);
+ }
+ }
+ }
+ }
+}
+
+SDLTensor<4> prepare_nhwc(tvm::runtime::DeviceAPI* device_api, const DLTensor* nhwc_flat,
+ bool copy_data) {
+ tvm::runtime::String vtcm_scope = "global.vtcm";
+
+ // Allocate blocks for activations. We will use the block pointers
+ // directly from the allocated area.
+ int n = nhwc_flat->shape[0];
+ int h = round_up(nhwc_flat->shape[1], 8);
+ int w = round_up(nhwc_flat->shape[2], 4);
+ int c = round_up(nhwc_flat->shape[3], 32);
+ int64_t shape_2d[2] = {(n * h * w * c) / (8 * 4 * 32), 8 * 4 * 32};
+ void* nhwc_vtcm =
+ device_api->AllocDataSpace(hexagon_device, 2, shape_2d, nhwc_flat->dtype, vtcm_scope);
+ if (copy_data) {
+ blockize_hwc_16b(nhwc_vtcm, nhwc_flat->data, nhwc_flat->shape[1], nhwc_flat->shape[2],
+ nhwc_flat->shape[3]);
+ }
+
+ return SDLTensor<4>(nhwc_vtcm, nhwc_flat->dtype, nhwc_vtcm, {n, h / 8, w / 4, c / 32});
+}
+
+SDLTensor<4> prepare_hwio(tvm::runtime::DeviceAPI* device_api, const DLTensor* hwio_flat,
+ int num_chunks, void** ptr_table) {
+ tvm::runtime::String vtcm_scope = "global.vtcm";
+
+ // Allocate one block for filter data. We will need to create our own
+ // pointer table. The reason is that filter chunks cannot be padded
Review Comment:
Why can they not be padded in height or width?
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
+ * @param zero_block A block filled with zeros
+ *
+ * @return
+ */
+void conv_layer_fp16_hvx(DLTensor& cr_out, const DLTensor& cr_act, // NOLINT(*)
+ const DLTensor& cr_filt, const DLTensor& out_shape,
+ const DLTensor& act_shape, const DLTensor& bias_flat,
+ const DLTensor& filt_shape, const DLTensor& pad_shape, bool relu,
+ int stride_h, int stride_w, uintptr_t zero_block) {
+ int64_t filt_height = filt_shape.shape[0];
+ int64_t filt_width = filt_shape.shape[1];
+ int64_t filt_idepth = filt_shape.shape[2];
+ (void)filt_idepth;
+
+ DEBUG_BLOCK(int pad_top = pad_shape.shape[0]; int pad_left = pad_shape.shape[1];)
Review Comment:
What is the purpose of DEBUG_BLOCK?
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
+ * @param zero_block A block filled with zeros
+ *
+ * @return
+ */
+void conv_layer_fp16_hvx(DLTensor& cr_out, const DLTensor& cr_act, // NOLINT(*)
+ const DLTensor& cr_filt, const DLTensor& out_shape,
+ const DLTensor& act_shape, const DLTensor& bias_flat,
+ const DLTensor& filt_shape, const DLTensor& pad_shape, bool relu,
+ int stride_h, int stride_w, uintptr_t zero_block) {
+ int64_t filt_height = filt_shape.shape[0];
+ int64_t filt_width = filt_shape.shape[1];
+ int64_t filt_idepth = filt_shape.shape[2];
+ (void)filt_idepth;
+
+ DEBUG_BLOCK(int pad_top = pad_shape.shape[0]; int pad_left = pad_shape.shape[1];)
+
+ debug("filt_height=%" PRId64 ", filt_width=%" PRId64 ", filt_idepth=%" PRId64
+ ", pad_top=%d, pad_left=%d\n",
+ filt_height, filt_width, filt_idepth, pad_top, pad_left);
+
+ assert(pad_top < 8 && pad_left < 4);
+
+ DEBUG_BLOCK(int a_height = cr_act.shape[1]; int a_width = cr_act.shape[2];
+ int a_depth = cr_act.shape[3];
+
+ int w_height = cr_filt.shape[0]; int w_width = cr_filt.shape[1];
+
+ int o_depth = cr_out.shape[3]; int b_depth = bias_flat.shape[0];)
+
+ int o_height = cr_out.shape[1];
+ int o_width = cr_out.shape[2];
+
+ int out_height = out_shape.shape[1];
+ int out_width = out_shape.shape[2];
+
+ debug("a: %dx%dx%dx%d, w: %dx%dx%dx%d, o: %dx%dx%dx%d, b: %d, out_shape: %dx%d\n", 1, a_height,
+ a_width, a_depth, w_height, w_width, static_cast<int>(cr_filt.shape[2]),
+ static_cast<int>(cr_filt.shape[3]), 1, o_height, o_width, o_depth, b_depth, out_height,
+ out_width);
+
+ assert(a_depth == cr_filt.shape[2]);
+ assert(o_depth == cr_filt.shape[3]);
+
+ int rd = round_down(filt_width, 4);
+ int wgt_chunk_thin_width = filt_width - rd;
+
+ /*
+ * Compute the output vector of either 1 or 2 elements along the width and max 32 elements along
+ * the depth to constitue a maximum of 64 elements
+ *
+ * The weights are loaded directly in the order they're stored, which results
+ * in 2 input channels and 32 output channels
+ *
+ * Weights vector illustration:
+ * ------- ------ ------------
+ * weights_vec = [0-0,0-1,1-0,1-1,2-0,2-1,3-0,3-1,4-0,4-1,...,31-0,31-1] -> This is the
+ * vector representation of weights, where the elements are represented as
+ * "out_channel-input_channel"
+ *
+ *
+ * Same 2 input channels have to be multiplied across all output channels in the weights.
+ *
+ * Activations vector would thus be:
+ * ----------- ------ ----- ---- --
+ * act_vec = [i0,i1,i0,i1,i0,i1,...,i0,i1] - 2 elements of the input channels broadcasted 32 times
Review Comment:
I think the mapping is `lambda n, h, w, c: [n, h//8, w//4, c//32, h%8, (w%4)//2, c%32, w%2]`
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
+ * @param zero_block A block filled with zeros
+ *
+ * @return
+ */
+void conv_layer_fp16_hvx(DLTensor& cr_out, const DLTensor& cr_act, // NOLINT(*)
+ const DLTensor& cr_filt, const DLTensor& out_shape,
+ const DLTensor& act_shape, const DLTensor& bias_flat,
+ const DLTensor& filt_shape, const DLTensor& pad_shape, bool relu,
+ int stride_h, int stride_w, uintptr_t zero_block) {
+ int64_t filt_height = filt_shape.shape[0];
+ int64_t filt_width = filt_shape.shape[1];
+ int64_t filt_idepth = filt_shape.shape[2];
+ (void)filt_idepth;
+
+ DEBUG_BLOCK(int pad_top = pad_shape.shape[0]; int pad_left = pad_shape.shape[1];)
+
+ debug("filt_height=%" PRId64 ", filt_width=%" PRId64 ", filt_idepth=%" PRId64
+ ", pad_top=%d, pad_left=%d\n",
+ filt_height, filt_width, filt_idepth, pad_top, pad_left);
+
+ assert(pad_top < 8 && pad_left < 4);
+
+ DEBUG_BLOCK(int a_height = cr_act.shape[1]; int a_width = cr_act.shape[2];
+ int a_depth = cr_act.shape[3];
+
+ int w_height = cr_filt.shape[0]; int w_width = cr_filt.shape[1];
+
+ int o_depth = cr_out.shape[3]; int b_depth = bias_flat.shape[0];)
+
+ int o_height = cr_out.shape[1];
+ int o_width = cr_out.shape[2];
+
+ int out_height = out_shape.shape[1];
+ int out_width = out_shape.shape[2];
+
+ debug("a: %dx%dx%dx%d, w: %dx%dx%dx%d, o: %dx%dx%dx%d, b: %d, out_shape: %dx%d\n", 1, a_height,
+ a_width, a_depth, w_height, w_width, static_cast<int>(cr_filt.shape[2]),
+ static_cast<int>(cr_filt.shape[3]), 1, o_height, o_width, o_depth, b_depth, out_height,
+ out_width);
+
+ assert(a_depth == cr_filt.shape[2]);
+ assert(o_depth == cr_filt.shape[3]);
+
+ int rd = round_down(filt_width, 4);
+ int wgt_chunk_thin_width = filt_width - rd;
+
+ /*
+ * Compute the output vector of either 1 or 2 elements along the width and max 32 elements along
+ * the depth to constitue a maximum of 64 elements
+ *
+ * The weights are loaded directly in the order they're stored, which results
+ * in 2 input channels and 32 output channels
+ *
+ * Weights vector illustration:
+ * ------- ------ ------------
+ * weights_vec = [0-0,0-1,1-0,1-1,2-0,2-1,3-0,3-1,4-0,4-1,...,31-0,31-1] -> This is the
Review Comment:
It will be good to document how this packing should be represented in terms of an index map which will be required for integration with TVM scheduling. For example,
`weight_index_map = lambda h, w, i, o: [h, w, i//16, o//32, (i%16)//2, o%32, i%2]`
Not sure if this is the correct mapping for FP16 please correct it if not and if possible include a similar mapping for the activation in the comments.
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
Review Comment:
Neat
##########
src/runtime/hexagon/ops/conv2d.h:
##########
@@ -0,0 +1,144 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_farf.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <cassert>
+
+#ifndef TVM_RUNTIME_HEXAGON_OPS_CONV2D_H_
+#define TVM_RUNTIME_HEXAGON_OPS_CONV2D_H_
+
+#ifdef DEBUG_CONV
+#define DEBUG_BLOCK(X) \
+ { X }
+#define debug(...) FARF(ALWAYS, ##__VA_ARGS__)
+#else
+#define DEBUG_BLOCK(X)
+#define debug(...)
+#endif
+
+#define HAP_CALL(hap_fn, ...) \
+ { \
+ int rc = hap_fn(__VA_ARGS__); \
+ if (rc != 0) { \
+ debug("%s failed: rc=%x", #hap_fn, rc); \
+ } \
+ }
Review Comment:
Agreed with @cconvey here, and hexagon_common.cc implements similar functionality into TVM's runtime logger which we should aim to use for debugging in the runtime. This can be done with `LOG(INFO/WARNING/FATAL) << "message"` which is automatically directed to FARF.
##########
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:
##########
@@ -0,0 +1,256 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import numpy as np
+import pytest
+import sys
+
+import tvm
+import tvm.contrib.hexagon
+from tvm.topi.testing import conv2d_nhwc_python
+
+
+def build_conv2d(target):
+ an, ah, aw, ac = (
+ tvm.te.var("an"),
+ tvm.te.var("ah"),
+ tvm.te.var("aw"),
+ tvm.te.var("ac"),
+ )
+ fh, fw, fo = tvm.te.var("fh"), tvm.te.var("fw"), tvm.te.var("fo")
+ off_l, off_t = tvm.te.var("off_l"), tvm.te.var("off_t")
+ stride_h, stride_w = tvm.te.var("stride_h"), tvm.te.var("stride_w")
+
+ act_flat = tvm.te.placeholder(shape=(an, ah, aw, ac), dtype="float16", name="act_flat")
+ wgt_flat = tvm.te.placeholder(shape=(fh, fw, ac, fo), dtype="float16", name="wgt_flat")
+
+ out_flat = tvm.te.extern(
+ shape=(an, (ah - fh) // stride_h + 1, (aw - fw) // stride_w + 1, fo),
+ inputs=[act_flat, wgt_flat],
+ fcompute=lambda ins, outs: tvm.tir.call_cpacked(
+ "conv2d_packed", # Function from TVM runtime
+ ins[0],
+ ins[1],
+ off_t,
+ off_l,
+ stride_h,
+ stride_w,
+ outs[0],
+ tvm.runtime.const(0), # resource_handle (unused)
+ ),
+ dtype="float16",
+ )
+
+ s = tvm.te.create_schedule(out_flat.op)
+
+ func_name = "extern_conv"
+ with tvm.transform.PassContext(opt_level=3):
+ module = tvm.build(
+ s,
+ [act_flat, wgt_flat, off_t, off_l, stride_h, stride_w, out_flat],
+ target=target,
+ name=func_name,
+ )
+
+ return module
+
+
+shape_parameters = [
+ {
+ "act_shape": (1, 8, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 10, 14, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (2, 2, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 64),
+ "wgt_shape": (3, 3, 64, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 4, 4, 40),
+ "wgt_shape": (3, 3, 40, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 4, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 5, 5, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 6, 6, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 7, 7, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 64),
+ "wgt_shape": (2, 2, 64, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (2, 2, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+]
+
+
+def gen_id(param):
+ """Utility function to generate useful ids for shape_parameters"""
+
+ dims = lambda vals: "x".join(map(str, vals))
+
+ act_shape = param["act_shape"]
+ wgt_shape = param["wgt_shape"]
+ inp_stride = param["inp_stride"]
+ return f"nhwc{dims(act_shape)}-hwio{dims(wgt_shape)}-stride{dims(inp_stride)}"
+
+
+@tvm.testing.requires_hexagon
+@pytest.mark.parametrize("shapes", shape_parameters, ids=map(gen_id, shape_parameters))
Review Comment:
After using tvm.testing.parameter this decoration can be deleted and test_conv2d can take these params as function arguments directly.
##########
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:
##########
@@ -0,0 +1,256 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import numpy as np
+import pytest
+import sys
+
+import tvm
+import tvm.contrib.hexagon
+from tvm.topi.testing import conv2d_nhwc_python
+
+
+def build_conv2d(target):
+ an, ah, aw, ac = (
+ tvm.te.var("an"),
+ tvm.te.var("ah"),
+ tvm.te.var("aw"),
+ tvm.te.var("ac"),
+ )
+ fh, fw, fo = tvm.te.var("fh"), tvm.te.var("fw"), tvm.te.var("fo")
+ off_l, off_t = tvm.te.var("off_l"), tvm.te.var("off_t")
+ stride_h, stride_w = tvm.te.var("stride_h"), tvm.te.var("stride_w")
+
+ act_flat = tvm.te.placeholder(shape=(an, ah, aw, ac), dtype="float16", name="act_flat")
+ wgt_flat = tvm.te.placeholder(shape=(fh, fw, ac, fo), dtype="float16", name="wgt_flat")
+
+ out_flat = tvm.te.extern(
+ shape=(an, (ah - fh) // stride_h + 1, (aw - fw) // stride_w + 1, fo),
+ inputs=[act_flat, wgt_flat],
+ fcompute=lambda ins, outs: tvm.tir.call_cpacked(
+ "conv2d_packed", # Function from TVM runtime
+ ins[0],
+ ins[1],
+ off_t,
+ off_l,
+ stride_h,
+ stride_w,
+ outs[0],
+ tvm.runtime.const(0), # resource_handle (unused)
+ ),
+ dtype="float16",
+ )
+
+ s = tvm.te.create_schedule(out_flat.op)
+
+ func_name = "extern_conv"
+ with tvm.transform.PassContext(opt_level=3):
+ module = tvm.build(
+ s,
+ [act_flat, wgt_flat, off_t, off_l, stride_h, stride_w, out_flat],
+ target=target,
+ name=func_name,
+ )
+
+ return module
+
+
+shape_parameters = [
+ {
+ "act_shape": (1, 8, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 10, 14, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (2, 2, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 14, 6, 64),
+ "wgt_shape": (3, 3, 64, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 4, 4, 40),
+ "wgt_shape": (3, 3, 40, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 4, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 5, 5, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 6, 6, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 7, 7, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 8, 64),
+ "wgt_shape": (2, 2, 64, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (1, 1),
+ },
+ {
+ "act_shape": (1, 8, 4, 3),
+ "wgt_shape": (3, 3, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (3, 3, 3, 64),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 14, 6, 3),
+ "wgt_shape": (5, 5, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
+ {
+ "act_shape": (1, 8, 8, 3),
+ "wgt_shape": (2, 2, 3, 3),
+ "inp_offset": (0, 0),
+ "inp_stride": (2, 2),
+ },
Review Comment:
These are all very small convolutions. Can we add some that are representative of real workloads and mark them with the `slow` decorator, or skip.
##########
src/runtime/hexagon/ops/conv2d.h:
##########
@@ -0,0 +1,144 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_farf.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <cassert>
+
+#ifndef TVM_RUNTIME_HEXAGON_OPS_CONV2D_H_
+#define TVM_RUNTIME_HEXAGON_OPS_CONV2D_H_
+
+#ifdef DEBUG_CONV
+#define DEBUG_BLOCK(X) \
+ { X }
+#define debug(...) FARF(ALWAYS, ##__VA_ARGS__)
+#else
+#define DEBUG_BLOCK(X)
+#define debug(...)
+#endif
+
+#define HAP_CALL(hap_fn, ...) \
+ { \
+ int rc = hap_fn(__VA_ARGS__); \
+ if (rc != 0) { \
+ debug("%s failed: rc=%x", #hap_fn, rc); \
+ } \
+ }
+
+namespace detail {
+static constexpr auto hexagon_device = DLDevice{static_cast<DLDeviceType>(kDLHexagon), 0};
+
+// Standalone DLTensor: the standalone-ness means that this object owns the shape
+// (as opposed to a DLTensor).
+template <size_t N>
+class SDLTensor : public DLTensor {
+ public:
+ SDLTensor(void* data_ptr, DLDataType data_type, void* data_space, const int64_t* data_dims)
+ : SDLTensor(data_ptr, data_type, data_space) {
+ for (size_t i = 0; i != N; ++i) dims[i] = data_dims[i];
+ }
+
+ SDLTensor(void* data_ptr, DLDataType data_type, void* data_space,
+ std::initializer_list<int64_t> data_dims)
+ : SDLTensor(data_ptr, data_type, data_space, data_dims.begin()) {}
+
+ void* GetDataSpace() const { return data_space; }
+
+ private:
+ SDLTensor(void* data_ptr, DLDataType data_type, void* data_space) : data_space(data_space) {
+ data = data_ptr;
+ device = hexagon_device;
+ ndim = N;
+ dtype = data_type;
+ shape = dims;
+ strides = nullptr;
+ byte_offset = 0;
+ }
+
+ void* data_space = nullptr;
+ int64_t dims[N];
+};
+
+inline void* to_ptr(uintptr_t v) { return reinterpret_cast<void*>(v); }
+
+inline uintptr_t to_uint(void* ptr) { return reinterpret_cast<uintptr_t>(ptr); }
+
+inline constexpr int xyc_to_sm_16b(int y, int x, int c) {
+ // Map y,x,c coordinates within a block to the offset (in 16-bit elements)
+ // from the beginning of the block in spatial-major layout.
+ // 10-bit spatial mask: yyyxcccccx
+ return y << 7 | (x & 2) << 5 | c << 1 | (x & 1);
+}
+
+inline constexpr int hwio_to_sm_16b(int width, int y, int x, int i, int o) {
+ // Map y,x,i,o coordinates within a chunk (assuming the origin at the
+ // top-left spatial corner) to the offset (in 16-bit elements) from the
+ // beginning of the chunk in spatial-major layout.
+ // Spatial mask: p..piiiioooooi, where p..p are position bits.
+ int p = y * width + (width - 1 - x);
+ return p << 10 | (i & 0x1e) << 5 | o << 1 | (i & 1);
+}
+
+inline constexpr int round_up(int v, int p2) { return (v + p2 - 1) & -p2; }
+
+constexpr uintptr_t nhwc_at(const DLTensor& a, int n, int y, int x, int c) {
+ if (y < 0 || y >= a.shape[1]) return uintptr_t(0);
+ auto p = static_cast<uintptr_t*>(a.data);
+ assert(n == 0);
+ return p[y * a.shape[2] * a.shape[3] + x * a.shape[3] + c];
+}
+
+constexpr uintptr_t hwio_at(const DLTensor& f, int y, int x, int i, int o) {
+ auto p = static_cast<uintptr_t*>(f.data);
+ return p[y * f.shape[1] * f.shape[2] * f.shape[3] + x * f.shape[2] * f.shape[3] + i * f.shape[3] +
+ o];
+}
+
+constexpr uint32_t* bias_at(const DLTensor& b, int d) {
+ auto p = static_cast<uint32_t*>(b.data);
+ return p + d;
+}
+
+void blockize_hwc_16b(void* out, void* inp_flat, int height, int width, int depth);
+
+void deblockize_hwc_16b(void* out_flat, void* inp, int height, int width, int depth);
+
+void chunkify_hwio_16b(void** out_ptr, int out_ptr_size, void* out, void* inp, int height,
+ int width, int idepth, int odepth);
Review Comment:
Agreed, I made a comment about this above and after reading further I think even my definitions are not accurate.
##########
tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py:
##########
@@ -0,0 +1,256 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+import numpy as np
+import pytest
+import sys
+
+import tvm
+import tvm.contrib.hexagon
+from tvm.topi.testing import conv2d_nhwc_python
+
+
+def build_conv2d(target):
+ an, ah, aw, ac = (
+ tvm.te.var("an"),
+ tvm.te.var("ah"),
+ tvm.te.var("aw"),
+ tvm.te.var("ac"),
+ )
+ fh, fw, fo = tvm.te.var("fh"), tvm.te.var("fw"), tvm.te.var("fo")
+ off_l, off_t = tvm.te.var("off_l"), tvm.te.var("off_t")
+ stride_h, stride_w = tvm.te.var("stride_h"), tvm.te.var("stride_w")
+
+ act_flat = tvm.te.placeholder(shape=(an, ah, aw, ac), dtype="float16", name="act_flat")
+ wgt_flat = tvm.te.placeholder(shape=(fh, fw, ac, fo), dtype="float16", name="wgt_flat")
+
+ out_flat = tvm.te.extern(
+ shape=(an, (ah - fh) // stride_h + 1, (aw - fw) // stride_w + 1, fo),
+ inputs=[act_flat, wgt_flat],
+ fcompute=lambda ins, outs: tvm.tir.call_cpacked(
+ "conv2d_packed", # Function from TVM runtime
+ ins[0],
+ ins[1],
+ off_t,
+ off_l,
+ stride_h,
+ stride_w,
+ outs[0],
+ tvm.runtime.const(0), # resource_handle (unused)
+ ),
+ dtype="float16",
+ )
+
+ s = tvm.te.create_schedule(out_flat.op)
+
+ func_name = "extern_conv"
+ with tvm.transform.PassContext(opt_level=3):
+ module = tvm.build(
+ s,
+ [act_flat, wgt_flat, off_t, off_l, stride_h, stride_w, out_flat],
+ target=target,
+ name=func_name,
+ )
+
+ return module
+
+
+shape_parameters = [
Review Comment:
Instead using tvm's helpers, you can either do
```
act_shape, wgt_shape, inp_offset, inp_stride = tvm.testing.parameters(((1,8,4,3),(3,3,3,3),(0,0),(1,1)), ((1,10,14,3), (3,3,3,3), (0,0), (1,1)), ...)
```
or
```
act_shape = tvm.testing.parameter((1,8,4,3), (1,10,14,3), ...)
wgt_shape = tvm.testing.parameter((3,3,3,3), (3,3,3,64), ...)
inp_offset = tvm.testing_parameter((0, 0))
inp_stride = tvm.testing_parameter((1,1), (2,2))
```
which will do the cartesian product of the individually set parameters.
```
##########
src/runtime/hexagon/ops/conv2d_hvx.cc:
##########
@@ -0,0 +1,468 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <HAP_compute_res.h>
+#include <hexagon_types.h>
+#include <hvx_hexagon_protos.h>
+#include <tvm/runtime/c_runtime_api.h>
+#include <tvm/runtime/device_api.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cinttypes>
+
+#include "conv2d.h"
+
+// Current limitations:
+// - N in NHWC must be 1
+// - dilated convolutions are not supported
+// - Bias is not accepted
+// - Optional "relu" is not performed
+
+// Packed arguments:
+// 0: DLTensor activations (NHWC)
+// 1: DLTensor weights (HWIO)
+// 2: int offset_top
+// 3: int offset_left
+// 4: int stride_h
+// 5: int stride_w
+// 6: DLTensor output (NHWC)
+extern "C" int conv2d_packed(TVMValue* args, int* type_codes, int num_args, TVMValue* out_val,
+ int out_code, void* res_handle);
+
+namespace detail {
+
+inline uint16_t* getElementPtr(int block_out_y, int block_out_x, int block_out_c, int yi, int xio,
+ int ci, int xii, const DLTensor& block) {
+ auto block_ptr = nhwc_at(block, 0, block_out_y, block_out_x, block_out_c);
+ auto block_offset = yi * 128 + xio * 64 + ci * 2 + xii;
+ auto first_element_ptr = reinterpret_cast<uint16_t*>(block_ptr);
+ return first_element_ptr + block_offset;
+}
+
+/**
+ * @brief Compute 2 vectors with ones in the even and odd lanes
+ *
+ * Output vectors are:
+ * vector 1 = [0xFFFF,0x0000,0xFFFFF,0x0000,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * vector 2 = [0x0000,0xFFFF,0x0000,0xFFFFF,...,0xFFFF,0x0000]
+ * vector lanes = [ 0 , 2 , 3 , 4 ,..., 62 , 63 ]
+ *
+ * @return Return the 2 vectors
+ */
+inline std::pair<HVX_Vector, HVX_Vector> getOddEvenOnes() {
+ HVX_Vector v0 = Q6_V_vzero();
+ HVX_Vector v1 = Q6_Vh_vsplat_R(0xFFFF);
+
+ HVX_Vector v1e = Q6_Vh_vshuffe_VhVh(v0, v1);
+ HVX_Vector v1o = Q6_V_vnot_V(v1e);
+ return {v1e, v1o};
+}
+
+/**
+ * @brief Return the input vector filled with the 2 channel elements(which is the 1st and 3rd
+ * element) from base_ptr filled up 32 times to get 64 elements
+ *
+ * 1. It's generated by first creating 2 vectors "splatted" with the 2 required elements
+ * 2. Then we andd it with vectors containing all ones (0xFFFF) in the even and odd lanes
+ * 3. Finally those 2 vectors are OR'ed together
+ *
+ * @param base_ptr pointer to the first of the 2 channel elements to be filled
+ *
+ * @return input vector
+ */
+inline HVX_Vector getInputVector(uint16_t* base_ptr) {
+ HVX_Vector v1 = Q6_Vh_vsplat_R(base_ptr[0]);
+ HVX_Vector v2 = Q6_Vh_vsplat_R(base_ptr[2]);
+
+ auto oddEvenOnes = getOddEvenOnes();
+ auto v1e = oddEvenOnes.first;
+ auto v1o = oddEvenOnes.second;
+
+ HVX_Vector v_even_vals = Q6_V_vand_VV(v1, v1e);
+ HVX_Vector v_odd_vals = Q6_V_vand_VV(v2, v1o);
+
+ return Q6_V_vor_VV(v_even_vals, v_odd_vals);
+}
+
+/**
+ * @brief Return the Output vector which contains the 32 output channels in the even lanes
+ *
+ * The output vector is commputed as:
+ * 1. vector multiply(vmpy) of input and weights
+ * 2. Rotate the vector right by 1 element and add with the first vector to add the 2 input channels
+ * 3. Then convert the results back from qfloat16 to IEEE half-precision float
+ * 4. The added values are in even lanes, so zero out the odd lanes by anding with ones in even
+ * lanes and return
+ *
+ * @param act_vec Input activations vector
+ * @param wgt_vec Weights vector
+ *
+ * @return output vector with 32 output channels even lanes
+ */
+inline HVX_Vector computeOuputVector(HVX_Vector act_vec, HVX_Vector wgt_vec) {
+ HVX_Vector v_res = Q6_Vqf16_vmpy_VhfVhf(act_vec, wgt_vec); // result is in qf16
+ HVX_Vector v_rot = Q6_V_vror_VR(v_res, 2);
+ HVX_Vector v_reduced = Q6_Vqf16_vadd_Vqf16Vqf16(v_res, v_rot);
+ HVX_Vector v_hf = Q6_Vhf_equals_Vqf16(v_reduced);
+ HVX_Vector v1e = getOddEvenOnes().first;
+ HVX_Vector v_reduced_even_lanes = Q6_V_vand_VV(v_hf, v1e);
+ return v_reduced_even_lanes;
+}
+
+static int round_down(int v, int base) { return v - (v % base); }
+
+/**
+ * @brief Compute the convolution of inputs from cr_act, and weights from
+ * cr_filt to update the output to cr_out. The goal is to have an efficient
+ * HVX implementation
+ *
+ * Assumptions:
+ * -----------
+ * - This implementation right now assumes that the dilation is 1
+ * - there is zero padding or the input was already pre-padded.
+ * - block specific spatial padding is only expected at the end and hence
+ * pad_top and pad_left are not yet used
+ * - Relu activation is not used
+ * - Bias add is not done
+ *
+ * @param cr_out blockized output tensor with zeros already filled in
+ * @param cr_act blockized activations
+ * @param cr_filt Chunkified weights as returned from output of prepare_hwio
+ * @param out_shape Original output shape of the tensor before blockization
+ * @param act_shape Original input shape
+ * @param bias_flat Flat bias values and are not used right now
+ * @param filt_shape Original filter shape
+ * @param pad_shape Pad top and pad left shape
+ * @param relu Whether to apply relu after convolution, not done right now
Review Comment:
Please add a `# TODO(author): Add support for relu` and similar for all the params which are not supported but planned to be.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org