From db00668890e1aba956e02fa02e1383b54dfd1435 Mon Sep 17 00:00:00 2001 From: steniu01 Date: Wed, 9 Aug 2017 16:26:22 +0100 Subject: COMPMID-478 Implemnt CL direct convolution 5x5 Change-Id: I4b975aff310cda9964d8c5dcee182d5d5c82741b Reviewed-on: http://mpd-gerrit.cambridge.arm.com/83474 Tested-by: Kaizen Reviewed-by: Gian Marco Iodice --- .../CL/kernels/CLDirectConvolutionLayerKernel.h | 1 + src/core/CL/CLKernelLibrary.cpp | 5 + src/core/CL/cl_kernels/direct_convolution1x1.cl | 3 +- src/core/CL/cl_kernels/direct_convolution3x3.cl | 10 +- src/core/CL/cl_kernels/direct_convolution5x5.cl | 149 +++++++++++++++++++++ .../CL/kernels/CLDirectConvolutionLayerKernel.cpp | 12 +- tests/datasets_new/ShapeDatasets.h | 2 +- tests/validation_new/CL/DirectConvolutionLayer.cpp | 17 ++- 8 files changed, 183 insertions(+), 16 deletions(-) create mode 100644 src/core/CL/cl_kernels/direct_convolution5x5.cl diff --git a/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h b/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h index aa6ecd6631..e225b64bae 100644 --- a/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h +++ b/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h @@ -53,6 +53,7 @@ public: * @note: DirectConvolution only works in the following configurations: * 1x1 convolution with stride_x = 1/2/3, stride_y = 1/2/3 * 3x3 convolution with stride_x = 1/2, stride_y = 1/2 + * 5x5 convolution with stride_x = 1/2, stride_y = 1/2 * * @param[in] input The input tensor to convolve. 3 lower dimensions represent a single input [width, height, IFM], * while every optional dimension from 4 and above represent a batch of inputs. Data types supported: F16/F32. diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 435e19a22b..1647a37ce0 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -147,6 +147,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "dilate", "dilate.cl" }, { "direct_convolution1x1", "direct_convolution1x1.cl" }, { "direct_convolution3x3", "direct_convolution3x3.cl" }, + { "direct_convolution5x5", "direct_convolution5x5.cl" }, { "erode", "erode.cl" }, { "fast_corners", "fast_corners.cl" }, { "fill_image_borders_constant", "fill_border.cl" }, @@ -358,6 +359,10 @@ const std::map CLKernelLibrary::_program_source_map = { "direct_convolution3x3.cl", #include "./cl_kernels/direct_convolution3x3.clembed" + }, + { + "direct_convolution5x5.cl", +#include "./cl_kernels/direct_convolution5x5.clembed" }, { "erode.cl", diff --git a/src/core/CL/cl_kernels/direct_convolution1x1.cl b/src/core/CL/cl_kernels/direct_convolution1x1.cl index 66c618e033..2aa999a80f 100644 --- a/src/core/CL/cl_kernels/direct_convolution1x1.cl +++ b/src/core/CL/cl_kernels/direct_convolution1x1.cl @@ -33,6 +33,7 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) #else /* FIXED_POINT_POSITION */ +#undef CONVERT_SAT #define ADD_OP(a, b) ((a) + (b)) #define MUL_OP(a, b) ((a) * (b)) @@ -205,4 +206,4 @@ __kernel void direct_convolution1x1( vstore8(CONVERT_SAT(pixels, VEC_DATA_TYPE(DATA_TYPE, 8)), 0, (__global DATA_TYPE *)dst.ptr); } -#endif // defined(DATA_TYPE) && defined(DATA_SIZE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) \ No newline at end of file +#endif // defined(DATA_TYPE) && defined(DATA_SIZE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) diff --git a/src/core/CL/cl_kernels/direct_convolution3x3.cl b/src/core/CL/cl_kernels/direct_convolution3x3.cl index 4da7c39e26..28da544f89 100644 --- a/src/core/CL/cl_kernels/direct_convolution3x3.cl +++ b/src/core/CL/cl_kernels/direct_convolution3x3.cl @@ -50,8 +50,8 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) #define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ ({ \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = vload4(0, weights_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 3) \ + weights_values0 = vload3(0, weights_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 8) \ src0 = vload8(0, src_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 2) \ @@ -64,8 +64,8 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) #define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ ({ \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = vload4(0, weights_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 3) \ + weights_values0 = vload3(0, weights_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 16) \ src0 = vload16(0, src_row_ptr); \ DATA_TYPE src1 = *(src_row_ptr + 16); \ @@ -152,4 +152,4 @@ __kernel void direct_convolution3x3( vstore8(CONVERT_SAT(pixels0, VEC_DATA_TYPE(DATA_TYPE, 8)), 0, (__global DATA_TYPE *)dst.ptr); } -#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) \ No newline at end of file +#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) diff --git a/src/core/CL/cl_kernels/direct_convolution5x5.cl b/src/core/CL/cl_kernels/direct_convolution5x5.cl new file mode 100644 index 0000000000..d8c0d891d7 --- /dev/null +++ b/src/core/CL/cl_kernels/direct_convolution5x5.cl @@ -0,0 +1,149 @@ +/* + * Copyright (c) 2016, 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +#undef CONVERT_SAT + +#if STRIDE_X == 1 +#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) +#elif STRIDE_X == 2 /* STRIDE_X == 1 */ +#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) +#else /* STRIDE_X not equals 1 or 2 */ +#error "STRIDE_X larger than 2 is not supported" +#endif /* STRIDE_X == 2 */ + +#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + VEC_DATA_TYPE(DATA_TYPE, 4) \ + weights_values0 = vload4(0, weights_row_ptr); \ + DATA_TYPE weights_value1 = *(weights_row_ptr + 4); \ + VEC_DATA_TYPE(DATA_TYPE, 8) \ + src0 = vload8(0, src_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 4) \ + src1 = vload4(0, src_row_ptr + 8); \ + \ + acc += src0 * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1234, src0.s567, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s234, src0.s567, src1.s01) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s345, src0.s67, src1.s012) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s45, src0.s67, src1.s0123) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ + }) + +#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + VEC_DATA_TYPE(DATA_TYPE, 4) \ + weights_values0 = vload4(0, weights_row_ptr); \ + DATA_TYPE weights_value1 = *(weights_row_ptr + 4); \ + VEC_DATA_TYPE(DATA_TYPE, 16) \ + src0 = vload16(0, src_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 4) \ + src1 = vload4(0, src_row_ptr + 16); \ + acc += src0.even * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1357, src0.s9BDF) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s2468, src0.sACE, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ + \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s3579, src0.sBDF, src1.s1) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ + acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s468a, src0.sCE, src1.s02) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ + }) + +/** This kernel performs a direct convolution to convolve the low three dimensions. + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH + * @note In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row. + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 + * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p weights_ptr + * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) + * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) + * @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes) + * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor + * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr + * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) + * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor + * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension + */ +#if defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) +__kernel void direct_convolution5x5( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + TENSOR3D_DECLARATION(weights), +#ifdef HAS_BIAS + VECTOR_DECLARATION(biases), +#endif /* defined(HAS_BIAS) */ + unsigned int weights_stride_w) +{ + Image src = CONVERT_TO_IMAGE_STRUCT(src); + Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + + VEC_DATA_TYPE(DATA_TYPE, 8) + pixels0 = 0; + + __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + const int kernel_index = get_global_id(2); + weights_addr += kernel_index * weights_stride_w; + + for(int d = 0; d < WEIGHTS_DEPTH; ++d) + { + CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)src_addr, (__global DATA_TYPE *)weights_addr); + CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y)); + + src_addr += src_stride_z; + weights_addr += weights_stride_z; + } + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + + pixels0 += (VEC_DATA_TYPE(DATA_TYPE, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, kernel_index))); +#endif /* defined(HAS_BIAS) */ + + vstore8(pixels0, 0, (__global DATA_TYPE *)dst.ptr); +} +#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) diff --git a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp index c5fdb77a4a..1620d545c7 100644 --- a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp +++ b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp @@ -53,14 +53,14 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QS8, DataType::QS16, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); ARM_COMPUTE_ERROR_ON_MSG(weights->info()->dimension(0) != weights->info()->dimension(1), - "Only kernel sizes 1x1 and 3x3 are supported"); - ARM_COMPUTE_ERROR_ON_MSG(weights->info()->dimension(0) != 1 && weights->info()->dimension(0) != 3, - "Only kernel sizes 1x1 and 3x3 are supported"); + "Weights should have same width as length"); + ARM_COMPUTE_ERROR_ON_MSG(weights->info()->dimension(0) != 1 && weights->info()->dimension(0) != 3 && weights->info()->dimension(0) != 5, + "Kernel sizes other than 1x1, 3x3 or 5x5 are not supported"); ARM_COMPUTE_ERROR_ON(weights->info()->dimension(2) != input->info()->dimension(2)); ARM_COMPUTE_ERROR_ON(weights->info()->dimension(0) != weights->info()->dimension(1)); ARM_COMPUTE_ERROR_ON(weights->info()->num_dimensions() > 4); ARM_COMPUTE_ERROR_ON_MSG((weights->info()->dimension(0) == 1) && std::get<0>(conv_info.stride()) > 3, "Strides larger than 3 not supported for 1x1 convolution."); - ARM_COMPUTE_ERROR_ON_MSG((weights->info()->dimension(0) == 3) && std::get<0>(conv_info.stride()) > 2, "Strides larger than 2 not supported for 3x3 convolution."); + ARM_COMPUTE_ERROR_ON_MSG((weights->info()->dimension(0) == 3 || weights->info()->dimension(0) == 5) && std::get<0>(conv_info.stride()) > 2, "Strides larger than 2 not supported for 3x3 convolution."); if(biases != nullptr) { @@ -138,9 +138,9 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL // Configure kernel window Window win = calculate_max_window(*output->info()); - bool is_kernel3x3_stride2 = ((kernel_size == 3) && (_conv_stride_x == 2)); + bool is_stride2 = ((kernel_size != 1) && (_conv_stride_x == 2)); - const unsigned int num_elems_read_per_iteration_x = 8 + 2 * (kernel_size / 2) + (is_kernel3x3_stride2 ? 7 : 0); + const unsigned int num_elems_read_per_iteration_x = 8 + 2 * (kernel_size / 2) + (is_stride2 ? 6 + kernel_size / 2 : 0); const unsigned int num_elems_read_per_iteration_y = kernel_size; const unsigned int num_elems_written_per_iteration_x = 8; const unsigned int num_elems_written_per_iteration_y = 1; diff --git a/tests/datasets_new/ShapeDatasets.h b/tests/datasets_new/ShapeDatasets.h index 14f7851621..f6cd3f2d0e 100644 --- a/tests/datasets_new/ShapeDatasets.h +++ b/tests/datasets_new/ShapeDatasets.h @@ -115,7 +115,7 @@ public: SmallDirectConvolutionShapes() : ShapeDataset("InputShape", { - TensorShape{ 3U, 3U, 3U, 2U, 4U, 5U }, + TensorShape{ 5U, 5U, 3U, 2U, 4U, 5U }, TensorShape{ 32U, 37U, 3U }, TensorShape{ 13U, 15U, 8U, 3U } }) diff --git a/tests/validation_new/CL/DirectConvolutionLayer.cpp b/tests/validation_new/CL/DirectConvolutionLayer.cpp index d82f535136..1c698ace0f 100644 --- a/tests/validation_new/CL/DirectConvolutionLayer.cpp +++ b/tests/validation_new/CL/DirectConvolutionLayer.cpp @@ -50,6 +50,17 @@ constexpr AbsoluteTolerance tolerance_qs8(0); /**< Tolerance for fixed constexpr AbsoluteTolerance tolerance_qs16(0); /**< Tolerance for fixed point tests */ /** Direct convolution data set. */ +const auto data_quantized = combine(datasets::SmallDirectConvolutionShapes(), + combine(framework::dataset::make("StrideX", 1, 3), + combine(framework::dataset::make("StrideY", 1, 3), + combine(concat(combine(framework::dataset::make("PadX", 0), + combine(framework::dataset::make("PadY", 0), + framework::dataset::make("KernelSize", 1))), + combine(framework::dataset::make("PadX", 0, 2), + combine(framework::dataset::make("PadY", 0, 2), + framework::dataset::make("KernelSize", { 3 })))), + framework::dataset::make("NumKernels", { 1, 4, 8, 16 }))))); + const auto data = combine(datasets::SmallDirectConvolutionShapes(), combine(framework::dataset::make("StrideX", 1, 3), combine(framework::dataset::make("StrideY", 1, 3), @@ -58,7 +69,7 @@ const auto data = combine(datasets::SmallDirectConvolutionShapes(), framework::dataset::make("KernelSize", 1))), combine(framework::dataset::make("PadX", 0, 2), combine(framework::dataset::make("PadY", 0, 2), - framework::dataset::make("KernelSize", 3)))), + framework::dataset::make("KernelSize", { 3, 5 })))), framework::dataset::make("NumKernels", { 1, 4, 8, 16 }))))); } // namespace @@ -93,7 +104,7 @@ using CLDirectConvolutionLayerFixedPointFixture = DirectConvolutionValidationFix TEST_SUITE(Quantized) TEST_SUITE(QS8) -FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture, framework::DatasetMode::ALL, combine(combine(data, framework::dataset::make("DataType", DataType::QS8)), +FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS8)), framework::dataset::make("FractionalBits", 2, 7))) { // Validate output @@ -102,7 +113,7 @@ FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture, f TEST_SUITE_END() TEST_SUITE(QS16) -FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture, framework::DatasetMode::ALL, combine(combine(data, framework::dataset::make("DataType", DataType::QS16)), +FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS16)), framework::dataset::make("FractionalBits", 2, 15))) { // Validate output -- cgit v1.2.1